背景
随着深度学习的发展,并行计算的需求也越来越多,不论是算法工程师还是搞性能优化的,知道点cuda编程也是当前的必备技能之一。但是目前关于gpu资源很杂,重点不突出,我们需要在极短的时间内get到cuda的精华,本系列就由此而诞生。
文章框架
cuda精讲系列文章主要由三部分构成:
- cuda基础
- GPU架构
- cuda编程优化
每一部分都是通过知识点的方式将重点提炼出来,方便快速查看。
cuda基础
- cuda的基本概念
host: CPU
device: gpu
从软件层面讲:
(1)kernel:就是开发者开发的一个在gpu上要执行的函数代码
(2)thread:用来执行一个kernel的线程
(3)block:多个thread的组合
(4)grid:多个block的组合
(5)warp:block以一个warp为调度单位进行线程调度。因此block中的线程数尽量是32的整数倍。每32个连续线程号的线程被安排在同一个wrap里
从硬件层面讲:
(1)SP(core):一个物理上的SP对应一个逻辑上的thread
(2)SM(multicore processor):多个sp的组成,处于SM内的sp可以共享shared_mem,并且同一个block里的线程可以被同步。一个物理上的SM可以对应一个或多个逻辑上的block
(3)Device:就是一个GPU设备,一个GPU包括了多个SM。
- cuda编程框架
__global__ void my_kernel(...){
...
}
int main() {
...
cudaMalloc(...)
// host to device
cudaMemcpy(...)
...
my_kernel<<<nblock, blocksize>>>(...)
...
// device to host
cudaMemcpy(...)
...
}
- cuda编程常用语法:
- kernel launch:就是gpu上执行的kernel函数,因为执行kernel通常要从cpu登录到gpu因此称为launch,kernel函数语法为:
kernelFunc<<<nB, nT>>>
nB:block的数量
nT:thread的数量
- 内建变量
threadIdx;blockIdx;blockDim;gridDim - 同步语句
__syncthreads() - 生存周期生命:
global void kernelFunc(…),运行在device中,在cpu中调用
device void GlobalVar; 设备变量
shared void SharedVar; 定义每个block的shared memory中定义的变量
- cuda函数的异步性
cuda中的大部分函数其实都是异步的,比如:
- kernel launches
- 内存拷贝中可以指定option:cudaMemcpyAsync,cudaMemsetAsync
- cudaEvent function
- CPU与GPU的同步语句:
- Device based: cudaDeviceSynchronize(),在某一个设备GPUkernel执行结束后在执行cpu语句。
- Context(Thread) based: cudaThreadSynchronize(),所有设备的所有kernel都执行完毕后在执行。
- Stream based: cudaStreamSynchronize(stream-id)
- Event based:某个设备的kernel事件执行完毕后cpu语句再执行。
- 多GPU编程,后续补上
- 单机多卡:利用pthread、openMP
- 多机多卡:MPI,计算交给GPU,具体的通信细节交给MPI
- GPU之间的数据共享
(1)直接拷贝 host to device
(2)零拷贝:将device的内存直接映射到host的内存当中,但是在host中被共享的数据需要被page-locked(pinned):由此扩展的hostAPI有,
- cudaMallocHost(),分配pinned的host mem
(3)p2p内存拷贝:
直接拷贝数据从GPUA到GPUB当中,API为:cudaMemcpyPeer(void dst, int
dstDevice, const void src, int srcDevice,size_t count)
- 动态并行(Dynamic parallelism, dp)
避免从cpu到gpu的多次kernel登录,可以实现从gpu到gpu之间进行kernel登录
通过dp可实现动态的blocksize与gridsize分配: