CUDA变量存储与原子操作

C++的变量存储与赋值

变量类型

1. 局部变量
   • 函数中定义的变量是该函数的局部变量，他们在其他函数的语句中是不可见的
   • 函数的形参也是局部变量，其作用域仅限于函数体
   • 只有在定义的函数正在执行时，局部变量才存在
2. 全局变量
   • 在所有函数之外的所有变量是全局变量，全局变量可以被所有定义在全局变量之后的函数访问
   • 只需在一个源文件中定义，就可以作用于所有的源文件
   • 其他不包含全局变量定义的源文件需要用extern 关键字再次声明这个全局变量
3. 静态局部变量
   • 定义在函数内部，只能在函数内部访问，并且其生命周期始终与程序相关
   • 当函数被调用时，静态局部变量的值会被保留，即使函数已经退出
4. 静态全局变量
   • 只会被初始化一次，并在整个程序生命周期中保持不变
   • 值在每次程序运行时都是相同的，而全局变量的值在每次程序运行时都是不同的
   • 如果程序包含多个文件的话，作用于定义它的文件里，不能作用到其它文件里

变量存储

C语言经过编译后将内存分为以下几个区域：

1. stack 栈：
   • 由编译器进行管理，自动分配和释放
   • 存放函数调用过程中的各种参数、局部变量、返回值以及函数返回地址
2. heap 堆：
   • 用于程序动态申请分配和释放空间，若没有释放则程序结束时系统自动回收
   • malloc和free、new和delete
3. 全局（静态）存储区：
   • DATA段（全局初始化区）：存放初始化的全局变量和静态变量
   • BSS段（全局未初始化区）：存放未初始化的全局变量和静态变量
4. 文字常量区：存放常量字符，程序结束后由系统释放
5. 程序代码区：存放程序的二进制代码

变量赋值

赋值的过程其实就是复制对象的值，并将该值赋给另一个对象。

CUDA 内存模型

CUDA变量类型

1. __device__
2. __constant__
3. __shared__
4. __managed__
5. __restrict__

GPU 内存设备

• CUDA 中每个线程都有自己私有的本地内存
• 线程块有自己的共享内存，对线程块内所有的线程可见
• 所有的线程都能访问读取常量内存与纹理内存（只读不写）
• 对于一个应用而言，全局内存、常量内存与纹理内存具有相同的生命周期

• 寄存器

  • 寄存器对于每个线程是私有的

  • 需要考虑由于线程中的变量过多以至于寄存器发生溢出的情况

    // 方法一：在核函数中配置额外信息辅助编译器优化
    __global__ void
    __lauch_bounds__(maxThreadaPerBlock,minBlocksPerMultiprocessor)
    kernel(...) {
        /* kernel code */
    }
    

    # 方法二：在编译选项中控制编译单元中所有核函数使用的最大数量
    -maxrregcount=32
    

• 共享内存

  • 在核函数中使用__share__修饰
  • 共享内存是片上内存，跟主存相比，速度要快很多，也即是延迟低，带宽高。
  • 类似于一级缓存，但是可以被编程

• 本地内存

  • 核函数中符合存储在寄存器中但不能进入被核函数分配的寄存器空间中的变量将存储在本地内存中
  • 编译器可能存放在本地内存中的变量类型
    1. 使用未知索引引用的本地数组
    2. 可能会占用大量寄存器空间的较大本地数组或者结构体
    3. 任何不满足核函数寄存器限定条件的变量
  • 高延迟，低带宽

• 常量内存

  • 在核函数外使用__constant__修饰
  • host 端可以初始化常量内存，被 host 端初始化后不能被核函数修改

• 纹理内存

  • 通过指定的缓存访问的全局内存

• 全局内存

  • GPU 上最大的内存空间，延迟最高
  • 一般在主机端代码里定义，也可以在设备端定义（需要加修饰符）
  • 不销毁则与应用程序同生命周期

GPU 线程与存储

• 所有通过cudaMalloc分配的存储器都是全局内存

CUDA原子操作

atomic 函数可以对全局/共享内存的变量执行读-修改-写的原子操作

要点

• 当线程束的多个线程向同一个内存地址写数据时，程序的正确性是无法保证的，此时需要使用 CUDA 提供的 atomic 函数
• 原子操作可以保证程序的正确性，但是会造成线程束中线程的串行化 serialization ，执行时间比并行执行要长。这是因为在一个线程完成该操作之前，其他线程将不能对相同的内存地址进行读写操作
• 通过使用共享内存，可以减少对原子操作的使用，从而提高程序的性能

atomic 函数

• atomicAdd()

  • 操作：address是old的在全局/共享内存的地址，进行(old + val)的操作，并且将答案赋值回address对应的地址下
  • 返回：old

  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);
  

• atomicSub()

  • 操作：*address = (old - val)
  • 返回：old

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

• atomicExch()

  • 操作：*address = val
  • 返回：old

  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);
  

• atomicMin()

  • 操作：*address = min(old, val)
  • 返回：old

  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);
  

• atomicMax()

  • 操作：*address = max(old, val)
  • 返回：old

  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);
  

• atomicInc()

  • 操作：*address = ((old >= val) ? 0 : (old + 1))
  • 返回：old

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

• atomicDec()

  • 操作：*address = (((old == 0) || (old > val)) ? val : (old - 1))
  • 返回：old

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

• atomicCAS()

  • 操作：*address = (old == compare ? val : old)
  • 返回：old( Compare And Swap)

  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);
  

• atomicAnd()

  • 操作：*address = (old & val)
  • 返回：old

  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);
  

• atomicOr()

  • 操作：*address = (old | val)
  • 返回：old

  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);
  

• atomicXor()

  • 操作：*address =(old ^ val)
  • 返回：old

  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);