本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/,仅用于学习。
CUDA是并行计算的平台和类C编程模型,我们能很容易的实现并行算法,就像写C代码一样。只要配备的NVIDIA GPU,就可以在许多设备上运行你的并行程序,无论是台式机、笔记本抑或平板电脑。熟悉C语言可以帮助你尽快掌握CUDA。
CUDA编程
CUDA编程允许你的程序执行在异构系统上,即CUP和GPU,二者有各自的存储空间,并由PCI-Express 总线区分开。因此,我们应该先注意二者术语上的区分:
- Host:CPU and itsmemory (host memory)
- Device: GPU and its memory (device memory)
代码中,一般用h_前缀表示host memory,d_表示device memory。
kernel是CUDA编程中的关键,他是跑在GPU的代码,用标示符__global__注明。
host可以独立于host进行大部分操作。当一个kernel启动后,控制权会立刻返还给CPU来执行其他额外的任务。所以,CUDA编程是异步的。一个典型的CUDA程序包含由并行代码补足的串行代码,串行代码由host执行,并行代码在device中执行。host端代码是标准C,device是CUDA C代码。我们可以把所有代码放到一个单独的源文件,也可以使用多个文件或库。NVIDIA C编译器(nvcc)可以编译host和device生成可执行程序。
这里再次说明下CUDA程序的处理流程:
- 从CPU拷贝数据到GPU。
- 调用kernel来操作存储在GPU的数据。
- 将操作结果从GPU拷贝至CPU。
Memory操作
cuda程序将系统区分成host和device,二者有各自的memory。kernel可以操作device memory,为了能很好的控制device端内存,CUDA提供了几个内服务器托管网存操作函数:
为了保证和易于学习,CUDA C 的风格跟C很接近,比如:
cudaError_t cudaMalloc ( void** devPtr, size_t size )
我们主要看看cudaMencpy,其函数原型为:
cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )
其中cudaMemcpykind的可选类型有:
- cudaMemcpyHostToHost
- cudaMemcpyHossToDevice
- cudaMemcpyDeviceToHost
- cudaMemcpuDeviceToDevice
具体含义很好懂,就不多做解释了。
对于返回类型cudaError_t,如果正确调用,则返回cudaSuccess,否则返回cudaErrorMemoryAllocation。可以使用char* cudaGetErrorString(cudaError_t error)将其转化为易于理解的格式。
组织线程
掌握如何组织线程是CUDA编程的重要部分。CUDA线程分成Grid和Block两个层次。
由一个单独的kernel启动的所有线程组成一个grid,grid中所有线程共享global memory。一个grid由许多block组成,block由许多线程组成,grid和block都可以是一维二维或者三维,上图是一个二维grid和二维block。
这里介绍几个CUDA内置变量:
- blockIdx:block的索引,blockIdx.x表示block的x坐标。
- threadIdx:线程索引,同理blockIdx。
- blockDim:block维度,上图中blockDim.x=5.
- gridDim:grid维度,同理blockDim。
一般会把grid组织成2D,block为3D。grid和block都使用dim3作为声明,例如:
dim3 block(3);
// 后续博文会解释为何这样写grid
dim3 grid((nElem+block.x-1)/block.x);
需要注意的是,dim3仅为host端可见,其对应的device端类型为uint3。
启动CUDA kernel
CUDA kernel的调用格式为:
kernel_name>>(argument list);
其中grid和服务器托管网block即为上文中介绍的类型为dim3的变量。通过这两个变量可以配置一个kernel的线程总和,以及线程的组织形式。例如:
kernel_name>>(argumentt list);
该行代码表明有grid为一维,有4个block,block为一维,每个block有8个线程,故此共有4*8=32个线程。
注意,不同于c函数的调用,所有CUDA kernel的启动都是异步的,当CUDA kernel被调用时,控制权会立即返回给CPU。
函数类型标示符
__device__ 和__host__可以组合使用。
kernel的限制:
- 仅能获取device memory 。
- 必须返回void类型。
- 不支持可变数目参数。
- 不支持静态变量。
- 不支持函数指针。
- 异步。
代码分析
#include
#include
#define CHECK(call)
{
const cudaError_t error = call;
if (error != cudaSuccess)
{
printf("Error: %s:%d, ", __FILE__, __LINE__);
printf("code:%d, reason: %sn", error, cudaGetErrorString(error));
exit(1);
}
}
void checkResult(float *hostRef, float *gpuRef, const int N) {
double epsilon = 1.0E-8;
bool match = 1;
for (int i=0; i epsilon) {
match = 0;
printf("Arrays do not match!n");
printf("host %5.2f gpu %5.2f at current %dn",hostRef[i],gpuRef[i],i);
break;
}
}
if (match) printf("Arrays match.nn");
}
void initialData(float *ip,int size) {
// generate different seed for random number
time_t t;
srand((unsigned) time(&t));
for (int i=0; i>>(d_A, d_B, d_C);
printf("Execution configuration >>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);
}
服务器托管,北京服务器托管,服务器租用 http://www.fwqtg.net
Relocation(重定位)是一种将程序中的一些地址修正为运行时可用的实际地址的机制。在程序编译过程中,由于程序中使用了各种全局变量和函数,这些变量和函数的地址还没有确定,因此它们的地址只能暂时使用一个相对地址。当程序被加载到内存中运行时,这些相对地址需要被…