cuda并行计算|cuda 核函数的定义和使用 cuda并行计算|cuda

1.cu文件一般写cuda的核函数 2.在.vscode/settings.json中配置*.cu : cuda-cpp，可以使得代码被正确解析 3.Makefile中，cu交给nvcc进行编译，nvcc是一个编译器，用来编译cuda的c程序的 4.cu文件可以当做正常cpp写即可，他是cpp的超集，兼容支持cpp的所有特性 5.cu文件中引入了一些新的符号和语法： 5.1_global__标记，核函数标记
1）调用方必须是host
2）返回值必须是void
3）例如：global void kernel(const float* pdata, int ndata)
4）必须以
kernel<<>>(pdata, ndata)的方式启动，只有__global__ 修饰的函数才可以用<<<>>>的方式调用
其参数类型是：
<<>>
dim3有构造函数dim3(int x, int y=1, int z=1)，因此当直接赋值为int时，实则定义了dim.x = value, dim.y = 1, dim.z = 1。value就是参数传进来的那个值。bytesSharedMemorySize参数的单位为字节。
其中gridDim, blockDim, bytesSharedMemory, stream是线程layout（线程布局）参数（下文第8点中有详细解释），gridDim, blockDim用来告诉这个函数需要启动多少个线程，可以理解为启动了gridDim个block，每个block有blockDim个线程。 gridDim和blockDim的类型都是dim3，dim3是一个结构体，启动的线程总数量nthreads就等于

dim3 gridDim; dim3 blockDim; int nthreads = gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z

线程总数量nthreads的上限是可以通过runtime API或者deviceQuery可以查询到。
如果指定了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.内存索引（线程的绝对索引即在线程总数里这段程序是属于第几个线程）的计算公式方便记忆的办法，是左乘右加，无论tensor维度多复杂，这个方法都适用，核函数里把blockDim和gridDim看做shape，把threadIdx和blockIdx看做index，下面代码中的dims和indexs如下图所示：

position = 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为高低顺序
【cuda并行计算|cuda 核函数的定义和使用】8）layout是设置核函数执行的线程数，要明白最大值、block最大线程数、warpsize取值
maxGridSize对应gridDim的取值最大值
maxThreadsDim对应blockDim的取值最大值
warpSize对应线程束中的线程数量
maxThreadsPerBlock对应blockDim元素乘积最大值