一、CUDA编程模型基础
在进入CUDAC代码之前，CUDA的那些新知识将受益于CUDA编程模型的基本描述以及所使用的一些术语。
CUDA编程模型是一种异构模型，其中同时使用了CPU和GPU。在CUDA中，主机是指CPU及其内存，而设备是指GPU及其内存。在主机上运行的代码可以管理主机和设备上的内存，还可以启动内核，这些内核是在设备上执行的功能。这些内核由许多GPU线程并行执行。
鉴于CUDA编程模型的异构性质，CUDAC程序的典型操作序列为：
（1）、声明并分配主机和设备内存。
（2）、初始化主机数据。
（3）、将数据从主机传输到设备。
（4）、执行一个或多个内核。
（5）、将结果从设备传输到主机。
牢记此操作顺序，让我们看一下CUDAC示例。
1、第一个CUDAC程序
函数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中的主机和设备具有单独的内存空间，这两个内存空间都可以通过主机代码进行管理（CUDAC内核也可以在支持该内存的设备上分配设备内存）。
主机上的变量做初始化操作，我们设置数组x的所有元素为1，并设置数组y的所有元素为2。
for(inti=0;iN;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+)/,(N,2.0,d_x,d_y);
三个尖括号之间的信息是对GPU进行配置，该配置决定了有多少个设备线程并行执行该内核。在CUDA中，软件中存在线程层次结构，该层次结构模仿了如何在GPU上对线程处理器进行分组。在CUDA编程模型中，我们谈到启动带有线程块网格的内核。执行配置中的第一个参数指定网格中线程块的数量，第二个参数指定线程块中线程的数量。
通过为这些参数传递dim3（由CUDA定义的简单结构，带有x和y，以及z成员）值，可以将线程块和网格制成一维，二维或三维，但是对于这个简单的示例，我们只需要一个维，因此我们可以传递整数。在这种情况下，我们使用包含个线程的线程块启动内核，并使用整数算法确定处理N数组的所有元素（(N+)/）所需的线程块数。
对于无法通过线程块大小将数组中的元素数整除的情况，内核代码必须检查越界内存访问。即设置每个线程块中存在个线程，但是我们需要确定需要有多少个线程块。我们有N的数据需要并行处理，那么我们通过（N+）/来确定多少个线程块，对于多出来的线程（即最后一个块可能不满）我们必须做越界处理。
4、释放分配的所有资源
程序运行结束之后，需要释放我们之前分配的内存空间。对于设备上分配的内存，我们用cudaFree函数释放；对于主机上分配的内存，我们用free函数释放。
cudaFree(d_x);//释放设备内存
cudaFree(d_y);
free(x);//释放主机内存
free(y);
5、设备代码
__global__voidsaxpy(intn,floata,float*x,float*y)
{
inti=blockIdx.x*blockDim.x+threadIdx.x;
if(in)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分别包含其线程块内的线程的索引和网格内的线程块的索引。表达方式：
inti=blockDim.x*blockIdx.x+threadIdx.x
生成用于访问数组元素的全局索引。在此示例中，我们没有使用gridDim，gridDim包含在启动的第一个执行配置参数中指定的网格尺寸。
在使用此索引访问数组元素之前，将对它的值与元素数n进行检查，以确保没有越界内存访问。如果无法通过线程块大小将数组中的元素数整除，并且此内核启动的线程数大于数组大小，则需要进行此检查。内核的第二行执行saxpy的元素级工作，除了边界检查以外，它与saxpy主机实现的内部循环类似。
6、编译并运行代码
配置用的是VisualStudio配置的cuda10.1
运行的结果为：Maxerror:0.
7、总结
通过SAXPY的简单CUDAC实现的演练，了解了CUDAC编程的基本知识。只需要对C进行一些扩展，即可将C代码“移植”到CUDAC__global__启动内核时使用的执行配置；和内置设备变量blockDim，blockIdx以及threadIdx用于鉴定和并行执行内核分化GPU线程。
异构CUDA编程模型的一个优点是，可以将现有代码从C移植到CUDAC的过程可以逐步完成，一次可以一个内核。