CUDA新特性：原子操作与动态并行

原子操作

由于原子操作由GPU内存控制器实现，它们只在本地设备内存位置工作。

atomicAdd()
//读取位于全局或共享存储器中地址address处的32位或64位字old，计算(old+val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。只有全局存储器支持64位字。
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);

atomicSub()
//读取位于全局或共享存储器中地址address 处的32位字old，计算(old-val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address,unsigned int val);

atomicExch()
//读取位于全局或共享存储器中地址address处的32位或64位字old，并将val 存储在存储器的同一地址中。这两项操作在一次原子事务中执行。该函数将返回old。只有全局存储器支持64位字。
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处的32位字old，计算old 和val的最大值，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
atomicMin()
int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address,unsigned int val);
atomicMax()
int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address,unsigned int val);

atomicInc()
//读取位于全局或共享存储器中地址address处的32位字old，计算 ((old >= val) ? 0:(old+1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
unsigned int atomicInc(unsigned int* address,unsigned int val);
atomicDec()
//读取位于全局或共享存储器中地址address处的32位字old，计算( ((old == 0)|(old > val)) ? val : (old-1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
unsigned int atomicDec(unsigned int* address,unsigned int val);

//位逻辑函数
atomicAnd()
//读取位于全局或共享存储器中地址address 处的32位字old，计算 (old & val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,unsigned int val);
atomicOr()
//读取位于全局或共享存储器中地址address 处的32 位字old，计算 (old | val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,unsigned int val);
atomicXor()
//读取位于全局或共享存储器中地址address 处的32 位字old，计算 (old ^ val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,unsigned int val);

除了自主 原子操作可以使用在线程块之间的同步。 CUDA硬件支持同步的主要抽象：“对比和交换”(或CAS，compare and swap)。在CUDA上，对比和交换(也可为compare and exchange–即x86中的CMPXCHG指令)被定义如下：

//读取位于全局或共享存储器中地址address处的32位或64位字old，计算(old==expect?val: old)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old（比较并交换）。只有全局存储器支持64 位字。
int atomicCAS(int *address, int expect, int val);
unsigned int atomicCAS(unsigned int* address, unsigned int expect, unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address, unsigned long long int expect, unsigned long long int val);

用atomicCAS实现自旋锁

 class cudaSpinlock{
 public:
	cudaspinlock(int *p); 
 	void acquire();
 	void release();
 private:
 	int *m_p;
 };
 inline __device__ cudaSpinlock::cudaspinlock(int *p)
 {
 	m_p=p;
 }

 inline __device__ void cudaSpinlock::acquire()
 {
 	while( atomicCAS(m_p,0,1));
 }

 inline __device__ void cudaSpinlock::release()
 {
 	atomicExch(m_p,0);
 }
 

动态并行

动态并行是计算能力为3.5及以上的设备支持的一个新特性。该特性支持从设备端启动 kernel ,从而超越了简单的递归方式。这就意味着分而治之策略可以方便地用 CUDA 实现。一个 kernel 可以使用同主机端一样的语法(即,使用<<<>>>操作符),实现其他 kernel 的异步开启。
一个 CUDA 线程启动的线程网格是该线程网络的子线程。该线程网格的启动者被称为父线程网格。父线程网格开启的子线程网格异步执行,就如同在主机端启动一样。同时,满足典型的嵌套规则,即,直到所有启动的子线程网格都执行完毕后,该线程网格才执行完毕。

设备启动的 kernel 只能被同启动线程属于同一个线程块的线程监管。通过主机端使用的 API 的子集,可以使用流和事件进行同步以及控制不同 kernel 间的依赖。然而,在多 GPU 系统中,设备端启动的 kernel 不能运行在另一个 GPU 上,这仍然是主机的特权。
至于内存管理,子线程网络可以传递对全局内存的引用,但不同传递对共享内存的引用。这是因为共享内存是 SM 私有的。相同的规则也适用于私有内存,因为私有内存是线程私有的。