一、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的过程可以逐步完成,一次可以一个内核。
,