一、CUDA编程模型基础
在进入CUDA C代码之前，CUDA的那些新知识将受益于CUDA编程模型的基本描述以及所使用的一些术语。
CUDA编程模型是一种异构模型，其中同时使用了CPU和GPU。在CUDA中，主机是指CPU及其内存，而设备是指GPU及其内存。在主机上运行的代码可以管理主机和设备上的内存，还可以启动内核，这些内核是在设备上执行的功能。这些内核由许多GPU线程并行执行。
鉴于CUDA编程模型的异构性质，CUDA C程序的典型操作序列为：
（1）、声明并分配主机和设备内存。
（2）、初始化主机数据。
（3）、将数据从主机传输到设备。
（4）、执行一个或多个内核。
（5）、将结果从设备传输到主机。
牢记此操作顺序，让我们看一下CUDA C示例。
1、第一个CUDA C程序
函数 saxpy 是在GPU上并行运行的内核， main 函数是主机代码。让我们从主机代码开始对程序分析。
2、主机代码
main函数声明了两对数组
float *x, *y, *d_x, *d_y;
x = (float*)malloc(N*sizeof(float));
y = (float*)malloc(N*sizeof(float));
cudaMalloc(&d_x, N*sizeof(float));
cudaMalloc(&d_y, N*sizeof(float));
指针 x 和 y 通过malloc函数分配指定大小的内存。
指针 d_x 和 d_y 通过CUDA分配内存的函数cudaMalloc来分配GPU上的资源。
CUDA中的主机和设备具有单独的内存空间，这两个内存空间都可以通过主机代码进行管理（CUDA C内核也可以在支持该内存的设备上分配设备内存）。
主机上的变量做初始化操作，我们设置数组 x 的所有元素为1，并设置数组 y 的所有元素为2。
for (int i = 0; i < N; i ) {
x[i] = 1.0f;
y[i] = 2.0f;
}
初始化设备（GPU）数据，通过cuda函数cudaMemcpy将CPU上的数据拷贝到GPU上，即将主机上的数据拷贝到设备上。
cudaMemcpy函数的总共需要四个参数。
第一个参数：目标数据指针
第二个参数：源数据指针
第三个参数：需要拷贝数据的大小
第四个参数：拷贝的方向
（1）从主机拷贝到设备cudaMemcpyHostToDevice；
（2）从设备拷贝到主机cudaMemcpyDeviceToHost。
cudaMemcpy （d_x ，x ，N * sizeof （float ），cudaMemcpyHostToDevice ）;
cudaMemcpy （d_y ，y ，N * sizeof （float ），cudaMemcpyHostToDevice ）;
运行内核之后，我们将设备中计算的结果拷贝到主机上。
cudaMemcpy（y，d_y，N * sizeof（float），cudaMemcpyDeviceToHost）;
3、启动内核
saxpy 内核由以下语句启动：
saxpy<<<(N 255)/256, 256>>>(N, 2.0, d_x, d_y);
三个尖括号之间的信息是对GPU进行配置，该配置决定了有多少个设备线程并行执行该内核。在CUDA中，软件中存在线程层次结构，该层次结构模仿了如何在GPU上对线程处理器进行分组。在CUDA编程模型中，我们谈到启动带有线程块网格的内核 。执行配置中的第一个参数指定网格中线程块的数量，第二个参数指定线程块中线程的数量。
通过为这些参数传递dim3（由CUDA定义的简单结构，带有x和y，以及z成员）值，可以将线程块和网格制成一维，二维或三维，但是对于这个简单的示例，我们只需要一个维，因此我们可以传递整数。在这种情况下，我们使用包含256个线程的线程块启动内核，并使用整数算法确定处理N数组的所有元素（(N 255)/256）所需的线程块数 。
对于无法通过线程块大小将数组中的元素数整除的情况，内核代码必须检查越界内存访问。即设置每个线程块中存在256个线程，但是我们需要确定需要有多少个线程块。我们有N的数据需要并行处理，那么我们通过（N 255）/ 256来确定多少个线程块，对于多出来的线程（即最后一个块可能不满）我们必须做越界处理。
4、释放分配的所有资源
程序运行结束之后，需要释放我们之前分配的内存空间。对于设备上分配的内存，我们用cudaFree函数释放；对于主机上分配的内存，我们用free函数释放。
cudaFree(d_x); // 释放设备内存
cudaFree(d_y);
free(x); // 释放主机内存
free(y);
5、设备代码
__global__ void saxpy(int n, float a, float *x, float *y)
{
int i = blockIdx.x*blockDim.x threadIdx.x;
if (i < n) y[i] = a*x[i] y[i];
}
在CUDA中，我们使用__global__声明说明符定义了如saxpy的内核。设备代码中定义的变量不需要指定为设备变量，因为它们被假定驻留在设备上。在这种情况下 n，a 和 i 变量将由每个线程存储在寄存器中，并且指针 x 和 y 必须是指向设备内存地址空间的指针。这确实是正确的，因为从主机代码启动 d_x 和 d_y 时将它们传递给内核。但是，前两个参数 n 和 a 没有以主机代码显式传输到设备。由于在C / C 中，函数参数默认是按值传递的，因此CUDA运行时可以自动处理这些值到设备的传输。CUDA运行时API的此功能使在GPU上启动内核非常自然和容易-与调用C函数几乎相同。
saxpy内核中只有两行。 如前所述，内核由多个线程并行执行。如果希望每个线程处理结果数组的元素，则需要一种区分和标识每个线程的方法。 CUDA定义了变量 blockDim，blockIdx 和 threadIdx。这些预定义变量的类型为dim3，类似于主机代码中的执行配置参数。预定义的变量 blockDim 包含内核启动的第二个执行配置参数中指定的每个线程块的尺寸。预定义变量 threadIdx 和 blockIdx 分别包含其线程块内的线程的索引和网格内的线程块的索引。 表达方式：
int i = blockDim.x * blockIdx.x threadIdx.x
生成用于访问数组元素的全局索引。在此示例中，我们没有使用gridDim，gridDim包含在启动的第一个执行配置参数中指定的网格尺寸。
在使用此索引访问数组元素之前，将对它的值与元素数n进行检查，以确保没有越界内存访问。如果无法通过线程块大小将数组中的元素数整除，并且此内核启动的线程数大于数组大小，则需要进行此检查。内核的第二行执行saxpy的元素级工作，除了边界检查以外，它与saxpy主机实现的内部循环类似。
6、编译并运行代码
配置用的是Visual Studio 2017配置的cuda 10.1
运行的结果为：Max error: 0.000000
7、总结
通过SAXPY的简单CUDA C实现的演练，了解了CUDA C编程的基本知识。只需要对C进行一些扩展，即可将C代码“移植”到CUDA C __global__。启动内核时使用的执行配置；和内置设备变量blockDim，blockIdx以及threadIdx用于鉴定和并行执行内核分化GPU线程。
异构CUDA编程模型的一个优点是，可以将现有代码从C移植到CUDA C的过程可以逐步完成，一次可以一个内核。
,
