CUDA C编程手册: 编程模型

核函数

CUDA C 是对C的一中扩展， 允许程序员定义自己的C函数， 并且将之称为 核。与传统C函数不同的是， 这样的核被调用的时候， 会被CUDA 线程并行地执行N次。

核函数定义的时候使用 _global_声明标识符。同时， 在调用这样的核函数的时候， 会使用一个特定的语法
<<< … >>>来指定执行的配置。每一个线程执行这个核函数的时候， 都会被分配一个独一无二的线程ID， 这个线程ID可以通过内置的变量 threadIdx来获得。

下面是一个示例代码，演示两个大小为N的向量A和B相加，结果存放在C中。 其中每个线程都会执行VecAdd()核函数， 使得每一对元素相加。

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
	int i = threadIdx.x;
	C[i] = A[i] + B[i];
}
int main()
{
	...
	// Kernel invocation with N threads
	VecAdd<<<1, N>>>(A, B, C);
	...
}

线程体系

为了简单起见， threadIdx是一个三元素的向量， 因此可以通过一个一维/二维/三维的线程索引来组成一个一维/二维/三维的线程块(thread block)。对于向量、矩阵和体数据的计算来说， 这是一种非常自然且合适的方法。

线程的索引和其ID之间的关系是非常直观的：对于一维的线程块， 他们是一样的， 即索引就是线程的ID； 对于大小为($D_x$, $D_y）$的线程块来说， 线程索引为$(x,y)$的ID是$(x+y\cdot D_x)$; 对于一个三维$(D_x,D_y,D_z)$的线程块来说， 线程的索引为$(x,y,z)$的ID是$(x+y\cdot D_x+z\cdot D_x\cdot D_z)$。

下面简单的示例代码，表示的是两个$N\times N$的矩阵A和B相加，结果存储在C中。

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
	int i = threadIdx.x;
	int j = threadIdx.y;
	C[i][j] = A[i][j] + B[i][j];
}
int main()
{
	...
	// Kernel invocation with one block of N * N * 1 threads
	int numBlocks = 1;
	dim3 threadsPerBlock(N, N);
	MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
	...
}

每个线程块中的线程数量都有上限，因为每个线程块中的线程都会驻留在同一个处理核心中且同时共享有限的处理核心上的内存资源。对于现在发行的GPU， 一个线程块中线程的上限是1024。然而， 一个核函数， 一个核函数可以被多个线程块执行， 所以线程的总数量为每个线程块中线程数量与线程块个数的乘积。

线程块可以被组成一维的， 二维的或者三维的线程块网格grid。线程块网格中线程块的数量通常根据需要处理数据大小而得来， 或者是系统的处理器的数量， 这样能获得较大的加速效果。

线程块的执行是无序的， 即他们的执行时相互独立的， 没有依赖关系。同一个线程块内的线程可以通过共享内存shared memory来进行合作， 或者通过线程同步来“有序”的访问内存。更确切地说， 可以通过在核函数中调用*__syncthreads()*来指明的同步点， 这个函数类似于一个栅栏， 只有当块中的所有线程都到达了该点才可以继续往下执行，否则先到达该点的线程都需要停下来等待其他线程。

内存体系

CUDA中的线程在执行的过程中， 可以访问多种内存空间中的数据。每个线程都有自己的私有的本地内存。 每个线程块都有对所有线程块可见的共享内存且该内存与线程块具有同样的生命周期；所有的线程都可以访问同样的全局内存空间。

存在两种可以被所有线程访问的额外的只读内存空间：常量内存空间和纹理内存空间。全局、常量和纹理内存空间都存在一些优化的使用方法。纹理内存也提供了一些特殊的寻址模式和对一些特殊数据格式的数据滤波。全局、常量和纹理内存空间贯穿于整个应用。

异构编程

CUDA编程模型假定在物理分离的设备device上执行的CUDA线程作为在主机hostC程序的协助者。也就是说，核函数运行在GPU上而剩下的C代码运行在CPU上。

CUDA编程模型也同时假定了主机和设备同时在DRAM上分开持有他们自己的内存空间， 分别称作host memory和device memeory。因此，一个程序同构调用CUDA 运行时（runtime）来管理着对所有线程可见的全局、常量和纹理内存空间。这些包括对设备内存的分配和释放以及主机和设备之间的数据传输。

统一内存通过managed memory来连接主机和设备两个内存空间。managed memory可以通过一个使用共同的寻址空间来被系统中的所有GPUs和CPUs所访问。这就使得设备内存使用可以“过载”且可以通过取消数据的在设备端和主机端进行镜像以极大地简化应用的移植难度。

#计算力
设备的计算力compute capability一般是由一个版本号表示，有时候也称作SM version。 这个版本号揭示了GPU硬件所支持的特性和决定了应用程序在当前GPU运行时哪些特性或者指令是可用的。

计算力的版本号通过一个主版本X和一个小版本Y构成， 即 X.Y。拥有相同主版本的设备具有通用的核心架构， 其中7表示Volta架构， 6表示Pascal架构， 5表示Maxwell架构， 3表示Kepler架构， 2表示Fermi架构， 1表示Tesla架构。小版本表示在主架构上进行了一些小提升，或者是一些新特性。