asc_ldca

产品支持情况

产品	是否支持
Ascend 950PR/Ascend 950DT	√
Atlas A3 训练系列产品/Atlas A3 推理系列产品	x
Atlas A2 训练系列产品/Atlas A2 推理系列产品	x

功能说明

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

函数原型

__simt_callee__ inline long int asc_ldca(__gm__ long int* address)
__simt_callee__ inline unsigned long int asc_ldca(__gm__ unsigned long int* address)
__simt_callee__ inline long long int asc_ldca(__gm__ long long int* address)
__simt_callee__ inline unsigned long long int asc_ldca(__gm__ unsigned long long int* address)
__simt_callee__ inline long2 asc_ldca(__gm__ long2* address)
__simt_callee__ inline ulong2 asc_ldca(__gm__ ulong2* address)
__simt_callee__ inline long4 asc_ldca(__gm__ long4* address)
__simt_callee__ inline ulong4 asc_ldca(__gm__ ulong4* address)
__simt_callee__ inline longlong2 asc_ldca(__gm__ longlong2* address)
__simt_callee__ inline ulonglong2 asc_ldca(__gm__ ulonglong2* address)
__simt_callee__ inline longlong4 asc_ldca(__gm__ longlong4* address)
__simt_callee__ inline ulonglong4 asc_ldca(__gm__ ulonglong4* address)
__simt_callee__ inline signed char asc_ldca(__gm__ signed char* address)
__simt_callee__ inline unsigned char asc_ldca(__gm__ unsigned char* address)
__simt_callee__ inline char2 asc_ldca(__gm__ char2* address)
__simt_callee__ inline uchar2 asc_ldca(__gm__ uchar2* address)
__simt_callee__ inline char4 asc_ldca(__gm__ char4* address)
__simt_callee__ inline uchar4 asc_ldca(__gm__ uchar4* address)
__simt_callee__ inline short asc_ldca(__gm__ short* address)
__simt_callee__ inline unsigned short asc_ldca(__gm__ unsigned short* address)
__simt_callee__ inline short2 asc_ldca(__gm__ short2* address)
__simt_callee__ inline ushort2 asc_ldca(__gm__ ushort2* address)
__simt_callee__ inline short4 asc_ldca(__gm__ short4* address)
__simt_callee__ inline ushort4 asc_ldca(__gm__ ushort4* address)
__simt_callee__ inline int asc_ldca(__gm__ int* address)
__simt_callee__ inline unsigned int asc_ldca(__gm__ unsigned int* address)
__simt_callee__ inline int2 asc_ldca(__gm__ int2* address)
__simt_callee__ inline uint2 asc_ldca(__gm__ uint2* address)
__simt_callee__ inline int4 asc_ldca(__gm__ int4* address)
__simt_callee__ inline uint4 asc_ldca(__gm__ uint4* address)
__simt_callee__ inline float asc_ldca(__gm__ float* address)
__simt_callee__ inline float2 asc_ldca(__gm__ float2* address)
__simt_callee__ inline float4 asc_ldca(__gm__ float4* address)
__simt_callee__ inline bfloat16_t asc_ldca(__gm__ bfloat16_t* address)
__simt_callee__ inline bfloat16x2_t asc_ldca(__gm__ bfloat16x2_t* address)
__simt_callee__ inline half asc_ldca(__gm__ half* address)
__simt_callee__ inline half2 asc_ldca(__gm__ 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"头文件。

#include "simt_api/device_functions.h"

#include "simt_api/asc_fp16.h"

#include "simt_api/asc_bf16.h"

调用示例

__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_ldcg(src + idx);
}