asc_ldca
产品支持情况
产品 |
是否支持 |
|---|---|
Atlas 350 加速卡 |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
功能说明
首先从Data Cache加载缓存数据，若未命中，则尝试从L2 Cache加载。L2 Cache与Global Memory之间的数据一致性由硬件保证。如果Data Cache和L2 Cache中均未找到所需数据，则从Global Memory中读取数据，然后将其缓存到L2 Cache和Data Cache中。
函数原型
1 | inline long int asc_ldca(long int* address) |
1 | inline unsigned long int asc_ldca(unsigned long int* address) |
1 | inline long long int asc_ldca(long long int* address) |
1 | inline unsigned long long int asc_ldca(unsigned long long int* address) |
1 | inline long2 asc_ldca(long2* address) |
1 | inline ulong2 asc_ldca(ulong2* address) |
1 | inline long4 asc_ldca(long4* address) |
1 | inline ulong4 asc_ldca(ulong4* address) |
1 | inline longlong2 asc_ldca(longlong2* address) |
1 | inline ulonglong2 asc_ldca(ulonglong2* address) |
1 | inline longlong4 asc_ldca(longlong4* address) |
1 | inline ulonglong4 asc_ldca(ulonglong4* address) |
1 | inline signed char asc_ldca(signed char* address) |
1 | inline unsigned char asc_ldca(unsigned char* address) |
1 | inline char2 asc_ldca(char2* address) |
1 | inline uchar2 asc_ldca(uchar2* address) |
1 | inline char4 asc_ldca(char4* address) |
1 | inline uchar4 asc_ldca(uchar4* address) |
1 | inline short asc_ldca(short* address) |
1 | inline unsigned short asc_ldca(unsigned short* address) |
1 | inline short2 asc_ldca(short2* address) |
1 | inline ushort2 asc_ldca(ushort2* address) |
1 | inline short4 asc_ldca(short4* address) |
1 | inline ushort4 asc_ldca(ushort4* address) |
1 | inline int asc_ldca(int* address) |
1 | inline unsigned int asc_ldca(unsigned int* address) |
1 | inline int2 asc_ldca(int2* address) |
1 | inline uint2 asc_ldca(uint2* address) |
1 | inline int4 asc_ldca(int4* address) |
1 | inline uint4 asc_ldca(uint4* address) |
1 | inline float asc_ldca(float* address) |
1 | inline float2 asc_ldca(float2* address) |
1 | inline float4 asc_ldca(float4* address) |
1 | inline bfloat16_t asc_ldca(bfloat16_t* address) |
1 | inline bfloat16x2_t asc_ldca(bfloat16x2_t* address) |
1 | inline half asc_ldca(half* address) |
1 | inline half2 asc_ldca(half2* address) |
参数说明
参数名 |
输入/输出 |
描述 |
|---|---|---|
address |
输入 |
Global Memory的地址。 |
返回值说明
Data Cache或L2 Cache中的缓存数据，或输入的Global Memory地址上的数据。
约束说明
SIMT编程场景当前不支持使用该接口。
需要包含的头文件
使用除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"头文件。
1 | #include "simt_api/device_functions.h" |
1 | #include "simt_api/asc_fp16.h" |
1 | #include "simt_api/asc_bf16.h" |
调用示例
SIMD与SIMT混合编程场景：
1 2 3 4 5 6 | //混合编程场景，使用地址空间限定符__gm__表示Global Memory内存空间。 __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); } |