asc_ldca
Applicability
Product |
Supported |
|---|---|
Atlas 350 Accelerator Card |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
Function Usage
First loads cached data from data cache. If cache miss occurs, it attempts to load from L2 cache. Data consistency between L2 cache and global memory is guaranteed by hardware. If the required data is not found in either data cache or L2 cache, it is fetched from global memory and then cached into both L2 cache and data cache.
Prototype
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) |
Parameters
Parameter |
Input/Output |
Description |
|---|---|---|
address |
Input |
Address of the Global Memory |
Returns
Cached data in data cache or L2 cache, or data from the input Global Memory address
Restrictions
For SIMT programming, this API is not supported.
Header Files to Be Included
To use APIs of types other than half, half2, bfloat16_t, and bfloat16x2_t, include the simt_api/device_functions.h header file. To use APIs of the half and half2 types, include the simt_api/asc_fp16.h header file. To use APIs of the bfloat16_t and bfloat16x2_t types, include the simt_api/asc_bf16.h header file.
1 | #include "simt_api/device_functions.h" |
1 | #include "simt_api/asc_fp16.h" |
1 | #include "simt_api/asc_bf16.h" |
Examples
1 2 3 4 5 6 | // In hybrid programming scenarios, the address space qualifier __gm__ is used to indicate the Global Memory space. __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); } |