核函数

1. 核函数是cuda编程关键
2. 通过创建.cu创建cudac程序文件，然后交给nvcc编译
3. 加上__global__前缀的函数由host端调用，__device修饰的函数为device函数，由设备调用
4. __host修饰的函数为host函数，由设备调用
5. host调用核函数的方式是function<<<gridDim,blockDim,sharedMemorySize,stream>>>(args,…)

Cuda线程&线程id

CUDA里面用了grid和block作为线程的组织单位，一个grid可以包含多个block，一个block包含多个thread，grid和block都是三维向量，下段代码展示了如何定义一个grid和block：

dim3 grid(x.y,z);
dim3 block(x,y,z);

在计算线程索引的时候，经常会看到gridDim、blockDim、blockIdx、threadIdx。关于这些名词的解释如下：

• gridDim：block在gird里面的数量维度，
• blockDIm: threads在blcok中的维度
• blockIdx：block在grid中的索引
• threadIdx: thread在block中的索引

计算线程坐标

怎么从那么多的线程中找到目标线程，这就需要计算线程坐标，可以拆成以下三步去计算：

• 获取thread在block中的位置
• 获取block在grid中的位置
• 计算block有多少线程，求解位置索引

计算thread在block中的位置

$$threadInBlock=threaIdx.x+threadIdx.y\ast blockDim.x+threadIdx.z\ast blockDim.x\ast blockDim.y$$
一开始我并不理解这个计算公式，后来为了加强记忆，理解成下面这样（不对的话请指证）：
在三维block中计算线程的一维索引时，可以想像成将三维数据展平成一维，按照“优先顺序”的话，x维度是变化最快的，其次是y维度，最后是z维度。因此计算方式是：

• threadIdx.x 直接加到索引上，因为它是最内层的循环（X 维度）。
• threadIdx.y * blockDim.x 是因为每增加一个 Y 索引，你跳过了整个 X 维度的线程数。
• threadIdx.z * blockDim.x * blockDim.y 是因为每增加一个 Z 索引，你跳过了整个 X-Y 平面的线程数

计算block在grid中的位置

$$blockInGrid=blockIdx.x+blockIdx.y\ast girdDim.x+blockIdx.z\ast girdDim.x\ast gridDim.y$$

求解全局idx:

一个block中的thread总数为：
$perBlockSIze=blockDim.x\ast blockDim.y$
因此$idx=threadInBlock+blockIngrid\ast perBlockSize$
上式可以理解为：每个block有perBlockSize个thread，对应线程所在的block起始位置就是blockInGrid*perperBlockSize，然后加上偏移就是threadInblock