cuda 编程学习笔记

programming model

kernels

类似于c函数，函数定义使用__global__声明，使用<<<…>>>形式的execution configuration决定kernal执行的线程数，使用threadIdx变量可以获得每一个线程对应的id。

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}
int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

execution configuration定义形式为

<<< Dg, Db, Ns, S >>>
Dg: 表示grid的维度和大小，对应的block个数为Dg.x * Dg.y * Dg.z
Db: 表示block的维度和大小，对用的thread个数为Db.x * Db.y * Db.z
Ns: 表示shared memory的大小，可选参数
S : 表示cudaStream_t的类型，可选参数

thread hierarchy

thread：
block: 多个thread，个数类型可以是int/dim3，一般最多1024个thread
grid: 多个block，个数类型可以是int/dim3
block之间独立执行，block内的thread可以通过shared memory共享数据，通过调用__syncthreads()函数实现同步。

• threadIdx: block内thread的id
• blockIdx: grid内的block的id
• blockDim: 一个block的thread的个数

memory hierarchy

每个thread有local memory
每个block有shared memory，block内的thread都可以访问
所有的thread都可以访问global memory/const memory/texture memory，其中const memory和texture memory只读。

heterogeneous programming

host和device维护DRAM中不同的内存空间，分别称为host memory和device memory。

compute capability

表示GPU硬件的特性，也称为“SM version”，由两部分组成"X.Y"，同一个X的版本号对应的内核架构一致，其中5表示基于Maxwell架构，3表示Kepler架构，2表示Fermi架构，1表示Tesla架构。从CUDA7.0开始不再支持Tesla架构。

programming interface

NVCC编译

编译分为offline compilation或者just-in-time compilation

1. offline compilation
   cuda的代码可以是PTX或者c，都需要通过nvcc进行编译，编译流程

• 分离host code和device code
• 编译device code为PTX汇编格式或者cubin二进制格式
• 使用编译好的kernel替换host code中的kernel函数（以<<<…>>>语法标记的）
• 修改后的host code使用其他工具编译
• nvcc编译后的程序可以链接到编译好的host code或者直接通过CUDA driver API进行加载执行

2. just-in-time compilation
   运行时加载的PTX程序，使用device driver进一步编译成二进制程序，称为just-in-time compilation.
   这种编译方式增加了加载时间，但是可以使用新的device driver带来的特性。

-code指定二进制程序运行的目标设备的架构，比如-code=sm_35产生的二进制程序运行在compute capability 3.5
-arch指定的架构针对c代码到PTX汇编的编译过程
支持部分c/c++的语法格式
64-bit模式编译的device code只能被64-bit模式编译的host code支持；32-bit的同理。
32-bit的nvcc可以使用-m64选项编译64-bit的device code；64-bit的nvcc可以使用-m32选项编译32-bit的device code。

cuda c runtime

通过cudart库文件实现: libcudart.a或者libcudart.so。

initialization

第一次调用runtime函数的时候执行初始化，为device创建cuda context，所有的host线程共享。
cudaDeviceReset()函数销毁context。

device memory

分为linear memory和CUDA arrays两种。
linear memory是device上的40-bit的地址空间。

• cudaMalloc/cudaFree/cudaMemcpy
• cudaMallocPitch/cudaMalloc3D/cudaMemcpy2D/cudaMemcpy3D: 分配或者复制2D/3D数组，满足对齐要求/高效，尽量使用

shared memory

使用__shared__标示，shared memory比global memory快，尽量使用shared memory。

page-locked host memory

优点：

• page-locked host memory和device memory之间内存拷贝可以和kernel执行同时进行
• 可以map到device的地址空间
• 具有front-side bus的系统，page-locked host memory和device memory有更高的带宽

异步并行

以下操作相互独立，可以同步执行：

• host计算
• device计算
• host和device回见的内存转移
• device内部和device之间的内存转移

多设备系统

cudaGetDeviceCount: 获取设备数
cudaSetDevice: 设置使用设备

统一虚拟地址空间

对于64-bit的进程，host和compute capability 2.0及更高版本的device使用一个地址空间

进程间通信

使用Inter Process Communication API完成

version

两个版本需要关注：compute capability有关硬件设备的特性；CUDA driver api有关driver API和runtime。
driver API的版本使用CUDA_VERSION定义，支持向后兼容，也就是低版本的程序可以在高版本的driver上面运行。

compute modes

可以使用nvidia-smi设置为一下几种模式：

• default: 多个host线程可以同时使用device
• exclusive-process: 只能在device创建一个CUDA context，创建context的进程中的多个线程可以使用
• exclusive-process-and-thread: 只允许创建一个context，而且context一次只能被一个线程使用
• prohibit: 禁止创建CUDA context

c语言扩展

函数类型：

• _device_: 在device上执行，只能被device调用
• _global_: 标示kernel函数，在device上执行，可以被host/device调用
• _host_: 在host执行，只能被host调用
• _noinline_: 函数不inline
• _forceinline_: 函数inline

变量类型：

• _device_: device上的变量
• _constant_: constant memory space上的变量
• _shared_: shared memory，所有以这种形式声明的指针具有相同的起点
• _managed_: host和device都可以读写的地址
• _restrict_: 避免aliasing问题

内置vector

make_<type name>

vector是结构体，可以通过x/y/z/w来获取第1/2/3/4个元素。
dim3用来表示维度，等同于uint3。

内置变量

gridDim: grid维度
blockIdx: grid中的block index
blockDim: block维度
threadIdx: block中thread的index
warpSize: 线程中的warp size

性能优化

并行

并行库，比如cuBLAS/cuFFT
并行编译器，比如使用progma
并行代码

debug

CUDA_GDB

timing

1. CPU Timer
   注意CUDA API函数很多是异步的，在使用计时函数前需要调用cudaDeviceSynchronize()函数以同步CPU线程和GPU。

2. GPU Timer
   使用event计时

cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

参考

cuda-c-programming-guide
cuda-c-best-practices-guide

后面的技术分享转移到微信公众号上面更新了，【欢迎扫码关注交流】