1.cu文件一般写cuda的核函数,用来加速比如图像的前后处理,核函数的代码是在gpu上运行的,核函数里的打印是在gpu上做打印,但它打印的结果会回收到cpu上展示出来。核函数调用等于启动了很多个线程,每个线程独自去运行这一部分的代码
2.在.vscode/settings.json中配置*.cu : cuda-cpp,可以使得代码被正确解析
3.Makefile中,.cu文件交给nvcc进行编译,nvcc是一个编译器,用来编译cuda的c程序的
4.cu文件可以当做正常cpp写即可,他是cpp的超集,也就是说它涵盖了c++的所有特性,在c++的所有特性之上又增加了一些新的特性。
5.cu文件中引入了一些新的符号和语法:
5.1_global__标记,核函数标记
1)核函数的调用方必须是host,主机即cpu。然后就是如果想在核函数里调用其他函数,这个其他函数也必须是设备函数(带__ device__标记)。但比如像printf和exp这些函数也像host函数,那为什么在gpu里可以被正常调用呢?因为这些常规函数在gpu上都有对应的实现的即相对于的device版本,NVIDIA已经帮你封装好了。
2)返回值必须是void
3)例如__global__ void kernel(const float* pdata, int ndata)
4)必须以
kernel<<<gridDim, blockDim, bytesSharedMemorySize, stream>>>(pdata, ndata)
的方式启动,只有__global__ 修饰的函数才可以用<<<>>>的方式调用
其参数类型是:
<<<dim3 gridDim, dim3 blockDim, size_t bytesSharedMemorySize, cudaStream_t stream>>>
dim3这个类有构造函数dim3(int x, int y=1, int z=1),因此当直接赋值为int时,实则定义了dim.x = value, dim.y = 1, dim.z = 1,可以看到默认为1,value就是参数传进来的那个值。bytesSharedMemorySize参数的单位为字节。
其中gridDim, blockDim是线程layout(线程布局)参数(下文第8点中有详细解释),stream参数表示流,要做异步时用stream来控制它,bytesSharedMemory参数表示共享内存的大小,gridDim, blockDim用来告诉cuda需要启动多少个线程,可以理解为启动了gridDim个block(具体以下面代码为准!),每个block有blockDim个线程。为什么要搞这么复杂用gridDim和blockDim去描述线程数呢,是因为后期考虑到性能,通常遇到的问题都不是一维的而是多维的。因此有了gridDim和blockDim去描述能更加方便去使用。可以理解为最终目的是让每一个线程去处理一个维度的数据!!!比如在图像处理加速中就是图像分辨率有多大,启动的线程就有多少个。 gridDim和blockDim的类型都是dim3,dim3是一个结构体,启动的线程总数量nthreads就等于:
dim3 gridDim;
dim3 blockDim;
int nthreads = gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z
gridDim和blockDim的上限是可以通过runtime API或者deviceQuery可以查询到。gridDim的x,y,z分别的上限值为括号的(21亿,65536,65536),blockDim的x,y,z分别的上限值为(1024,64,64)。还有个约束就是blockDim.x * blockDim.y * blockDim.z <=1024。
如果指定了stream参数,则把核函数加入到stream中异步执行(异步就是执行了立马返回,不用等操作是否进行完毕)
pdata和ndata则是核函数的函数调用参数
函数调用参数必须传值,不能传引用等。参数可以是类类型等
5 核函数的执行无论stream是否为nullptr(默认流),都将是异步执行。 因此在核函数中进行printf操作,你必须进行等待,加一个设备同步或者流同步就行了,例如cudaDeviceSynchronize、或者cudaStreamSynchronize,否则你将无法看到打印的信息
5.2 __ device__标记,设备调用的函数,函数前加个__device__,把这个就看做一个标识符就行了,没什么别的用处
调用方必须是device
5.3 __host__标记,主机调用函数
调用方必须是主机
5.4 也可以__device__ __host__两个标记同时有,表明该函数可以设备也可以主机
5.5 __constant__标记,定义常量内存
5.6 __shared__标记,定义共享内存
6.通过cudaPeekAtLastError/cudaGetLastError函数,捕获核函数是否出现错误或者异常
7.内存索引(也叫线程索引(线程的绝对索引),即在线程总数里当前执行的这段代码是属于第几个线程)的计算公式
定位一个thread到底是多少的时候,通常会映射到指针的偏移量上,这个偏移量就是线程索引。计算方式就是这么做的,方便记忆的办法,是左乘右加,无论tensor维度多复杂6维7维100维,这个方法都适用,核函数里把blockDim和gridDim看做shape,把threadIdx和blockIdx看做index,下面代码中的dims和indexs如下图所示:
//把内存地址看做长条数组,则通用的内存定位计算如下
position = 0 //计算线程(内存)索引时,初始地址都看做0
for i in range(6):
position *= dims[i]
position += indexs[i]
8.buildin变量,即内置变量,通过ctrl+鼠标左键点进去查看定义位置(把blockDim和gridDim看做shape,把threadIdx和blockIdx看做shape对应的索引index)
1)所有核函数都可以访问,其取值由执行器维护和改变
2)gridDim[x, y, z]:网格维度,线程布局的大小,是核函数启动时指定的,gridDim对应的索引是blockIdx,比如gridDim的shape是3乘3,那么blockIdx范围就是0,1,2。即gridDim是多少,那么blockIdx就是在它范围之内的。
3)blockDim[x, y, z]:块维度,线程布局的大小,是核函数启动时指定的
4)blockIdx[x, y, z]:块索引,blockIdx的最大值是0到gridDim-1,由执行器根据当前执行的线程进行赋值,核函数内访问时已经被配置好,blockIdx一定是小于gridDim的,blockIdx就是gridDim为shape的索引。所以称blockIdx、threadIdx为索引,启动核函数后,枚举每一个维度值,不同线程取值不同。
5)threadIdx[x, y, z]:线程索引,threadIdx的最大值是0到blockDim-1,由执行器根据当前执行的线程进行赋值,核函数内访问时已经被配置好,threadIdx一定是小于blockDim的。threadIdx就是blockDim为shape的索引。所以称blockIdx、threadIdx为索引,启动核函数后,枚举每一个维度值,不同线程取值不同
6)Dim是固定的,启动后不会改变,并且是Idx的最大值
7)每个都具有x、y、z三个维度,分别以z、y、x为高低顺序
8)layout是设置核函数执行的线程数,要明白最大值、block最大线程数、warpsize取值
maxGridSize对应gridDim的取值最大值
maxThreadsDim对应blockDim的取值最大值
warpSize对应线程束中的线程数量
maxThreadsPerBlock对应blockDim元素乘积最大值
9)重点!!!
一般情况,为了简化问题,我们只需要用到 threadIdx.x,blockIdx.x,blockDim.x 这三个量即可,所以计算idx 的公式如下:int idx = threadIdx.x + blockIdx.x * blockDim.x。
如果是个2d的layout,什么是2d的layout呢,举个例子dim3 blockDim(32,32),dim3 gridsDim(13,13),即在blockDim和gridDim中传入了两个参数,对dim3的构造函数中的参数x和y都进行了传参。其实不用担心,我们可以看成两个1d的layout,1d的layout就是dim3 blockDim(32)这个样子。比如int idx = threadIdx.x + blockIdx.x * blockDim.x和int idy = threadIdx.y + blockIdx.y * blockDim.y; 其表示的含义是要求thread的1D idx,先得知道在第几个block里,再知道在这个block里的第几个thread。在图像处理加速中就是图像有多大启动的线程就有多少个
Ref:
1.2.CUDA核函数
cuda中threadIdx、blockIdx、blockDim和gridDim的使用
CUDA线程网络的分配
An introduction to CUDA in Python (Part 2)