asc_ldca

产 品 支 持 情 况

• Ascend 950PR/Ascend 950DT：支 持
• Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品：不 支 持
• Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品：不 支 持
• Atlas 200I/500 A2 推 理 产 品：不 支 持
• Atlas 推 理 系 列 产 品AI Core：不 支 持
• Atlas 推 理 系 列 产 品Vector Core：不 支 持
• Atlas 训 练 系 列 产 品：不 支 持

功 能 说 明

首 先 从Data Cache加 载 缓 存 数 据，若 未 命 中，则 尝 试 从L2 Cache加 载。L2 Cache与Global Memory之 间 的 数 据 一 致 性 由 硬 件 保 证。如 果Data Cache和L2 Cache中 均 未 找 到 所 需 数 据，则 从Global Memory中 读 取 数 据，然 后 将 其 缓 存 到L2 Cache和Data Cache中。

函 数 原 型

Text

inline long int asc_ldca(long int* address)

Text

inline unsigned long int asc_ldca(unsigned long int* address)

Text

inline long long int asc_ldca(long long int* address)

Text

inline unsigned long long int asc_ldca(unsigned long long int* address)

Text

inline long2 asc_ldca(long2* address)

Text

inline ulong2 asc_ldca(ulong2* address)

Text

inline long4 asc_ldca(long4* address)

Text

inline ulong4 asc_ldca(ulong4* address)

Text

inline longlong2 asc_ldca(longlong2* address)

Text

inline ulonglong2 asc_ldca(ulonglong2* address)

Text

inline longlong4 asc_ldca(longlong4* address)

Text

inline ulonglong4 asc_ldca(ulonglong4* address)

Text

inline signed char asc_ldca(signed char* address)

Text

inline unsigned char asc_ldca(unsigned char* address)

Text

inline char2 asc_ldca(char2* address)

Text

inline uchar2 asc_ldca(uchar2* address)

Text

inline char4 asc_ldca(char4* address)

Text

inline uchar4 asc_ldca(uchar4* address)

Text

inline short asc_ldca(short* address)

Text

inline unsigned short asc_ldca(unsigned short* address)

Text

inline short2 asc_ldca(short2* address)

Text

inline ushort2 asc_ldca(ushort2* address)

Text

inline short4 asc_ldca(short4* address)

Text

inline ushort4 asc_ldca(ushort4* address)

Text

inline int asc_ldca(int* address)

Text

inline unsigned int asc_ldca(unsigned int* address)

Text

inline int2 asc_ldca(int2* address)

Text

inline uint2 asc_ldca(uint2* address)

Text

inline int4 asc_ldca(int4* address)

Text

inline uint4 asc_ldca(uint4* address)

Text

inline float asc_ldca(float* address)

Text

inline float2 asc_ldca(float2* address)

Text

inline float4 asc_ldca(float4* address)

Text

inline bfloat16_t asc_ldca(bfloat16_t* address)

Text

inline bfloat16x2_t asc_ldca(bfloat16x2_t* address)

Text

inline half asc_ldca(half* address)

Text

inline half2 asc_ldca(half2* address)

参 数 说 明

表 1 参 数 说 明

参 数 名	输 入/输 出	描 述
address	输 入	Global Memory的 地 址。

返 回 值 说 明

Data Cache或L2 Cache中 的 缓 存 数 据，或 输 入 的Global Memory地 址 上 的 数 据。

约 束 说 明

无

需 要 包 含 的 头 文 件

使 用 除half、half2、bfloat16_t、bfloat16x2_t类 型 之 外 的 接 口 需 要 包 含"simt_api/device_functions.h"头 文 件，使 用half和half2类 型 接 口 需 要 包 含"simt_api/asc_fp16.h"头 文 件，使 用bfloat16_t和bfloat16x2_t类 型 接 口 需 要 包 含"simt_api/asc_bf16.h"头 文 件。

Text

#include "simt_api/device_functions.h"

Text

#include "simt_api/asc_fp16.h"

Text

#include "simt_api/asc_bf16.h"

调 用 示 例

• SIMT编 程 场 景：

  Text

  __global__ __launch_bounds__(1024) void kernel_asc_ldca(float* dst, float* src)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = asc_ldca(src + idx);
  }
  

• SIMD与SIMT混 合 编 程 场 景：

  SIMD与SIMT混 合 编 程 场 景，需 要 显 式 使 用 地 址 空 间 限 定 符 表 示 地 址 空 间：__gm__表 示Global Memory内 存 空 间。

  Text

  __simt_vf__ __launch_bounds__(1024) inline void kernel_asc_ldca(__gm__ float* dst, __gm__ float* src)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = asc_ldca(src + idx);
  }