CUDA C++ Programming Guide

【CUDA编程】【28】【6.C++ Language Extensions】【Part3】

Load Functions Using Cache Hints，Store Functions Using Cache Hints，Time Function，Atomic Functions，Address Space Predicate Functions

Posted by x-jeff on January 17, 2025

【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。

1.Load Functions Using Cache Hints

这些加载函数仅支持计算能力在5.0及以上的device。

1
2
3
4
5

T __ldcg(const T* address);
T __ldca(const T* address);
T __ldcs(const T* address);
T __ldlu(const T* address);
T __ldcv(const T* address);

以下是个人对这些函数的解释，如有不准确的地方，欢迎批评指正：

  • __ldcg()：仅缓存到全局级别（L2缓存，不包括L1缓存）。
  • __ldca()：缓存到所有级别（L1和L2缓存）。这是默认的加载模式，适用于频繁访问的数据。
  • __ldcs()：流式缓存，适用于只访问一次或两次的数据（流式数据处理）。
  • __ldlu()：最后一次使用，数据不会再次使用，因此避免写回或缓存污染。
  • __ldcv()：不缓存，每次加载都重新获取数据（无缓存加载）。

解释参考下表，Cache Operators for Memory Load Instructions：

类型T可以是：

  • char、signed char、short、int、long、long long、unsigned char、unsigned short、unsigned int、unsigned long、unsigned long long。
  • char2、char4、short2、short4、int2、int4、longlong2、uchar2、uchar4、ushort2、ushort4、uint2、uint4、ulonglong2。
  • float、float2、float4、double、double2。
  • 如果包含了cuda_fp16.h头文件，类型T还可以是__half或__half2。
  • 如果包含了cuda_bf16.h头文件，类型T还可以是__nv_bfloat16或__nv_bfloat162。

2.Store Functions Using Cache Hints

这些存储函数仅支持计算能力在5.0及以上的device。

1
2
3
4

void __stwb(T* address, T value);
void __stcg(T* address, T value);
void __stcs(T* address, T value);
void __stwt(T* address, T value);

个人解释（这些函数可以和第1部分的加载函数配合使用，需要注意缓存一致性以提升性能）：

  • __stwb()：缓存写回所有一致的缓存级别（L1和L2缓存），是默认的存储模式。
  • __stcg()：仅缓存到全局级别（L2缓存，不包括L1缓存）。
  • __stcs()：流式缓存存储，适用于仅访问一次或两次的流式数据。
  • __stwt()：写穿（write-through）存储，数据直接写到系统内存（通过L2缓存写入全局内存）。

参考下表，Cache Operators for Memory Store Instructions：

类型T可以是：

  • char、signed char、short、int、long、long long、unsigned char、unsigned short、unsigned int、unsigned long、unsigned long long。
  • char2、char4、short2、short4、int2、int4、longlong2、uchar2、uchar4、ushort2、ushort4、uint2、uint4、ulonglong2。
  • float、float2、float4、double、double2。
  • 如果包含了cuda_fp16.h头文件，类型T还可以是__half或__half2。
  • 如果包含了cuda_bf16.h头文件，类型T还可以是__nv_bfloat16或__nv_bfloat162。

3.Time Function

1
2
3

//记录的为时钟周期数
clock_t clock(); //返回的是一个32位整数，适合测量较短时间间隔
long long int clock64(); //返回的是一个64位整数，适合测量更长的时间间隔，避免32位计数器溢出

clock()和clock64()通常只能在同一个线程内使用。这是因为它们返回的计数器值是线程所在的多处理器上的局部计数器的值，而不是整个device上的全局计数器。

4.Atomic Functions

一个原子函数（atomic function）会对全局内存或共享内存中的一个字（word，可以是32位、64位或128位）执行“读-修改-写”的原子操作。在处理float2或float4类型时，会对驻留在全局内存中的向量中的每个元素都执行“读-修改-写”操作。例如，atomicAdd()函数会在全局内存或共享内存中的某个地址读取一个值，将一个数加到该值上，然后将结果写回相同地址。整个过程是原子的，不会被其他线程中断。原子函数只能在device函数中使用。

本部分描述的原子函数遵循cuda::memory_order_relaxed的内存顺序（意味着操作没有任何额外的同步或顺序约束，仅保证该原子操作本身是原子的），并且只在特定范围内是原子的：

  • 带有_system后缀的原子API（比如atomicAdd_system），如果满足特定条件，则它们在范围cuda::thread_scope_system内是原子的。
  • 不带后缀的原子API（比如atomicAdd），在范围cuda::thread_scope_device内是原子的。
  • 带有_block后缀的原子API（比如atomicAdd_block），在范围cuda::thread_scope_block内是原子的。

以下示例展示了CPU和GPU如何在地址addr上对一个整数值进行原子更新操作：

1
2
3
4
5
6
7
8
9
10
11
12

__global__ void mykernel(int *addr) {
  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
}

void foo() {
  int *addr;
  cudaMallocManaged(&addr, 4); //分配托管内存，即CPU和GPU可以共享的内存
  *addr = 0;

   mykernel<<<...>>>(addr);
   __sync_fetch_and_add(addr, 10);  // CPU atomic operation
}

请注意，任何原子操作都可以基于atomicCAS()（Compare And Swap，比较与交换）实现。例如，针对双精度浮点数的atomicAdd()在计算能力低于6.0的device上不可用，但可以通过以下方式实现：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        //atomicCAS()解释见第4.1.8部分
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));

    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif
  • 对于计算能力低于6.0的device，仅支持device范围的原子操作，不支持系统范围和块范围的原子操作。
  • 对于计算能力低于7.2的Tegra device，不支持系统范围的原子操作。

4.1.Arithmetic Functions

4.1.1.atomicAdd()

1
2
3
4
5
6
7
8
9
10
11
12
13

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
                                 unsigned long long int val);
float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);
__half2 atomicAdd(__half2 *address, __half2 val);
__half atomicAdd(__half *address, __half val);
__nv_bfloat162 atomicAdd(__nv_bfloat162 *address, __nv_bfloat162 val);
__nv_bfloat16 atomicAdd(__nv_bfloat16 *address, __nv_bfloat16 val);
float2 atomicAdd(float2* address, float2 val);
float4 atomicAdd(float4* address, float4 val);

从全局内存或共享内存的指定地址address中读取old值（可以是16位、32位或64位），然后计算(old + val)，并将结果写回同一地址。这三个操作（读、加、写）是以原子事务的方式完成的，确保线程安全。函数返回old。

举个例子：

1
2
3
4
5

int a = 5;
int b = 1;
int* pa = &a;
int s;
s = atomicAdd(pa, b); //s等于5，*pa等于6

32位浮点版本的atomicAdd()仅支持计算能力在2.x及以上的device。

64位浮点版本的atomicAdd()仅支持计算能力在6.x及以上的device。

32位__half2浮点版本的atomicAdd()仅支持计算能力在6.x及以上的device。对于__half2 atomicAdd(__half2 *address, __half2 val);来说，对__half2的整体操作并不能保证是原子的，而对其包含的两个__half元素的操作是被保证为原子的。简单解释下，__half2包含两个元素：__half[0]和__half[1]，硬件可以保证对这两个元素的操作分别是原子的，但多个线程可以同时尝试对__half2中的不同元素进行更新，此时更新的最终结果可能会因为竞争条件而不一致。这种情况也适用于__nv_bfloat16、float2、float4等复合类型。

float2和float4浮点向量版本的atomicAdd()仅支持计算能力在9.x及以上的device，并且仅支持全局内存地址。

16位__half浮点版本的atomicAdd()仅支持计算能力在7.x及以上的device。

16位__nv_bfloat16浮点版本的atomicAdd()仅支持计算能力在8.x及以上的device。

4.1.2.atomicSub()

1
2
3

int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address,
                       unsigned int val);

从全局内存或共享内存的指定地址address中读取old值（32位），然后计算(old - val)，并将结果写回同一地址。这三个操作（读、减、写）是以原子事务的方式完成的，确保线程安全。函数返回old。

4.1.3.atomicExch()

1
2
3
4
5
6

int atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address,
                        unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address,
                                  unsigned long long int val);
float atomicExch(float* address, float val);

从全局内存或共享内存的指定地址address中读取old值（32位或64位），然后将val值写回到同一地址。这两个操作（读、写）在一次原子事务中完成。函数返回old。

1

template<typename T> T atomicExch(T* address, T val);

从全局内存或共享内存的指定地址address中读取old值（128位），然后将val值写回到同一地址。这两个操作（读、写）在一次原子事务中完成。函数返回old。类型T必须满足以下要求：

1
2
3
4
5

sizeof(T) == 16
alignof(T) >= 16
std::is_trivially_copyable<T>::value == true
// for C++03 and older
std::is_default_constructible<T>::value == true

128位版本的atomicExch()仅支持计算能力在9.x及以上的device。

4.1.4.atomicMin()

1
2
3
4
5
6
7

int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicMin(unsigned long long int* address,
                                 unsigned long long int val);
long long int atomicMin(long long int* address,
                                long long int val);

从全局内存或共享内存的指定地址address中读取old值（32位或64位），然后计算old和val的最小值，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

64位版本的atomicMin()仅支持计算能力在5.0及以上的device。

4.1.5.atomicMax()

1
2
3
4
5
6
7

int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicMax(unsigned long long int* address,
                                 unsigned long long int val);
long long int atomicMax(long long int* address,
                                 long long int val);

从全局内存或共享内存的指定地址address中读取old值（32位或64位），然后计算old和val的最大值，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

64位版本的atomicMax()仅支持计算能力在5.0及以上的device。

4.1.6.atomicInc()

1
2

unsigned int atomicInc(unsigned int* address,
                       unsigned int val);

从全局内存或共享内存的指定地址address中读取old值（32位），然后计算((old >= val) ? 0 : (old+1))，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

4.1.7.atomicDec()

1
2

unsigned int atomicDec(unsigned int* address,
                       unsigned int val);

从全局内存或共享内存的指定地址address中读取old值（32位），然后计算(((old == 0) || (old > val)) ? val : (old-1)，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

4.1.8.atomicCAS()

1
2
3
4
5
6
7
8
9
10

int atomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address,
                       unsigned int compare,
                       unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address,
                                 unsigned long long int compare,
                                 unsigned long long int val);
unsigned short int atomicCAS(unsigned short int *address,
                             unsigned short int compare,
                             unsigned short int val);

从全局内存或共享内存的指定地址address中读取old值（16位、32位或64位），然后计算(old == compare ? val : old)，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。CAS的意思是Compare And Swap。

1

template<typename T> T atomicCAS(T* address, T compare, T val);

从全局内存或共享内存的指定地址address中读取old值（128位），然后计算(old == compare ? val : old)，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

类型T必须满足以下要求：

1
2
3
4
5

sizeof(T) == 16
alignof(T) >= 16
std::is_trivially_copyable<T>::value == true
// for C++03 and older
std::is_default_constructible<T>::value == true

128位版本的atomicCAS()仅支持计算能力在9.x及以上的device。

4.2.Bitwise Functions

4.2.1.atomicAnd()

1
2
3
4
5

int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAnd(unsigned long long int* address,
                                 unsigned long long int val);

从全局内存或共享内存的指定地址address中读取old值（32位或64位），然后计算(old & val)，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

64位版本的atomicAnd()仅支持计算能力在5.0及以上的device。

4.2.2.atomicOr()

1
2
3
4
5

int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,
                      unsigned int val);
unsigned long long int atomicOr(unsigned long long int* address,
                                unsigned long long int val);

从全局内存或共享内存的指定地址address中读取old值（32位或64位），然后计算(old | val)，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

64位版本的atomicOr()仅支持计算能力在5.0及以上的device。

4.2.3.atomicXor()

1
2
3
4
5

int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicXor(unsigned long long int* address,
                                 unsigned long long int val);

从全局内存或共享内存的指定地址address中读取old值（32位或64位），然后计算(old ^ val)，并将其写回到同一地址。这三个操作在一次原子事务中完成。函数返回old。

64位版本的atomicXor()仅支持计算能力在5.0及以上的device。

5.Address Space Predicate Functions

本部分的函数如果参数为null指针，则其行为未定义。

5.1.__isGlobal()

1

__device__ unsigned int __isGlobal(const void *ptr);

如果ptr是全局内存空间中对象的通用地址，则返回1，否则返回0。

5.2.__isShared()

1

__device__ unsigned int __isShared(const void *ptr);

如果ptr是共享内存空间中对象的通用地址，则返回1，否则返回0。

5.3.__isConstant()

1

__device__ unsigned int __isConstant(const void *ptr);

如果ptr是常量内存空间中对象的通用地址，则返回1，否则返回0。

5.4.__isGridConstant()

1

__device__ unsigned int __isGridConstant(const void *ptr);

如果ptr是带有__grid_constant__的kernel参数的通用地址，则返回1，否则返回0。仅支持计算能力在7.x及以上的device。

5.5.__isLocal()

1

__device__ unsigned int __isLocal(const void *ptr);

如果ptr是局部内存空间中对象的通用地址，则返回1，否则返回0。

CATALOG
    FEATURED TAGS
    README Machine Learning Series Mathematics Series Image Segmentation Python Series Fundamentals of Computers OpenCV Series C++ Series Deep Learning Series Conda Tensorflow Series Algorithms Series AI Papers Object Recognition Natural Language Processing Aha! Algorithms Camera Calibration How Program Works Optical Flow Estimation How Networks Work Computer Vision Source Code CUDA C++ Programming Guide C++ Concurrency IN ACTION Large Language Models