asc_ldcg
Applicability
Product |
Supported |
|---|---|
Atlas 350 Accelerator Card |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
Function Usage
Loads cached data from the L2 cache. If the cache hit occurs, the data is returned directly. If the cache miss occurs, the data is preloaded from Global Memory into L2 Cache and then returned.
Prototype
1 | inline long int asc_ldcg(long int* address) |
1 | inline unsigned long int asc_ldcg(unsigned long int* address) |
1 | inline long long int asc_ldcg(long long int* address) |
1 | inline unsigned long long int asc_ldcg(unsigned long long int* address) |
1 | inline long2 asc_ldcg(long2* address) |
1 | inline ulong2 asc_ldcg(ulong2* address) |
1 | inline long4 asc_ldcg(long4* address) |
1 | inline ulong4 asc_ldcg(ulong4* address) |
1 | inline longlong2 asc_ldcg(longlong2* address) |
1 | inline ulonglong2 asc_ldcg(ulonglong2* address) |
1 | inline longlong4 asc_ldcg(longlong4* address) |
1 | inline ulonglong4 asc_ldcg(ulonglong4* address) |
1 | inline signed char asc_ldcg(signed char* address) |
1 | inline unsigned char asc_ldcg(unsigned char* address) |
1 | inline char2 asc_ldcg(char2* address) |
1 | inline uchar2 asc_ldcg(uchar2* address) |
1 | inline char4 asc_ldcg(char4* address) |
1 | inline uchar4 asc_ldcg(uchar4* address) |
1 | inline short asc_ldcg(short* address) |
1 | inline unsigned short asc_ldcg(unsigned short* address) |
1 | inline short2 asc_ldcg(short2* address) |
1 | inline ushort2 asc_ldcg(ushort2* address) |
1 | inline short4 asc_ldcg(short4* address) |
1 | inline ushort4 asc_ldcg(ushort4* address) |
1 | inline int asc_ldcg(int* address) |
1 | inline unsigned int asc_ldcg(unsigned int* address) |
1 | inline int2 asc_ldcg(int2* address) |
1 | inline uint2 asc_ldcg(uint2* address) |
1 | inline int4 asc_ldcg(int4* address) |
1 | inline uint4 asc_ldcg(uint4* address) |
1 | inline float asc_ldcg(float* address) |
1 | inline float2 asc_ldcg(float2* address) |
1 | inline float4 asc_ldcg(float4* address) |
1 | inline bfloat16_t asc_ldcg(bfloat16_t* address) |
1 | inline bfloat16x2_t asc_ldcg(bfloat16x2_t* address) |
1 | inline half asc_ldcg(half* address) |
1 | inline half2 asc_ldcg(half2* address) |
Parameters
Parameter |
Input/Output |
Description |
|---|---|---|
address |
Input |
Address of the Global Memory |
Returns
Cached data in 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_ldcg(__gm__ float* dst, __gm__ float* src) { int idx = threadIdx.x + blockIdx.x * blockDim.x; dst[idx] = asc_ldcg(src + idx); } |