Kernels（内核）
CUDA C ++通过允许程序员定义称为kernels（内核）的C ++函数来扩展C ++，这些函数在被调用时由N个不同的CUDA线程并行执行N次，而不是像常规C ++函数那样仅执行一次。
CUDA使用global声明定义kenels，并且当执行给定内核调用的内核时，需要使用新的<<< ... >>>执行配置语法来指定CUDA线程数。 每个执行内核的线程都有一个唯一的线程ID，该ID在内核中可以通过内置变量来访问。
作为示例，以下代码使用内置变量threadIdx将两个大小为N的向量A和B相加，并将结果存储到向量C中：
// 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);
...
在这里，执行VecAdd（）的N个线程中的每一个都执行一对成对加法。
线程层次结构
为方便起见，threadIdx是3分量向量，因此可以使用一维，二维或三维线程索引来标识线程，形成一个一维，二维或三维块, 称为线程块。 这提供了一种自然的方式来调用域中元素（例如矢量，矩阵或体积）进行计算。
线程的索引及其线程ID以直接的方式相互关联：
对于一维块，它们是相同的。 对于大小为（Dx，Dy）的二维块，索引为（x，y）的线程的线程ID为（x + y Dx）； 对于大小为（Dx，Dy，Dz）的三维块，索引为（x，y，z）的线程的线程ID为（x + y Dx + z Dx Dy）。
例如，以下代码将两个大小为NxN的矩阵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个线程。
但是，内核可以由多个大小相同的线程块执行，因此线程的总数等于每个块的线程数乘以块数。
如图所示，将块组织为一维，二维或三维的线程块网格。网格中的线程块数通常由所处理数据的大小决定，通常超过 系统中的处理器数量。
每个块的线程数和每个网格的块数在<<< ... >>>语法中指定, 可以是int或dim3(三维向量)类型。 二维块或网格可以像上面的示例中那样指定。
网格中的每一个块都可以被标识为一维、二维或三维的唯一索引，此索引可以通过内核内置的blockIdx变量来访问。线程块的维度可以在内核内通过内置的blockDim变量来访问。
以下代码示例通过扩展前面的MatAdd（）示例来处理多个块：
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
常见的线程块大小为16x16(256个线程)。参考线程块网格图，以上代码同样确保网格中有足够的线程块，使得每个快中的矩阵元素只包含一个线程。为简单起见，此示例假定每个维度中每个网格的线程数可以被该维度中每个块的线程数平均除尽，当然实际情况并非必须如此。
线程块需要能够以任意顺序，并行或串行的独立执行。这种独立性要求允许线程块在任意数量的内核之间以任意顺序进行调度，如下图所示，从而使程序员可以编写随内核数量扩展的代码。
一个线程块中的线程间要进行协作，需要通过共享内存来共享数据并且同步他们的执行来协调内存访问。更准确地说，可以通过调用__syncthreads（）内在函数来指定内核中的同步点。 __syncthreads（）充当屏障，在该屏障中，块中的所有线程必须等待，然后才能继续执行任何线程。 共享内存给出了使用共享内存的示例。 除了__syncthreads（）之外，合作组API还提供了一组丰富的线程同步原语。
为了进行有效的协作，共享内存应该是每个处理器核心附近的低延迟内存（非常类似于L1缓存），而__syncthreads（）应该是轻量级的。
内存层次结构
CUDA线程在执行期间可能会从多个内存空间访问数据，如下图所示。每个线程都有专用的本地内存。 每个线程块具有对该块的所有线程可见的共享内存，并且具有与该块相同的生存期。所有线程都可以访问相同的全局内存。
所有线程还可以访问两个附加的只读存储空间：常量存储空间和纹理存储空间。 全局，常量和纹理内存空间针对不同的内存使用进行了优化。 纹理存储器还为某些特定的数据格式提供了不同的寻址模式以及数据过滤。
全局，常量和纹理存储空间在同一应用程序的内核启动期间是持久的。
异构编程
如下图所示，CUDA编程模型假定CUDA线程在物理上独立的设备上执行，该设备充当运行C ++程序所在主机的协处理器。 例如，当内核在GPU上执行而其余C ++程序在CPU上执行时，就是这种情况。
CUDA编程模型还假定主机和设备都在DRAM中维护自己的独立存储空间，分别称为主机存储器和设备存储器。 因此，程序通过调用CUDA运行时间（CUDA runtime）来管理内核可见的全局、常量和纹理存储空间。其中包括设备内存的分配和释放以及主机与设备内存之间的数据传输。统一内存提供托管内存，以桥接主机和设备内存空间。 系统中的所有CPU和GPU都可以将托管内存作为具有公共地址空间的单个一致内存映像进行访问。 此功能可消除设备内存的超额分配，并且无需在主机和设备上显式镜像数据，从而可以大大简化移植应用程序的工作量。设备的计算能力由版本号表示，有时也称为“ SM版本”。 此版本号标识GPU硬件支持的功能，并且当应用程序在运行时可以用于确定当前GPU上可用的硬件功能和指令。
