CUDA语言模型通过对C/C++语言的简单标记，就可以将不同代码指定
到异构计算系统（所谓异构，就是指不同硬件，包含cpu与gpu）上执行。
对于异构计算系统的不同硬件之间（cpu与gpu）使用PCIE进行数据交换。
（由于PEIC的传输效率慢，之后也会有更多的技术对此进行改进）
所以CUDA的内存模型分为两个部分：
（1） Host memory：CPU的内存
（2） Device memory：GPU的内存
CUDA的整体代码架构也分为两个部分：
（1）CPU运行部分：主要做逻辑控制及内存显存的申请销毁。
（2）GPU运行部分：也叫kernel，是真正启动CUDA并行计算之处。
一个完整的CUDA代码基本都是由以下代码组件构成：
（1）在CPU代码部分申请cpu与gpu的内存
（2）在CPU代码部分将数据从cpu拷贝到gpu
（3）在CPU代码部分启动kernel（launch kernel）
（4）在GPU代码部分（kernel）执行相应的操作
（5）在CPU代码部分将数据从gpu拷贝到cpu
（6）释放申请的内存
CUDA的编程模型是假设一个异构系统是由host和device组成，
并且两者之间都有各自分开独立的内存空间。使用device的内存也
需要像使用host内存一样进行申请使用，并且都在host端进行申请。
为了统一编码风格，CUDA将申请释放内存接口设计成类似C语言：
其中，申请内存的函数原型为：
cudaError_t cudaMalloc ( void** devPtr, size_t size )。
这个函数会申请一块连续的内存，内存的头指针为devPtr，大小为size
个字节。
内存拷贝的函数原型为：
cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )。
这个函数会拷贝数据，从dst拷贝到src，长度为count字节。至于最后一个
参数 cudaMemcpyKind 是一个枚举值，可选的有：
（1） cudaMemcpyHostToHost
（2） cudaMemcpyHostToDevice
（3） cudaMemcpyDeviceToHost
（4） cudaMemcpyDeviceToDevice
从字面上可以清楚看到，这个枚举值是控制数据是从何处拷贝到何处。
以上两个函数的返回值都是cudaError_t（是一个枚举）。如果操作执行成
功，那么返回 cudaSuccess。如果不成功，可以用下面这个函数将具体信息打印
出来： char* cudaGetErrorString(cudaError_t error)
为了能够更加直观的体验整个过程，以下是一个例子。在例子里实现的功能是
把a向量和b向量相加到c向量。
先看看C代码是怎么书写：
#include <stdlib.h>
#include <string.h>
#include <time.h>
void sumArraysOnHost(float *A, float *B, float *C, const int N) {
for (int idx=0; idx<N; idx++) {
C[idx] = A[idx] + B[idx];
}
}
void initialData(float *ip,int size) {
// 初始化数据
time_t t;
srand((unsigned int) time(&t));
for (int i=0; i<size; i++) {
ip[i] = (float)( rand() & 0xFF )/10.0f;
}
}
int main(int argc, char **argv) {
int nElem = 1024;
size_t nBytes = nElem * sizeof(float);
float *h_A, *h_B, *h_C;
h_A = (float *)malloc(nBytes);
h_B = (float *)malloc(nBytes);
h_C = (float *)malloc(nBytes);
initialData(h_A, nElem);
initialData(h_B, nElem);
sumArraysOnHost(h_A, h_B, h_C, nElem);
free(h_A);
free(h_B);
free(h_C);
return(0);
}
这是个纯C代码，所以既可以用C的编译器进行编译，也可以用CUDA编译器 nvcc进行编译。
$ nvcc -Xcompiler -std=c99 sumArraysOnHost.c –o sum
$ ./sum
现在修改以上代码，使用cudaMalloc进行对GPU的内存申请。
float *d_A, *d_B, *d_C;
cudaMalloc((float**)&d_A, nBytes);
cudaMalloc((float**)&d_B, nBytes);
cudaMalloc((float**)&d_C, nBytes);
接着需要把CPU上的数据传输到GPU：
cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);
现在GPU上数据已经具备了，在CPU端可以进行launch kernel，启动计算了。需要注意的是：
CPU端launch kernel之后，控制权又回到了CPU端，CPU可以继续执行后续的代码流程，也就是GPU
在kernel计算和CPU端的后续代码是异步执行的。当GPU在kernel中计算完毕，会将数据保存到d_C中，
如果需要在CPU中打印，需要再次使用以下函数将数据搬运回CPU：
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
最后，使用cudaFree进行释放GPU内存。
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
需要注意的是：不能直接用gpuRef = d_C，因为gpuRef是CPU的地址，而d_C是GPU的地址，
二者不能直接赋值。
当host端发起kernel的launch，device端（gpu）就要开始分配资源并且执行kernel。
了解device的线程组织是CUDA编程中很重要的一环。CUDA的线程层次体系能够让使用者自己
来组织线程。下面是一张2维的线程组织方式：
在一次kernel启动中，所有的线程都被分配到一个叫做grid的单元中。在一个grid中的
所有线程有一个共享的全局内存区域，叫做global memory。一个grid是由很多线程块
（thread blocks）组成。一个线程块之间的所有的线程是可以互相通信，有两种方式：
（1）线程块级别的同步
（2）线程块的共享内存（shared memory）
在不同线程块中的线程之间是无法通信。
线程是由下面两个维度的坐标进行唯一确定的：
（1） blockIdx （属于grid中的哪一个block）
（2） threadIdx （属于block中的哪一个线程）
这两个参数是由CUDA提前内置的，每个线程都会有。这种CUDA内置的参数的类型是uint3，它是
个包含三个无符号整型的结构体，可以通过
blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z
来获取对应坐标值。此外，还有两个内置的变量：blockDim和gridDim。这两种结构的数据类型
是dim3，其构成和unit3是一样的。主要用来指示block的大小和grid的大小。在host端启动kernel
的时候就应该指定。特别的，如果有些维度没有显示指定，那么就是为1。
接着以上的例子，对于指定grid和block的大小应该根据数据量的多少来指定。上述例子中，
int nElem = 6;因此我们可以定义1D的block包含三个线程，1D的grid包含2个block：
dim3 block(3);
dim3 grid((nElem+block.x-1)/block.x);
可以通过在host端打印以下信息来显示grid和block的大小：
printf("grid.x %d grid.y %d grid.z %d\n",grid.x, grid.y, grid.z);
printf("block.x %d block.y %d block.z %d\n",block.x, block.y, block.z);
在device端，可以打印以下内容来确定：
printf("threadIdx:(%d, %d, %d) blockIdx:(%d, %d, %d) blockDim:(%d, %d, %d) "
"gridDim:(%d, %d, %d)\n", threadIdx.x, threadIdx.y, threadIdx.z,
blockIdx.x, blockIdx.y, blockIdx.z, blockDim.x, blockDim.y, blockDim.z,
gridDim.x,gridDim.y,gridDim.z);
现在准备工作都做好，可以进行启动kernel。对于kernel的调用，是通过三重尖括号进行
调用：kernel_name <<<grid, block>>>(argument list);
例如：kernel_name<<<4, 8>>>(argument list);其组织形式如下图：
kernel的调用和host的线程是异步执行的，当kernel被启动，控制权又回到host端，
但是可以通过以下函数让host强制停止，等待kernel完成。
cudaError_t cudaDeviceSynchronize(void);
前文提到，需要有一些标识区分kernel和C代码。这个标识就是__global__。
例如：__global__ void kernel_name(argument list);
下面是一些CUDA的标识符：
其中__device__和__host__可以同事使用，表明该函数需要被同时编译成host代码和
device代码。
那么以上的例子最终被改写成kernel就应该如下所示：
__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
启动的时候就应该是：sumArraysOnGPU<<<1,32>>>(float *A, float *B, float *C);
现在同时有了CPU和GPU的结果，两者之间应该是一致的。需要对最后的结果进行比对。
void checkResult(float *hostRef, float *gpuRef, const int N) {
double epsilon = 1.0E-8;
int match = 1;
for (int i = 0; i < N; i++) {
if (abs(hostRef[i] - gpuRef[i]) > epsilon) {
match = 0;
printf("Arrays do not match!\n");
printf("host %5.2f gpu %5.2f at current %d\n",
hostRef[i], gpuRef[i], i);
break;
}
}
if (match) printf("Arrays match.\n\n");
return;
}
CUDA的API大都会返回一个cudaError_t的枚举，对是否成功完成进行标识，可以定义
一个宏，来对错误结果进行展示：
#define CHECK(call) \
{ \
const cudaError_t error = call; \
if (error != cudaSuccess) \
{ \
printf("Error: %s:%d, ", __FILE__, __LINE__); \
printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
exit(1); \
} \
}
比如：CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));
#include <cuda_runtime.h>
#include <stdio.h>
#include <time.h>
#define CHECK(call) \
{ \
const cudaError_t error = call; \
if (error != cudaSuccess) \
{ \
printf("Error: %s:%d, ", __FILE__, __LINE__); \
printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
exit(1); \
} \
}
void checkResult(float *hostRef, float *gpuRef, const int N) {
double epsilon = 1.0E-8;
int match = 1;
for (int i = 0; i < N; i++) {
if (abs(hostRef[i] - gpuRef[i]) > epsilon) {
match = 0;
printf("Arrays do not match!\n");
printf("host %5.2f gpu %5.2f at current %d\n",
hostRef[i], gpuRef[i], i);
break;
}
}
if (match) printf("Arrays match.\n\n");
return;
}
void sumArraysOnHost(float *A, float *B, float *C, const int N) {
for (int idx=0; idx<N; idx++) {
C[idx] = A[idx] + B[idx];
}
}
__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
void initialData(float *ip,int size) {
// 初始化数据
time_t t;
srand((unsigned int) time(&t));
for (int i=0; i<size; i++) {
ip[i] = (float)( rand() & 0xFF )/10.0f;
}
}
int main(int argc, char **argv) {
printf("%s Starting...\n", argv[0]);
// set up device
int dev = 0;
cudaSetDevice(dev);
// set up data size of vectors
int nElem = 32;
printf("Vector size %d\n", nElem);
// malloc host memory
size_t nBytes = nElem * sizeof(float);
float *h_A, *h_B, *hostRef, *gpuRef;
h_A = (float *)malloc(nBytes);
h_B = (float *)malloc(nBytes);
hostRef = (float *)malloc(nBytes);
gpuRef = (float *)malloc(nBytes);
// initialize data at host side
initialData(h_A, nElem);
initialData(h_B, nElem);
memset(hostRef, 0, nBytes);
memset(gpuRef, 0, nBytes);
// malloc device global memory
float *d_A, *d_B, *d_C;
cudaMalloc((float**)&d_A, nBytes);
cudaMalloc((float**)&d_B, nBytes);
cudaMalloc((float**)&d_C, nBytes);
// transfer data from host to device
cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);
// invoke kernel at host side
dim3 block (nElem);
dim3 grid (nElem/block.x);
sumArraysOnGPU<<< grid, block >>>(d_A, d_B, d_C);
printf("Execution configuration <<<%d, %d>>>\n",grid.x,block.x);
// copy kernel result back to host side
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
// add vector at host side for result checks
sumArraysOnHost(h_A, h_B, hostRef, nElem);
// check device results
checkResult(hostRef, gpuRef, nElem);
// free device global memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
// free host memory
free(h_A);
free(h_B);
free(hostRef);
free(gpuRef);
return(0);
}
1157
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
