01 CUDA简介
什么是CUDA
CUDA(Compute Unified Device Architecture),NVIDIA推出的的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。
将GPU视为数据并行计算设备,在其上进行计算的分配和管理,而无需将其映射到图形API。
CUDA包含一个让开发者能够使用C作为高级编程语言的软件环境。其它的语言和应用编程接口(API)也被支持,如CUDA FORTRAN,OpenCL和Direct Compute。
主机与设备
GPU可以看做并行执行非常多个线程的计算设备(compute device)。
CPU作为主机(host),一个系统可以存在一个主机和若干个设备。
CPU、GPU各自拥有相互独立的存储地址空间:主机端的内存和设备端的显存。
并行线程的组织结构
Thread:并行的基本单位,具有IDs。
Thread block:互相合作的线程组
Grid:一组thread block
Kernel:在GPU上执行的核心程序
一个kernel函数中只有一个grid
异构计算
- 运行在GPU上的CUDA并行计算函数称为kernel(内核函数)。
- 一个完整的CUDA程序是由一系列的设备端kernel函数并行步骤和主机端的串行处理步骤共同组成。
GPU线程映射关系
线程由SP(Scalar Processor)执行。
Thread blocks在SM上执行
Thread blocks do not migrate
一个内核函数作为一个grid启动
同一时刻GPU上只能运行一个内核函数
Grid、Block及Kernel函数
__global__ void KernelFunc(...);
dim3 DimGrid(4,8); //2*2 个Block
dim3 DimBlock(4,4); ////每个Block有 4*2*2 个线程
size_t SharedMemBytes = 32;//64 bytes of shared memory
KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>();
Kernel函数的定义与调用
运行在GPU上的程序成为kernel,必须通过global函数类型限定定义。
只能在host代码中调用。
//代码演示对长度为N的两个向量A和B求和,结果存在向量C中。
//kernel定义
__global__ void VecAdd(float* A,float* B,float* C) {
int i=threadIdx.x;
C[i]=A[i]+B[i];
}
int main() {
//kernel调用
VecAdd<<< 1, N >>>(A,B,C);
}
线程层次
由于块内的所有线程必须存在于同一个处理器核心中且共享该核心有限的存储器资源,因此,一个块内的线程数目是有限的。在目前的GPU上,一个线程块可以包含多达1024个线程。
一个内核可被多个同样大小的线程块执行,所以总的线程数等于每个块内的线程数乘以线程块数。
//将长度为N*N的两个矩阵A和B相加,存入矩阵C。
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){
int i = threadIdx.x;
int j = threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
}
int main() {
// Kernel 调用每个blockN*N*1个线程
int numBlocks = 1;
dim3 threadsPerBlock(N, N);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}
线程结构
线程被组织成1、2、3维线程网格。
对于一维的block,线程的threadID就是threadIdx.x
对于大小为(Dx,Dy)的二维block,线程的threadID是
(threadIdx.x + threadIdx.y*Dx)大小为(Dx,Dy,Dz)的三维block,线程的threadID是
(threadIdx.x + threadIdx.yDx + threadIdx.zDx*Dy)
CUDA关键字
__device__
储存于GPU上的global memory 空间
和应用程序具有相同的生命期(lifetime)
可被grid中所有线程存取
__constant__
储存于GPU上的constant memory 空间
和应用程序具有相同的生命期(lifetime)
可被grid中所有线程存取,CPU代码通过runtime函数存取
__shared__
存储于GPU上的thread block内的共享存储器
和thread block 具有相同的生命期
只能被thread block内的线程存取
Local变量
存储于SM内的寄存器和local memory
和thread具有相同的生命期
thread私有
CUDA函数及变量定义
global定义kernel函数
必须返回voiddevice和host可以组合使用
则被定义的函数在CPU和GPU上都被编译device函数
不能被&运算符取地址
不支持递归调用
不支持静态变量
不支持可变长度参数函数调用
CUDA存储空间管理
每个线程可以:
读/写每个thread registers
读/写每个thread local memory
读/写每个thread shared memory
读/写每个grid global memory
读/写每个grid constant memory
读/写每个block texture memoryCPU可以读写global,constant和texture内存
Register与Local Memory
对每个线程来说,寄存器都是线程私有的。
如果寄存器被消耗完,数据将被存储在local memory。
Local memory是私有的,但是local memory中的数据是被保存在显存中,速度很慢。输入和中间输出变量将被保存在register或者local memory中。
Shared Memory
用于线程间通信的shared memory。
shared memory是一块可以被同一block中的所有thread访问的可读写存储器。访问shared memory几乎和访问register一样快,
是实现线程间通信的延迟最小的方法。shared memory可以实现许多不同的功能,
如用于保存公用的计数器或者block的公用结构。
Constant Memory,Texture Memory
利用GPU用于图形计算的专用单元发展而来的高速只读缓存。
速度与命中率有关,不命中时将进行对显存的访问。
constant memory空间较小(只有64k),支持随机访问。
从host端只写,从device端只读。texture memory尺寸则大的多,并且支持二维寻址
内存空间分配
GPU Global Memeory
cudaMalloc(void ** devPtr,size_t size):
分配显存中的global memory内存中
devPtr:对象数组指针;size:数组尺寸cudaFree(void* devPtr):
释放显存中的global memory-
cudaMemcpy(void *dst,const void *src,size_t count,enum cudaMemcpyKind kind) - 数据交换
dst - 目的存储器地址
src - 源存储器地址
count - 拷贝数据的大小
kind - 数据传输类型1.cudaMemcpyDeviceToHost将显存中的数据拷贝到内存中。
2.cudaMemcpyHostToDevice将内存中的数据拷贝到显存中。
3.cudaMemcpyDeviceToDevice将global memory中的数据拷贝到同一CUDA上下文的global memory的另一区域中。
int Block_size=64;
float* Md;
int size=BLOCK_SIZE*BLOCK_SIZE*sizeof(float);
cudaMalloc((void**)&Md,size);
cudaFree(Md);
CUDA编程框架
- 声明全局变量
- main()
- CPU函数
- 内核函数
//全局变量声明
__host__,...,__device__,...,__global__,__constant__,__texture__//函数原型声明
__global__void kernelOne(...)//内核函数
float handyFunc(...)//CPU函数
main(){
cudaMalloc(&d_GlblVarPtr,bytes);//在设备上分配空间
cudaMemcpy(d_GlblVarPtr,h_Gl...);//从主机端传输数据到设备端
kernelOne<<<arg1,arg2>>>(arg...);//arg1:grid参数;arg2:block参数
cudaMemcpy(h_GlblVarPtr...);//从设备端传输数据到主机端
}
__global__void kernelOne(type args,...){//内核函数
//局部变量声明
__local__,__shared__//自动变量被默认分配到register或local memory
}
CUDA函数库
- 目前CUDA中有CUFFT、CUBLAS和CUDPP三个函数库,提供了简单高效的常用函数。
- CUFFT库是一个利用GPU进行傅里叶变换的函数库。
- CUBLAS库是一个基本的矩阵与向量的运算库。
- CUDPP库提供了很多基本的常用的并行操作函数,如排序、搜索等。
线程同步
- void __syncthreads();
- 使得block内所有线程同步。
- 只有当所有线程都达到同步语句处,后续指令才能继续执行。
- 用于避免访问共享或全局内存时出现的RAW / WAR / WAW冲突。
- 在最理想的情况下,调用一次__syncthreads()需要至少四个时钟周期,但一般调用都需要更多的时钟周期,因此尽量避免或节约使用__syncthreads()。
GPU与CPU线程同步
- 在CUDA主机端代码中使用cudaThreadSynchronize(),可以实现GPU与CPU线程的同步。
- kernel启动后控制权将异步返回,利用该函数可以确定所有设备端线程均已结束。
- 类似函数:cudaStreamSynchronize()和cudaEventSynchronize(),它们阻塞所有的流或事件,直到此前的所有CUDA调用均已完成。
GPU计算单元
- SM代表流多处理器
- 每个SM包含8个或更多标量流处理器SP,以及少量的其它计算单元
每个SM具有单指令多数据(SIMD)架构
每个多处理器具有下列四种类型的片上存储器
每个处理器有一组本地32位寄存器。
并行数据高速缓存(shared memory),由所有处理器共享并实现共享内存空间。
只读constant cache由所有处理器共享并加速从常量内存空间的读取。
只读texture cache由所有处理器共享并加速从纹理内存空间的读取。
Warp
- 若干线程捆绑在一起执行成为warp。
- 每个block根据thread ID将thread分为多个warp。
- warp是SM内的基本调度单位。在任何时刻,硬件只选择一个warp执行。
- 划分warp的方式始终相同,每个warp包含thread ID连续递增的线程,其中第一个warp从0开始。
执行模型
- CUDA采用了SIMT(Single Instruction,Multiple Thread,单指令多线程)执行模型,是SIMD的一种改进。
- 重要的不同在于SIMD组织方法会向应用暴露SIMD宽度,而SIMT指定单线程的执行和分支行为。
- 与SIMD向量机相反,SIMT允许程序员为独立标量线程编写线程级并行代码,也为协作线程编写数据并行代码。
- 为了正确性,程序员可忽略SIMT行为;然而只要维护束内线程很少分支的代码就可显著提升性能。
SIMD (Single Instruction Multiple Data)
- 也可以认为是数据并行分割