最近在学习CUDA,感觉看完就忘,于是这里写一个导读,整理一下重点
主要内容来源于NVIDIA的官方文档《CUDA C Programming Guide》,结合了另一本书《CUDA并行程序设计 GPU编程指南》的知识。 因此在翻译总结官方文档的同时,会加一些评注,不一定对,望大家讨论指出。
另外,我才不会老老实实的翻译文档,因此细节还是需要从文档里看的。
看完两份文档总的来说,感觉《CUDA C Programming Guide》这本书作为一份官方文档,知识细碎且全面,且是针对最新的Maxwell、Pascal、Volta架构的阐述。但相对来说不够深入,且有关程序设计方面所述甚少。
而《CUDA并行程序设计 GPU编程指南》这本书,讲解的比较深入,不仅阐述了NVIDIA GPU的特性,并且在程序设计方面有比较深入的见解。美中不足的是该书是针对老旧的Tesla、Fermi架构GPU,没有涉及到新架构的新特性。
GPU是能够高度并行化、具有很多处理器核心的器件,具有很强的计算能力和内存带宽。下图是CPU和GPU在浮点运算上的性能对比发展趋势。
可以看到,NVIDIA的GPU在浮点运算能力上,吊打了Intel的CPU。其原因来自于CPU和GPU结构上的差异。
如下图所示,CPU仅仅具有有限的核心数量。相比于GPU,CPU的核心属于「少而精」的存在,核心数虽然很少,但是每个核心的性能很强,适合处理具有很多分支的复杂的逻辑。近些年来,CPU中集成了一些并行指令集,如SSE、AVX等,其中AVX可以同时处理256位(32个位元组),可以大大加速并行计算。但是相比于GPU,还是小巫见大巫。
GPU的设计理念与CPU不同,GPU具有庞大的核心数。以TITANX为例,流处理器(等同于CPU的核心)达到3072个之多。这意味著相比于8核的CPU处理器,TITANX可以同时并行处理384倍的任务。但是GPU的单个核心不如CPU的核心强大。因此相对于CPU,GPU更适合处理高度并行化的任务。
从另一个视角来看上图,在CPU晶元中,运算单元(ALU)所占的比例较小,CPU中更多的矽片被用来制作控制单元和缓存,以完成复杂的逻辑;而GPU的运算单元使用的矽片面积比例要大于CPU,以完成高强度的计算。因此CPU的侧重点在于逻辑控制,而GPU的侧重点在于计算。
深度学习,尤其是卷积神经网路中,有很多可以高度并行化的向量运算与矩阵运算。因此使用GPU进行深度学习运算,远比CPU快速。
另外还需注意的是,由于GPU的核心数量太多,即使GPU的内存(也称作显存)优于CPU的内存,其内存仍然是瓶颈。因此,GPU希望程序是 计算密集型 而不是 内存密集型。
CUDA(Compute Unified Device Architecture),是NVIDIA推出的通用并行计算平台和编程模型。CUDA是在底层API的基础上,封装了一层,使得程序员可以使用C语言来方便的编程。
CUDA还支持C++/Python等更高级的语言编程;此外,NVIDIA还提供了CuDNN、TensorRT、NPP等更高级的库函数。
从上图中也可以看出各个系列的GPU属于哪些架构、什么定位。例如GeForece 1000系列,就是使用Pascal架构的消费显卡。
CUDA的编程模型,使得同一个CUDA程序,可以在不同的显卡上运行。
如上图所示,CUDA程序一般会创建一些线程块(Block),线程块会被调度到空闲的流处理器簇(SM)上去。当线程块执行完毕后,线程块会退出SM,释放出SM的资源,以供其他待执行线程块调度进去。
因此,无论是只有2个SM的GPU,还是有4个SM的GPU,这些线程块都会被调度执行,只不过是执行的时间有长有短。因此,同样的程序,可以在具有不同SM数量上的GPU运行。
略
在讲解内核函数前,先讲解一下线程层级,不然有点难讲。
CUDA编程是一个多线程编程,数个线程(Thread)组成一个线程块(Block),所有线程块组成一个线程网格(Grid),如下图所示:
图中的线程块,以及线程块中的线程,是按照2维的方式排布的。实际上,CUDA编程模型允许使用1维、2维、3维三种方式来排布。另外,即使线程块使用的是1维排布,线程块中的线程也不一定要按照1维排,而是可以任意排布。
目前的GPU限制一个线程块中,最多可以安排1024个线程。
一个线程块用多少线程,以及一个线程网格用多少线程块,是程序员可以自由安排的。由于32个相邻的线程会组成一个线程束(Thread Warp),而一个线程束中的线程会运行同样的指令。因此一般线程块中线程的数量被安排为32的倍数,选用256是比较合适的。
在线程数定下来之后,一般根据数据的排布情况来确定线程块的个数。
例如:一个数组的长度为4096,安排每个线程处理一个元素。如果安排一个线程块为256个线程,则需要4096/256=16个线程块。
内核函数是CUDA 每个线程 执行的函数。CUDA使用扩展的C语言编写内核函数,关键字为global。内核函数返回值只能是void。
下面是一段简单的内核函数,用于求两个数组的和:
// 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); ... }
threadIdx.x是线程在所处线程块中的X方向的ID。由于本例中是定义的1维排布,因此X方向ID即为线程的ID。
threadIdx.x
由于GPU中的每个线程都会执行相同的VecAdd函数,因此不同的线程需要使用自己独有的ID来区分彼此,来获取不同的数据。这就是SIMT的概念,即「相同指令,不同线程」。
在main()函数中,我们注意到,VecAdd函数的调用使用了<<<blockPerGrid, threadsPerBlock>>>关键字。这是调用内核函数所独有的。程序员通过该关键字,制定网格中线程块和线程的排布方式。排布方式与数据息息相关。
<<<blockPerGrid, threadsPerBlock>>>
下面举一个2维排布的例子,用于做矩阵加法:
// Kernel definition __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) { int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i < N && j < N) C[j][i] = A[j][i] + B[j][i]; }
int main() { ... // Kernel invocation dim3 threadsPerBlock(16, 16); dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y); MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); ... }
首先看主函数,当排布不使用1维时,需要使用dim3数据类型。该程序每个线程块中线程为16x16排布,而线程块的排布依赖于数据的多少。
dim3
在内核函数中,i代表x方向上的ID,j代表y方向上的ID。blockDim代表当前线程块的尺寸。从程序中可以看到,x方向为行方向,y方向为列方向。(注意,这里官方文档里面写的有些错误)
blockDim
每个线程读取自己ID对应的数据A[j][i]和B[j][i],并将结果写回C[j][i]。其中A、B、C都存储在GPU的全局内存上(后面会提及)
同CPU一样,GPU也有不同层级的内存。越靠近核心的内存速度越快,但容量越小;反之,越远离核心的内存速度越慢,但容量较大。
上图是NVIDIA设备的硬体示意图。
这里想强调一下共享内存。
以上是从硬体的角度解读了一下GPU的内存层级。从编程角度来看,CUDA的线程网格、线程块、线程与各个内存见的关系如下图:
一种最简单的CPU/GPU混合编程如下图所示:
主机端(Host,即CPU)执行串列代码,然后调用内核函数,让设备端(Device,即GPU)执行并行代码。如此交错执行。
CPU和GPU的内存是独立的。因此在运行内核函数前,主机端需要调用内存拷贝函数,将数据通过PCI-E汇流排拷贝到设备端。内核运行结束后,需要CPU再次调用内存拷贝函数,将数据拷回主机端内存。
另一种方式是使用统一编址,将设备端的内存和主机端内存编到一起。这样主机就不需要显式的调用函数将数据拷贝到设备端内存了。
除了CPU/GPU交错执行代码的方式外,还可以通过使用事件(event)和流(stream)等方式,让CPU/GPU并行工作,提升整体的效率。
所谓的计算能力(Compute Capability),说白了就是GPU的版本号。有时也被称作SM Version。
不同版本的GPU具有不同的特性,因此程序编写也会有所差异。
计算能力为X.Y,其中主版本号X代表架构,各个架构如下表:
在CUDA的书籍文档中,我们经常能看到"1.x"、"5.x"等这样的字眼,代表第1代/第5代架构,也就是Tesla/Maxwell架构。
次版本号Y,代表在架构的基础上,有一定改进,或者有一些新特性的引入。
最新的图灵架构(Turing),实际上计算能力是7.5,也就是说还是属于Volta架构。
CUDA是软体平台,其版本(CUDA7.5 CUDA8.0 CUDA10.0)与计算能力基本没有关系。不过最新的CUDA一般会支持最新的架构。
初学者的话,上面内容能看明白的都是勇士。这里给一个例子,是我自己写的BGR转灰度图的程序,希望能让大家稍微明白一点CUDA程序如何写。
/* main.cu */ #include <iostream> #include <time.h> #include "opencv2/highgui.hpp" //实际上在/usr/include下 #include "opencv2/opencv.hpp" using namespace cv; using namespace std;
//内核函数 __global__ void rgb2grayincuda(uchar3 * const d_in, unsigned char * const d_out, uint imgheight, uint imgwidth) { const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; const unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
if (idx < imgwidth && idy < imgheight) //有的线程会跑到图像外面去,不执行即可 { uchar3 rgb = d_in[idy * imgwidth + idx]; d_out[idy * imgwidth + idx] = 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z; } }
//用于对比的CPU串列代码 void rgb2grayincpu(unsigned char * const d_in, unsigned char * const d_out, uint imgheight, uint imgwidth) { for(int i = 0; i < imgheight; i++) { for(int j = 0; j < imgwidth; j++) { d_out[i * imgwidth + j] = 0.299f * d_in[(i * imgwidth + j)*3] + 0.587f * d_in[(i * imgwidth + j)*3 + 1] + 0.114f * d_in[(i * imgwidth + j)*3 + 2]; } } }
int main(void) { Mat srcImage = imread("./test.jpg"); imshow("srcImage", srcImage); waitKey(0);
const uint imgheight = srcImage.rows; const uint imgwidth = srcImage.cols;
Mat grayImage(imgheight, imgwidth, CV_8UC1, Scalar(0));
uchar3 *d_in; //向量类型,3个uchar unsigned char *d_out;
//首先分配GPU上的内存 cudaMalloc((void**)&d_in, imgheight*imgwidth*sizeof(uchar3)); cudaMalloc((void**)&d_out, imgheight*imgwidth*sizeof(unsigned char));
//将主机端数据拷贝到GPU上 cudaMemcpy(d_in, srcImage.data, imgheight*imgwidth*sizeof(uchar3), cudaMemcpyHostToDevice);
//每个线程处理一个像素 dim3 threadsPerBlock(32, 32); dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x, (imgheight + threadsPerBlock.y - 1) / threadsPerBlock.y);
clock_t start, end; start = clock();
//启动内核 rgb2grayincuda<< <blocksPerGrid, threadsPerBlock>> >(d_in, d_out, imgheight, imgwidth);
//执行内核是一个非同步操作,因此需要同步以测量准确时间 cudaDeviceSynchronize(); end = clock();
printf("cuda exec time is %.8f ", (double)(end-start)/CLOCKS_PER_SEC);
//拷贝回来数据 cudaMemcpy(grayImage.data, d_out, imgheight*imgwidth*sizeof(unsigned char), cudaMemcpyDeviceToHost);
//释放显存 cudaFree(d_in); cudaFree(d_out);
imshow("grayImage", grayImage); waitKey(0);
return 0;
}
这里我对比了CUDA、CPU、OPENCV三种实现方式的执行时间:
速度:CUDA>OPENCV>CPU。其中OPENCV快于CPU的主要原因是OPENCV调用了并行运算指令,但慢于CUDA。
水水的介绍到此结束,剩下的我就要放飞自我了,看不懂别怪我
CUDA的编程介面由一系列C语言的扩展和运行库(runtime library)组成。
C语言的扩展在第二章「编程模型」中有所提及,如内核函数、线程网格和线程块等;
本章讲首先讲解CUDA程序的编译过程,之后会介绍CUDA运行库,最后会介绍程序兼容性等问题。
CUDA程序使用NVCC编译器。
NVCC进行离线编译的操作流程是:
PTX是一个虚拟汇编文件。其形式虽然很像汇编,但里面的每一条指令实际上是一个虚拟的指令,与机器码无法对应。需要编译器或设备驱动程序将其翻译成对应平台的汇编/机器码才能运行。
如果在编译过程中,NVCC不将设备端代码编译为cubin文件,即二进位代码,而是停在PTX代码上。设备驱动(device driver)会负责在运行时,使用PTX代码生成二进位代码。这个过程被称作在线编译(JIT Compilation, Just-In-Time Compilation)。
在线编译必然会使得程序启动的时间延长,不过设备驱动程序会自动缓存编译出来的二进位代码(也被称作compute cache)。
在线编译一方面的优势在于兼容性。另一方面的优势在于,当设备驱动程序有关编译的部分得到优化时,同样的PTX编出来的cubin文件同样会得到优化。也就是说,一段祖传的PTX代码,很有可能因为驱动程序不断的优化,而躺著得到了优化。而如果直接离线编译得到了cubin文件的话,则无法享受到这一优化。
二进位代码cubin是受到GPU计算能力的限制的。在编译时,需要使用-code来指定将代码编译到哪个计算能力平台上,如-code=sm_35代表生成的cubin代码是运行在计算能力为3.5的平台上的。
-code
-code=sm_35
二进位代码若要兼容,首先架构得一致。不同架构上的二进位代码不能互相兼容,如在Maxwell架构上编译出来的代码,不能在其他架构上运行。
另外需要说明的是,上述二进位代码的兼容性原则只限于桌面款显卡。
PTX代码的兼容性远强于二进位代码。只要不涉及到不同架构上的特性差异,PTX可以在任何架构上运行。
不过PTX代码在两种情况下其兼容性会受限:
在编译时,可以通过-arch来指定生成的PTX代码的版本,如-arch=compute_30。
-arch
-arch=compute_30
为了保证应用程序的兼容性,最好是将代码编译成PTX代码,然后依靠各个计算能力的驱动程序在线编译成对应平台的二进位代码cubin。
除了使用-arch和-code来分别指定C->PTX和PTX->cubin的计算能力外,还可以用-gencode关键字来操作,如下例:
-gencode
nvcc x.cu -gencode arch=compute_35,code=sm_35 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_60,code=compute_60,sm_60
使用上述编译指令后,会生成3.5/5.0/6.0的cubin文件,以及6.0的PTX代码。具体内容请参考nvcc user manual。
对于主机端代码,会自动编译,并在运行时决定调用哪一个版本的执行。对于上例,主机端代码会编译为:3.5/5.0/6.0的二进位文件,以及7.0的PTX文件。
另外,在程序中可以使用__CUDA_ARCH__宏来指定计算能力(只能用于修饰设备端代码)。计算能力3.5在程序中对应的__CUDA_ARCH__为350。
__CUDA_ARCH__
有一点需要注意的是,7.0以前,都是以线程束为单位在调度,线程束内指令永远是同步的。而Volta架构(计算能力7.x)引入了Independent Thread Scheduling,破坏了线程束内的隐式同步。因此,在不针对同步问题而修改代码的前提下,要想使用7.x设备,可以指定-arch=compute_60 -code=sm_70,即将PTX编到Pascal架构下以禁用Independent Thread Scheduling特性。
-arch=compute_60 -code=sm_70
另外,版本相关编译指令有缩写的情况,具体看手册。
对于主机端代码,nvcc支持C++的全部特性;而对于设备端代码,只支持C++的部分特性。具体查阅手册。
当且仅当主机端代码按照64位编译时,设备端代码才能编译为64位。当主机端代码编译为32位时,设备端代码只能编译成32位。即设备端代码的位数和主机端永远保持一致。
具体编译成32/64位的哪一种,取决于nvcc本身的版本。32位nvcc会自动编出32位的代码,不过可以使用-m64来编出64位代码。对于64位编译器亦然。
-m64
运行库实际上在cudart库内,可以使静态链接库cudart.lib/libcudart.a,或者动态链接库cudart.dll/cudart.so。
cudart.lib/libcudart.a
cudart.dll/cudart.so
所有程序的入口都是cuda。
cuda
CUDA运行库没有显式的初始化函数,在调用第一个函数时会自动初始化(设备和版本管理函数不行)。初始化时,会产生一个全局可见的设备上下文(device context)。
当主机端代码调用了cudaDeviceReset()函数,则会销毁掉这个上下文。注意,销毁的上下文是主机端正在操纵的设备。如要更换,需要使用cudaSetDevice()来进行切换。
cudaDeviceReset()
cudaSetDevice()
CUDA运行库提供了函数以分配/释放设备端的内存,以及与主机端内存传输数据。
这里的设备内存,指的是全局内存+常量内存+纹理内存。
设备内存有两种分配模式:线性存储(linear memory)、CUDA arrays。 其中CUDA arrays与纹理内存有关,本导读略去不谈。
线性内存是我们常用的内存方式,在GPU上用40位的地址线定址。线性内存可以用cudaMalloc()分配,用cudaFree()释放,用cudaMemcpy()复制数据,用cudaMemset()赋值。
cudaMalloc()
cudaFree()
cudaMemcpy()
cudaMemset()
对于2D或3D数组,可以使用cudaMallocPitch()和cudaMalloc3D()来分配内存。这两个函数会自动padding,以满足内存对齐的要求,提高内存读写效率。内存对齐的问题,会在第五章里详细阐述。
cudaMallocPitch()
cudaMalloc3D()
另外,如果要在设备内存中定义全局变数,则需要使用使用__constant__或__device__来修饰,并使用cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()来读写。如下例:
__constant__
__device__
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()
__constant__ float constData[256]; float data[256]; cudaMemcpyToSymbol(constData, data, sizeof(data)); cudaMemcpyFromSymbol(data, constData, sizeof(data));
__device__ float devData; float value = 3.14f; cudaMemcpyToSymbol(devData, &value, sizeof(float));
__device__ float* devPointer; float* ptr; cudaMalloc(&ptr, 256 * sizeof(float)); cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));
实际上,当使用__constant__关键字时,是申请了一块常量内存;而使用__device__时,是普通的全局内存。因此__device__申请的内存需要申请,而__constant__不用。不管是全局内存,还是常量内存,都是只读的,需要用带有Symbol的函数拷贝。
Symbol
不管是全局变数还是局部变数,都需要使用__shared__来修饰。不过需要注意的是,即使定义为全局变数,共享内存依旧只能被同一线程块内的线程可见。
__shared__
举个例子,对于如下代码,虽然是定义了一个全局的共享内存hist_shared,但实际上,在每一个线程块被调度到SM上时,都会在SM的共享内存区开一块内存。因此,每一个线程块都有一个hist_shared,且之间无法互相访问。
__shared__ unsigned int hist_shared[256]; //共享内存仅在线程块内共享
__global__ void getGrayHistincuda_usesharemem(unsigned char * const grayData, unsigned int * const hist, uint imgheight, uint imgwidth) //使用共享内存加速 { const unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x; const unsigned int idy = blockDim.y * blockIdx.y + threadIdx.y; const unsigned char inner_idx = threadIdx.y * blockDim.x + threadIdx.x;
hist_shared[inner_idx%256] = 0; //清空数据,由于每个块的inner_idx可以超过256,所以这样可以保证hist_shared被全部清零
__syncthreads(); //等待其他线程完成
if(idx < imgwidth && idy < imgheight) { const unsigned long pid = imgwidth * idy + idx; const unsigned char value = grayData[pid]; atomicAdd(&(hist_shared[value]), 1); }
__syncthreads();
if(threadIdx.y < 8) //每个线程块将自己共享内存中的值合并到全局内存中去 { atomicAdd(&(hist[inner_idx]), hist_shared[inner_idx]); }
当然,共享内存的声明放在内核函数里面也是可以的,效果一致。
使用共享内存,可以获得等同于L1 cache的访存速度,其速度远快于全局内存。
但是注意,并不是什么时候都可以使用共享内存来获取加速的。例如内核函数计算出来结果后,如果这个结果只需要传输回主机端,而不需要再次被用到时,直接写回全局内存会比较快。如果先写回共享内存,再写回全局内存,反而会比较缓慢。
强调一下,共享内存只能为线程块内的线程共享。如果需要整个线程网格中线程都能访问,则需要全局内存或常量内存。
另外,共享内存是一个稀缺资源。有些架构可以通过配置,分配L1 cache和共享内存的比例。
锁页内存指的是主机端上不会被换出到虚拟内存(位于硬碟)上的内存。
锁页内存的分配与释放:
cudaHostAlloc()
cudaFreeHost()
cudaHostRegister()
malloc()
NVIDIA官方给出的锁页内存相对于普通的内存的的好处是:
另一本书对于锁页内存之所以快的解释是:
注意,锁页内存在 non I/O coherent Tegra 设备上不支持
NVIDIA官方文档表示:上述所说的锁页内存的优点,只有在使用cudaHostAlloc()时,传入cudaHostAllocPortable flag,或者在使用cudaHostRegister()时传入cudaHostRegisterPortable flag,才能体现。否则锁页内存并不会有上述优点。
cudaHostAllocPortable
cudaHostRegisterPortable
《GPU编程指南》一书中是这么描述的:如果传入了cudaHostAllocPortable flag,则锁页内存在所有的CUDA上下文中变成锁页的和可见的。如果需要在CUDA上下文之间或者主机处理器的线程之间传递指针,则必须使用这个标志。
(好吧,从编程指南一书中确实没出来用Portable的必要性,不是很明白)
锁页内存默认是使用缓存的。如果将flag cudaHostAllocWriteCombined 传入到 cudaHostAlloc(),则可以将这块锁页内存指定为合并写内存。
cudaHostAllocWriteCombined
合并写内存不再使用主机端的L1&L2 cache,使得更多的cache可以供其他任务使用。
另外,对于通过PCI-E传输数据的情景,使用合并写内存不会被snooped (是不是指的是不会被缓存管?不理解这个snooped什么意思),可以提升40%的传输性能。
此外需要注意的是,由于合并写内存不使用缓存,因此读入CPU核的操作会非常的慢。因此合并写内存最好只用作向GPU传数据的内存,而不是传回数据的内存。
CUDA中的内存映射,指的是将CPU端的锁页内存,映射到GPU端。
通过向cudaHostAlloc()传入cudaHostAllocMapped flag,或向cudaHostRegister()传入cudaHostAllocMapped flag,来将一块内存指定为向GPU映射的内存。
cudaHostAllocMapped
映射的内存有两个地址,一个是CPU端访问的地址,一个是GPU端访问的地址。
cudaHostGetDevicePointer()
使用内存映射有以下好处:
使用内存映射必须要注意的几点:
cudaSetDeviceFlags()
cudaDeviceMapHost
canMapHostMemory
CUDA允许以下操作互相并行:
设备端的如下操作,可以跟主机端并行:
CUDA_LAUNCH_BLOCKING
Async
其中第3、4条说明,在使用cudaMemcpy()时,如果数据小于等于64KB,其实传输相对于CPU是非同步的。 如果数据多于64KB,则CPU会阻塞到数据传输完成。 这时使用带Async的内存传输函数,会释放CPU资源。
需要注意的是,如果没有使用锁页内存,即使使用了Async函数,内存传输也不是并行的(和CPU?还是GPU?)。
计算能力2.x及以上的设备,支持多个内核函数同时执行。(可以通过检查concurrentKernels来确定)
concurrentKernels
执行多个内核函数,需要主机端不同的线程启动。如果一个线程依次启动多个内核,则这些内核会串列执行。同一线程的内核函数返回时会触发隐式的同步。
另外,多个内核函数必须位于同一个CUDA上下文(CUDA context)上。不同CUDA上下文上的内核不能并行。这意味著,启动多个内核的多个线程必须使用相同的CUDA上下文。(如何传递CUDA上下文?)
一些设备支持数据传输(主机端/设备端、设备端/设备端)和内核执行并行,可通过检查asyncEngineCount来确认。
asyncEngineCount
一些设备支持设备端内部数据传输和内核执行/数据传输并行,可通过检查concurrentKernels来确认。
这一特性需要使用锁页内存。
计算能力2.x及以上的设备,支持数据传入和传出并行。
必须使用锁页内存。
在CUDA中,流(streams)指的是在GPU上一连串执行的命令。
不同的线程,可以向同一个流填入任务。
同一个流内的任务会按顺序执行;同一设备上不同的流有可能并行,其执行顺序不会有保证。
下述代码是一个流的创建和销毁的例子。该程序创建了两个流,分配了两个锁页内存传输数据,依次启动了两个内核,最后销毁了这两个流。
cudaStream_t stream[2]; for (int i = 0; i < 2; ++i) cudaStreamCreate(&stream[i]); float* hostPtr; cudaMallocHost(&hostPtr, 2 * size);
for (int i = 0; i < 2; ++i) { cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]); MyKernel <<<100, 512, 0, stream[i]>>> (outputDevPtr + i * size, inputDevPtr + i * size, size); cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]); }
for (int i = 0; i < 2; ++i) cudaStreamDestroy(stream[i]);
从上例中可以看到,流的创建需要定义cudaStream_t结构,并调用cudaStreamCreate()来初始化。
cudaStream_t
cudaStreamCreate()
cudaStreamDestroy()
当向流中添加内核函数任务时,<<<...>>>不再是<<<blocksPerGrid, threadsPerBlock>>>,而是<<<blocksPerGrid, threadsPerBlock, dynamic_shared_memory, stream>>>。
<<<...>>>
<<<blocksPerGrid, threadsPerBlock>>>
<<<blocksPerGrid, threadsPerBlock, dynamic_shared_memory, stream>>>
当设备还在执行流中的任务,而用户调用cudaStreamDestroy()函数时,函数会立刻执行(不会阻塞)。之后,当流中的任务完成后,与流相关的资源会自动释放。
另外需要注意的是,上例中主机端线程、数据拷贝和内核执行完全非同步,因此在"拷贝回主机端"这一操作完成之前,主机端的内存数据是不正确的。必须在数据返回的一步做同步操作,方能保证数据是正确的。
(需要了解一下流内部是如何实现的,为什么内核执行和内存拷贝能够非同步且重叠?什么样的操作又不能重叠?)
在调用内核函数时,不指定流或者将流指定为0,则代表使用了默认流(default stream)。
如果在编译时使用了--default-stream per-thread,或是在include任何cuda头文件前#define CUDA_API_PER_THREAD_DEFAULT_STREAM,则主机端的每一个线程都有自己专属的默认流。
--default-stream per-thread
#define CUDA_API_PER_THREAD_DEFAULT_STREAM
--default-stream legacy
可以使用如下函数进行显式同步:
cudaDeviceSynchronize()
cudaStreamSynchronize()
cudaStreamWaitEvent()
cudaStreamQuery()
注意,同步函数慎用,因为有可能会产生速度的下降。
一般来讲,不同流内的命令可以并行。但是当任何一个流执行如下的命令时,情况例外,不能并行:
操作的重叠程度,一方面取决于各个操作的顺序,另一方面取决于设备支持重叠的程度(是否支持内核执行并行/数据传输与内核执行并行/数据传输并行)
可以使用cudaStreamAddCallback()函数,向流中添加callback。该callback会在流中之前所有的任务完成后被调用。如果stream参数设为0,则代表之前的所有stream的任务执行完后就调用该callback。
cudaStreamAddCallback()
回调函数和cudaStreamWaitEvent()一样,对于在加在callback之后的指令,必须等待callback执行完成后,才会继续执行。
下例是一个使用回调的例子。该例中,两个stream将数据拷回主机端后,会调用回调函数。
void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){ printf("Inside callback %d ", (size_t)data); } ... for (size_t i = 0; i < 2; ++i) { cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]); MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size); cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]); cudaStreamAddCallback(stream[i], MyCallback, (void*)i, 0); }
回调函数中不能直接或间接的执行CUDA函数,否则会因为等待自己完成而造成死锁。 (原因尚不太明白)
可以通过cudaStreamCreateWithPriority()来在创建流时指定流的优先顺序。可以指定的优先顺序可由cudaDeviceGetStreamPriorityRange()来获得。
cudaStreamCreateWithPriority()
cudaDeviceGetStreamPriorityRange()
运行时,高优先顺序stream中的线程块不能打断正在执行的低优先顺序stream的线程块(即不是抢占式的)。但是当低优先顺序stream的线程块退出SM时,高优先顺序stream中的线程块会被优先调度进SM。
事件(Event)可以被压入流中以监视流的运行情况,或者用于精确计时。
如果向stream 0压入事件,则当压入事件前向所有流压入的任务完成后,事件才被触发。
cudaEvent_t start, stop; //创建 cudaEventCreate(&start); cudaEventCreate(&stop); ... cudaEventDestroy(start); //销毁 cudaEventDestroy(stop);
下例是一个使用Event计算时间的例子:
cudaEventRecord(start, 0); //记录事件(将事件压入流),流0则代表所有流完成任务后事件才会被触发 for (int i = 0; i < 2; ++i) { cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]); MyKernel<<<100, 512, 0, stream[i]>>>(outputDev + i * size, inputDev + i * size, size); cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]); } cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop); //获取两个事件发生的时间差(ms)
下例是如何枚举设备,并获取设备信息的例子:
int deviceCount; cudaGetDeviceCount(&deviceCount); //获取设备数量 int device; for (device = 0; device < deviceCount; ++device) { cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, device); printf("Device %d has compute capability %d.%d. ", device, deviceProp.major, deviceProp.minor); }
使用cudaSetDevice()选择设备,当不选择时,默认使用设备0。
注意,所有的内存分配、内核函数启动、流和事件的创建等,都是针对当前选择的设备的。
下例是一个设备选择的例子:
size_t size = 1024 * sizeof(float); cudaSetDevice(0); // Set device 0 as current float* p0; cudaMalloc(&p0, size); // Allocate memory on device 0 MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0 cudaSetDevice(1); // Set device 1 as current float* p1; cudaMalloc(&p1, size); // Allocate memory on device 1 MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
下面将讨论,如果对一个不属于当前设备的流或事件进行操作,哪些操作会成功,哪些操作会失败:
cudaSetDevice(0); // Set device 0 as current cudaStream_t s0; cudaStreamCreate(&s0); // Create stream s0 on device 0 MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0 cudaSetDevice(1); // Set device 1 as current cudaStream_t s1; cudaStreamCreate(&s1); // Create stream s1 on device 1 MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1
// This kernel launch will fail: MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0
另外需要注意,每个设备都有自己的默认流。因此在没有指定流的情况下,向不同设备分派的任务,实际上是压入了各个设备的默认流,他们之间是并行执行的。
计算能力2.0及以上的设备支持设备间对等内存访问,这意味著两个GPU之间的传输和访问可以不经过主机端中转,速度会有提升。查询cudaDeviceCanAccessPeer()可以得知设备是否支持这一特性。(官方文档说还需要一个条件:64位程序,存疑)
cudaDeviceCanAccessPeer()
需要使用cudaDeviceEnablePeerAccess()来使能这一特性。
cudaDeviceEnablePeerAccess()
对等设备的的地址是统一编址的,可用同一个指针访问,如下例:
cudaSetDevice(0); // Set device 0 as current float* p0; size_t size = 1024 * sizeof(float); cudaMalloc(&p0, size); // Allocate memory on device 0 MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0 cudaSetDevice(1); // Set device 1 as current cudaDeviceEnablePeerAccess(0, 0); // Enable peer-to-peer access with device 0
// Launch kernel on device 1 // This kernel launch can access memory on device 0 at address p0 MyKernel<<<1000, 128>>>(p0);
对等设备的地址是统一编址的,可以使用cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer、cudaMemcpy3DPeerAsync()来进行直接拷贝。无需先拷贝会主机端内存,再转到另一块卡上。如下例:
cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer、cudaMemcpy3DPeerAsync()
cudaSetDevice(0); // Set device 0 as current float* p0; size_t size = 1024 * sizeof(float); cudaMalloc(&p0, size); // Allocate memory on device 0 cudaSetDevice(1); float* p1; cudaMalloc(&p1, size); // Allocate memory on device 1 cudaSetDevice(0); // Set Device 0 as Current MyKernel<<<1000, 128>>>(p0); // Launch Kernel on Device 0 cudaSetDevice(1); // Set Device 1 as Current cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1 MyKernel<<<1000, 128>>>(p1); // Launch Kernel on Device 1
关于设备间的对等拷贝,如果使用的是NULL stream,则有如下性质:
(使用的如果不是NULL Stream,又会怎样呢?)
当程序是64位程序时,所有主机端内存,以及计算能力≥2.0的设备的内存是统一编址的。所有通过CUDA API分配的主机内存和设备内存,都在统一编址的范围内,有自己的虚拟地址。因此:
cudaPointerGetAttributes()
cudaMemcpy***()
cudaMemcpyKind
cudaMemcpyDefault
可以通过查询unifiedAddressing来查看设备是否支持统一虚拟编址。
unifiedAddressing
线程间通讯,可以很方便的通过共享的变数来实现。然而进程间通讯不行。
为了在进程间共享设备端内存的指针或者事件,必须使用IPC(Inter Process Communication) API。IPC API只支持64位程序,并且要求设备计算能力≥2.0。
通过IPC中的cudaIpcGetMemHandle(),可以得到设备内存指针的IPC句柄。该句柄可以通过标准的IPC机制(interprocess shared memory or files)传递到另一个进程,再使用cudaIpcOpenMemHandle()解码得到该进程可以使用的设备内存指针。
cudaIpcGetMemHandle()
cudaIpcOpenMemHandle()
所有的runtime function都会返回一个error code,可通过检查error code判断是否出错。
但是对于非同步函数,由于在执行前就会返回,因此返回的error code仅仅代表函数启动时的错误(如参数校验);非同步函数不会返回运行时出现的错误。如果运行时出了错,会被后面的某个函数捕获并返回。
检查非同步函数是否出错的唯一方式,就是在非同步函数启动后,进行同步。 如在非同步函数后,调用cudaDeviceSynchronize(),则非同步函数的错误会被cudaDeviceSynchronize()捕获到。
事实上,除了runtime function会返回error code之外,每一个主机端线程都会有一个初始化为cudaSuccess的变数,用于指示错误。一旦发生了错误,该变数也会被设置为相应的error code。
cudaSuccess
该变数不会被直接调用,但可以被cudaPeekAtLastError()和cudaGetLastError()访问到。不同的是,cudaGetLastError()在返回这一变数的同时,会把它重置为cudaSuccess。
cudaPeekAtLastError()
cudaGetLastError()
内核函数不会返回值,因此只能通过cudaPeekAtLastError()或cudaGetLastError()来知悉调用内核是否有错误。
另外需要注意的是,cudaStreamQuery()和cudaEventQuery()这类函数,有可能会返回cudaErrorNotReady。但这不被认为是错误,因此不会被cudaPeekAtLastError()和cudaGetLastError()捕获到。
cudaEventQuery()
cudaErrorNotReady
对于计算能力≥2.0的设备,可以通过cudaDeviceGetLimit()/cudaDeviceSetLimit()来查询/设置调用栈的大小。
cudaDeviceGetLimit()
cudaDeviceSetLimit()
有两个版本需要注意:计算能力,以及CUDA driver API的版本。其中计算能力及其兼容性在前面已有阐述。
CUDA driver API的版本定义在驱动的头文件中的CUDA_VERSION宏内。可以在程序中将该宏调出,以检查程序是否可以在目标设备上运行。
CUDA_VERSION
CUDA driver API不是向前兼容的。也就是说,针对新版本的CUDA driver API编译的程序、插件、库,并不能在旧版本的驱动上运行。
关于CUDA driver API,有几点需要注意:
NVIDIA的设备可以设置三种计算模式:
正常情况下,如果程序没有调用cudaSetDevice(),则会默认使用0号设备。但是如果0号设备被置成禁止模式,亦或是被其他进程所专属,则会在其他设备上创建上下文并使用。 可以向cudaSetValidDevices()函数输入一个设备列表,函数会在第一个可以使用的设备上创建上下文。
cudaSetValidDevices()
Pascal及以上架构,支持指令级的优先顺序调度。不再是以线程块为SM的最小调度单位,而是以指令为最小调度单位,且具有优先顺序。 这意味著具有冗长kernel的线程块不再会占据太多的计算资源,或是发生timeout。但是这也有缺点:当多个进程创建了上下文时,以往基于线程块的调度不会造成太多的上下文切换,但现在的指令级调度则会造成很多的上下文切换,降低效率。(注意跟GPU内线程的上下文切换不同,GPU内线程上下文切换几乎不浪费时间,直接换一个指针就好)。因此最好设置为Exclusive-process,这样只有一个进程会使用设备。(如果线程很多的话,效果不是一样吗?)
(讲道理好好看看上下文是什么鬼)
(这段需要再看一下Pascal架构的说明,看看指令级的调度是如何实现的,寄存器等又是如何分配的)
设备处于哪种计算模式,可通过检查computeMode来查看。
computeMode
GPU会将一些内存专门分配给primary surface,用于刷新显示设备。 一旦用户将显示模式切换,如增加解析度或增加彩色的位数,会造成primary surface所需的内存变多。此时系统会把原来分配给CUDA运算的内存,调拨给primary surface,从而造成CUDA runtime产生错误,并返回invalid context error。
(言外之意是说,跑cuda的时候不要切解析度?)
这份官方文档讲的硬体内容太少了,从另一本书里补一点过来,可能内容有点老旧,见谅。
(在一本比较老旧的书上找到的结构图)
下图是一款比较老旧的显卡(G80/GT200)的硬体结构图:
从图中可以看到,GPU由内存和一系列流处理器簇(Streaming Multiprocessors, SM)组成,不同GPU的具有不同的内存大小和SM数量。
流处理器簇(SM)内部结构如下图所示:
SM由一系列流处理器(Streaming Processor,SP)、寄存器文件(Register File)、共享内存(Shared Memory, SMem)、SPU(特殊运算单元),以及纹理/常量/L1缓存组成。L2 cache是由所有SM共享的。
线程块调度
另外需要强调一点的是,如果线程块需要的寄存器或共享内存太多,以至于SM连一个线程块都无法满足的时候,内核会启动失败。
线程束调度
当调度器选择调度线程束0时,第一个时钟周期会将半个线程束调度到第一个16xSP上,然后下一个周期会将另外半个线程束调度到另一个16xSP上。因此,调度一个完整的线程束0的指令,需要2个时钟周期。
线程束有时会处于等待状态(等待内存,或等待特殊计算的结果)。此时调度器会调度其它线程束到SP上运行。当等待的线程束等到了所需的元素后,会再次处于就位状态,等待调度器调度运行。
不同计算能力的平台,具有不同的调度器数量和SP数量。
NVIDIA GPU是典型的SIMT架构(Single-Instruction, Multiple-Thread Architecture,单指令多线程架构)。(半)线程束在同一时间内执行同样的指令(相同的PC),但每个线程有自己的数据空间(寄存器);可以同时做同样的事情,但是处理不同的数据。这样可以有效节省指令带宽(指令由线程束共享),实现高效并行。
当线程束中的线程,必须要执行不同的条件分支时,满足分支条件的线程会被激活并执行分支内的内容;不满足分支条件的线程会接收同样的指令,但不会被激活,不会实际执行,但也不能跳过去执行其他指令。换言之,当线程束中的线程遇到分支时,不论线程是否需要执行分支,都会消耗执行该分支的时间,因为线程束中的线程执行同样的指令,是高度同步的。
... if(条件) { 操作1 } else { 操作2 } ...
如果是CPU,如果不满足条件,则会直接跳转去执行操作2,不会执行操作1。 但是对于GPU同一个线程束内的线程,即使不满足条件,依旧会去跟其他线程一起去执行操作1。等到操作1执行完毕后,该线程又会跟其他线程一起执行操作2。这也就意味著,不管线程实际上需不需要执行某一分支,它都要跟其他线程一起跑完这一分支。
另外,GPU不支持分支预测和推测执行,只会老老实实的一条条执行各个分支里的内容。因此,对于CUDA程序来说,分支实际上是一种低效的行为。
但是有两种情况是例外:
但是从Volta架构开始,Independent Thread Scheduling被引入,线程束内的线程不再完全同步。每个线程都会有自己独立的PC。遇到分支时,不再像之前的架构一样,只有(半)线程束内的线程条件一致时,才会跳过分支;Volta架构的调度优化器会将线程束中的线程,按照分支条件是否满足,重新组合成SIMT单元,从而跳过分支。
Volta架构的Independent Thread Scheduling无疑是高效的,但是这是一个跟旧架构完全不同的特性。在编写旧架构的CUDA程序时,程序员会默认线程束内的线程一定会同步执行。Volta架构的新特性破坏了这一假设,无疑会给程序带来一些问题,需要注意。
在一个线程束中,参与当前指令的线程称为active threads,不参与的被称为inactive threads。造成线程inactive的原因有:
当线程束中所有线程,向同一个内存地址写数据时,不能保证哪个线程先写,哪个线程后写,即程序的正确性是无法保证的。 这时需要使用CUDA提供的原子操作(atomic)函数,如atomicAdd()。
atomicAdd()
原子操作可以保证程序的正确性,但是会造成线程束中线程的串列化(serialization),执行时间比并行执行要长。
另外需要注意的是,即使没使用原子操作,向同一个内存地址写数据,一样也会产生一定程度的串列化,串列化程度依架构而定。
GPU线程束中的线程的上下文(寄存器&PC等)都存在片内空间庞大的寄存器文件中,直到线程束执行完毕才会被释放(生命周期为整个线程束执行过程)。因此,不像CPU,GPU的线程束上下文切换十分迅速,没有损耗。 线程束调度器(warp scheduler)会选择ready状态的线程束,将其调度到SP上执行。
CUDA程序性能优化有三个原则:
在性能优化前,需要先分析程序性能的瓶颈,再针对瓶颈优化,否则收益会很低。
最大化利用率的方法就是并行。
从程序最高层来看,应该尽可能让主机端、设备端、PCI-E汇流排并行工作。对此可以使用非同步CUDA函数,以及流(Stream)来实现。
同步操作,以及内存的共享会影响程序的并行性。因此需要仔细设计演算法流程,尽量减少同步和内存共享。 如果一定需要同步和内存共享,尽量在线程块内完成(线程块同步——使用__syncthreads()涉及到的线程少,且可以通过SM内的共享内存共享数据。如果需要线程网格内同步,则需要两个内核调用,且共享数据只能通过全局内存,速度慢)。
__syncthreads()
可以通过流的方式,尽可能的让多个内核并行,提升利用率。
延迟(latency)指的是线程束(从上一个动作开始)到它处于ready状态的时钟数。 例如线程束先提交了一个内存访问请求,然后等了400个时钟周期,内存管理系统才返回数据,线程束可以继续执行。这400个时钟周期称为延迟。
当一个线程束发生延迟时,线程束调度器(warp scheduler)会将其他处于ready状态的线程束调度到SP上。等到延迟结束后,再将该线程调度回SP继续执行。这样一来,前一个线程束的延迟,就被另一个线程束的执行所隐藏了。 这一过程被称作延迟的隐藏(hidden latency)。
隐藏延迟是GPU编程的核心概念。由于GPU具有巨大的寄存器空间,线程的切换不存在损耗。因此,通过向GPU上分配足够多的线程,可以让这些线程延迟互相交错,以起到隐藏延迟的作用,提高硬体利用率。
造成线程(束)产生延迟的原因有:
通过配置线程网格、线程块、寄存器和共享内存用量,让SM可以运行尽可能多的线程束,以隐藏延迟。例如对于计算能力3.x的设备,为了完全隐藏全局内存读取的延迟(200-400时钟),需要大概40个线程束。
举个例子,设SM有32KB共享内存空间。程序每个线程需要32B共享内存,即一个线程束需要1KB共享内存,考虑下述两种方案:
虽然方案2在一个线程块上,有更多的线程束,但是实际上SM上运行的线程束减少了(32->18)。因此方案2隐藏延迟的能力弱于方案1,资源利用率较低。
此外,如果寄存器使用过多,超过了SM上的寄存器空间,则会使用本地内存作为寄存器。本地内存是存在在全局内存上的,速度很慢,会严重影响程序速度。因此需要严格考虑寄存器使用数量。 (这里官方文档和另一本书里说的有矛盾,难道是新的架构把本地内存给取消掉了?)
最后强调一点,线程块中的线程数量,最好是32的整数倍。这样,就不会有为了补齐线程束,而出现的永远不会激活的线程。这些不激活的线程也会占用SM的资源,降低资源利用率。
CUDA具有Occupancy Calculator,帮助程序员设计。
最大化内存吞吐,主要手段就是少用低带宽的内存。这意味著首先要尽可能减少主机端和设备端间的设备传输(PCI-E,特别慢),其次要尽可能减少全局内存的读写(快于PCI-E,但是相对于片内内存来说,还是挺慢的);尽可能的使用片内的内存(寄存器、cache、共享内存)。
这里需要强调一下cache和共享内存的事情。
共享内存是程序可控的高速缓存。一般情况下,共享内存的使用流程为: 将数据从全局内存拷贝到共享内存,或初始化共享内存 进行一个同步操作,确保共享内存全部被赋值 利用共享内存的数据,运行程序 如果出现了共享内存的写操作,一般需要进行一个同步操作,确保写操作全部完成后再进行下面的操作 * 将数据写回全局内存
这里有一点要强调,只有在数据需要反复读写的时候,共享内存才有意义。如果数据只会被读一次,处理完后又写回并不再处理。则直接从全局内存读出->寄存器运行->写回全局内存是最快的。在共享内存中转反而是慢的。
缓存(L1/L2 cache)是程序员无法显式编程的。但是如果了解缓存的特性的话,可以通过合适的程序设计,增加缓存命中率。
由于PCI-E传输并不快,因此要尽量减少主机端和设备端间的数据传输: 一种方式是让中间结果尽可能的在设备端产生,在设备端使用。 另一种方式是将很多小的数据,打包传输。 还有可以通过分配锁页内存来加快前端汇流排*系统的带宽。
当使用内存映射时,需要注意,每次内存访问都会启动一次PCI-E传输。因此,尽量保证数据只被读写一次,且尽可能合并访问以提升有效内存带宽。
有些GPU设备,主机端和设备端内存,在物理上就是同一块。这种情况下,主机端和设备端传输是不存在的。可通过标志integrated来查看。
integrated
全局内存支持合并访问,可以一次性传输连续的32、 64、 128位元组的数据。因此,在设计内核时,线程束内的线程尽量连续的访问内存。
考虑如下两个内核:
//假设gpuData是一个二维数组,尺寸为32x32 int gpuData[32][32]; //这样是不合法的,因为这么定义实际上是在主机端,还需要拷贝到设备端,这里只是为了方便说明问题
__global__ void Kernel1(int gpuData[][32]) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; int sum = 0; for(int i = 0; i < 32; i++) sum += gpuData[i][tid]; //行访问 ... }
__global__ void Kernel2(int gpu[][32]) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; int sum = 0; for(int i = 0; i < 32; i++) sum += gpuData[tid][i]; //列访问 ... }
上例中,执行Kernel1的线程束中的线程,在一次循环中,32个线程依次访问gpuData[0][0], gpuData[0][1], gpuData[0][2], ..., gpuData[0][31]。在内存中,这32个变数是连续存储的,因此可以被合并访问。这种访问被称为行访问。
gpuData[0][0], gpuData[0][1], gpuData[0][2], ..., gpuData[0][31]
gpuData[0][0], gpuData[1][0], gpuData[2][0], ..., gpuData[31][0]
上例中,列访问之所以效率低,原因有二:
因此,从上例中可以看到,好好安排内存排布,尽量使得内存访问可以合并,可以加速全局内存的读写。
当变数的尺寸为1/2/4/8/16位元组时,变数会对齐。但如果不是的话,变数无法对齐,会产生额外的内存访问。
C/C++内建的变数(int/float等),以及CUDA支持的向量(float2/float4等),是对齐的。
一些结构体可能会产生不对齐的情况,看下例:
struct struct1{ float x; float y; };
struct struct2{ float x; float y; float z; };
struct struct3 __align__(16){ float x; float y; float z; };
上例中,struct1是8位元组的结构体,自动会对齐; struct2具有12个位元组,无法对齐; struct3使用了__align__(16)关键字,显式指定对齐到16。
__align__(16)
使用各类malloc分配的设备内存,一定是256位元组对齐的。
当使用了自动变数(不明白,再看)时,有可能会将数据放到本地内存上:
通过看PTX代码,可以看到标记为.local的变数,就是本地内存。
.local
lmem
前面多次强调过了,一旦使用了本地内存,其速度会非常慢。不过本地内存在存储的时候,是按照32个线程连续存储的,因此可以合并访问。
共享内存实际上是被分为多个存储体(memory bank)。多个线程访问同一个存储体会造成串列化。
因此,编写内核时,需要认真设计,以避免存储体访问的冲突。
可以使用如下方法来最大化指令吞吐:
__restrict__
指令吞吐的定义:每个SP在每个时钟周期内执行的操作数。如果一个线程束在一个时钟周期内执行了N个操作,则指令吞吐为N/32。
官方文档这里比较混乱,但主要有如下几点: 不同架构的设备,不同指令有不同的指令吞吐,可以查表 有一些快速的内联(inline)函数,如使用__fdividef()(快速浮点数除法)来代替普通的除法来加速 整形的除法和取余会比较慢,可能需要20个机器周期;因此对于n为2的幂次的情况,使用i>>log2(n)代替i/n,使用i&(n-1)来代替i%n 半精度(浮点数)运算(Half Precision Arithmetic):可以使用half2数据类型,并使用对应的运算指令(如__hadd2, __hsub2, __hmul2, __hfma2等),来让一个周期内执行两次运算,以节省指令带宽。可以通过__halves2half2将两个半精度浮点数合并为half2数据类型。 (半精度又是咋定义的?) * 数据类型转换:当使用char或short,亦或是双精度常量与单精度变数相互操作时,会触发数据类型转换,需要一定执行时间(实际上,char和short,不管是存储在寄存器中,还是在运算时,都是以int型进行的)
__fdividef()
i>>log2(n)
i/n
i&(n-1)
i%n
half2
__hadd2, __hsub2, __hmul2, __hfma2
__halves2half2
尽量避免向线程束中引入分支。
此外,可以使用#pragma unroll宏,来进行循环展开,减少控制指令。
#pragma unroll
下表为不同计算能力的设备,同步指令__syncthreads()需要消耗的指令周期为:
__syncthreads()消耗的指令周期
注意,__syncthreads()会造成线程块中的线程等待,影响内核执行效率。