CUDA编程基础

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的可选类型有：

1. cudaMemcpyHostToHost
2. cudaMemcpyHossToDevice
3. cudaMemcpyDeviceToHost
4. 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<<<grid, block>>>(argument list);

其中grid和block即为上文中介绍的类型为dim3的变量。通过这两个变量可以配置一个kernel的线程总和，以及线程的组织形式。例如：

kernel_name<<<4, 8>>>(argumentt list);

该行代码表明有grid为一维，有4个block，block为一维，每个block有8个线程，故此共有4*8=32个线程。

注意，不同于c函数的调用，所有CUDA kernel的启动都是异步的，当CUDA kernel被调用时，控制权会立即返回给CPU。

函数类型标示符

__device__ 和__host__可以组合使用。

kernel的限制：

• 仅能获取device memory 。
• 必须返回void类型。
• 不支持可变数目参数。
• 不支持静态变量。
• 不支持函数指针。
• 异步。