【cuda】核函数的定义和使用

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)