uda内核函数
在GPU中调用的函数是CUDA内核函数(Kernel function )，该函数由GPU上的多个线程执行。 针对每个线程执行内核函数中的代码。 当然，线程编号不同，运行的代码路径也可能不同。 以下代码行是向量相加计算的CUDA内核函数。
__global__voidvectoradd(int*a、int *b、int *c、int n ) {
inti=block dim.x * blockidx.xthreadidx.x；
if(In ) {
c[i]=a[i] b[i]；
}
}
Vectoradd(d_a，d_b，d_c，n )；
从这段代码中可以看出CUDA内核函数的几个特征。
-声明标识符__global__位于函数的开头。 这个标识符表示这个函数可以在GPU上执行。 另外，尽管正在GPU中执行，但是从CPU侧被调用
-调用内核函数时，必须使用.符号指定线程配置
-在内核函数内部，可以调用threadIdx、blockDim等CUDA内置变量
-内核函数相对于CPU代码异步。 这意味着控制将在内核函数执行完成之前返回。 这样，CPU就可以继续执行后续的CPU代码，而无需等待内核函数完成
稍后将详细介绍线程的配置和内置变量。 除了以上几个特征外，CUDA内核函数还具有一些限制。
-在内核函数内部只能访问device存储器。 内核函数在设备端执行，因此只能访问设备端的内存。
必须返回void类型。 我们知道内核函数是在CPU端启动并在GPU上执行的函数。 内核函数内部的数据都在GPU上，如果内核函数有返回值，那么返回值就是GPU上的数据，不允许CPU直接接收这个数据。 因此，内核函数没有返回值。
-内核函数不支持可变参数
-内核函数不支持静态变量
-内核函数不支持函数指针
在CUDA编程中，除了__global__之外，常见的标识符包括：
_ _ device _ _ u
-具有标识符__device__的函数只能在GPU段上执行
只能在GPU段中调用，如__global__或__device__函数
-__global__和__device__不能同时使用
另一个常见的标识符是__host__ :
-只能在主机端执行
-只能在主机端调用
单独使用__host__时，该函数与普通CPU函数的性质和使用方法没有任何差异。 那么，为什么需要引入这个标识符呢？ 如果您期望同时在CPU和GPU上调用某个函数，则声明_ _ host _ _ _ device _ _ funforcpuandgpu (args )可以在CPU和GPU上执行此函数
线程配置
如上所述，调用内核函数时，必须在.中指定线程配置。 在具体说明之前，让我们先了解一些CUDA编程的基本概念。
线程(Thread )是CUDA程序的基本执行单位，各线程内的执行将依次执行。 所有线程执行相同的代码。 当然，它可能执行同一代码的不同分支。 所有线程都并行执行，没有优先级。
线程块(Thread Block )由一组线程组成。 具有内存-共享内存，可以在每个线程块中的线程之间进行协作，并可以共同访问。 每个线程块都在GPU上的一个流处理器(流多处理器，sm )中执行。
线程网格是线程块的集合。 线程网格中的线程块计划在GPU的多个SM上运行。 线程块之间没有同步机制，线程块执行的优先级没有确定。 线程块之间的通信很昂贵，需要通过全局存储器(global memory )来实现。
CDA线程层次结构。 每个内核函数对应一个线程网格，每个线程网格包含多个线程块，每个线程块包含多个线程。 线程网格和线程块可以是一维、二维或三维的。
调用内核函数时必须指定的线程配置是指每个线程网格中有多少个线程块，每个线程块中有多少个线程，以及它们的排列方式。 线程配置的示例如下。
dim 3网格(3，2，1 )、block (3，1 ) ) )。
Kernel_name(…)
线程网格和线程块的数据类型为dim3，实际上是一个结构，有三个变量用于描述x、y和z三个方向的长度。的第一个参数指定线程网格的结构，即每个线程网格中有多少个线程块。 在上面的示例中，每个线程网格包含321=6个线程块，它们在三个方向上分别排列为3、2、1。 第二个参数用于指定线程
块的结构，也就是每个线程块中有多少个线程，上面的例子中每个线程块中有531=15个线程，排布方式是三个方向上分别是5、3、1。<<<...>>>也可以接受整型变量，比如<<<6, 32>>>代表一个线程网格中有6个线程块，一维排布，一个线程网格块内有32个线程，同样一维排布。这样整个核函数内的总的线程数就是6*32=192。另外核函数内部可以使用CUDA的内置变量来获取线程号以及线程块号：
threadIdx.[x y z]指的是线程块内线程的编号
blockIdx.[x y z]指的是线程网格内线程块的编号
blockDim.[x y z]指的是线程块的维度，也就是线程块中每个方向上线程的数目
gridDim.[x y z]指的是线程网格的维度，也就是线程网格中每个方向上线程块的数目
下面我们来看一个简单的例子，线程网格有4个线程块，每个线程块内有8个线程，并且都是一维排布：
kernel_name<<<4, 8>>> (argument list)
具体的线程配置以及相应的内置变量的值如下图所示
内置变量均从0开始编号
从上图可以看出，我们可以很轻易的获取一个线程在线程块的位置。在核函数中，我们经常需要得知一个线程在一个线程网格中的位置，那么该怎么计算呢？同样来看一个简单的例子：
dim3 grid(4,1,1), block(4,1,1)
上面的例子中有4个线程块，每个线程块中4个线程，假设我们需要计算红色标记的线程在线程网格中的位置。观察上图，我们可以分成两个部分进程计算，首先计算该线程所在线程块前面总共有多少线程，然后在加上该线程在当前线程块的位置就可以获取在整个线程网格中的位置。该线程所在的线程块编号是blockIdx.x，每个线程块内的线程数是blockDim.x，那总的线程数是blockIdx.x * blockDim.x. 再加上该线程在当前线程块中的位置threadIdx.x，则有：
int idx = blockIdx.x * blockDim.x + threadIdx.x;
二维与三维的情况会稍微复杂些，但计算方法是一样的。下面是一段打印二维线程编号的核函数的例子，自己可以尝试编译运行，相信会有助于对线程位置的计算加深理解。
#include
#include
__global__ void printThreadIndex() {
int ix = threadIdx.x + blockIdx.x * blockDim.x;
int iy = threadIdx.y + blockIdx.y * blockDim.y;
unsigned int idx = iy*blockDim.x * gridDim.x + ix;
printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d, %d), global index %2d n",
threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, ix, iy, idx);
}
int main(void) {
dim3 grid(2, 3, 1), block(4, 8, 1);
printThreadIndex<<>>();
cudaResetDevice();
return 0;
}