06.CUDA编程模型概述（二）

4. CUDA核函数

  CUDA核函数：在GPU上并行执行的函数称为CUDA核函数（Kernel Function)，它属于CUDA编程中最为重要且核心的一个环节，也是我们重点要写的代码部分。我们先从启动一个核函数开始！

5. 启动一个CUDA核函数

C语言中函数调用语句：

function_name(argument list);

CUDA内核调用是对C语言函数调用语句的扩展，<<<>>>运算符内是核函数的执行配置，即需要指定网格和块的维度。

kernel_name<<<grid,block>>>(argument list);

执行配置的第一个值是网格维度，也就是启动块的数目。第二个值是块维度，也就是每个块中线程的数目。通过指定网格和块的维度，你可以进行以下配置：

• 内核中线程的数目
• 内核中使用的线程布局

  同一个块中的线程之间可以相互协作，不同块内的线程不能协作。对于一个给定的问题，可以使用不同的网格和块布局来组织你的线程。例如，假设你有32个数据元素用于计算，每8个元素一个块，需要启动4个块。我们可以使用dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化，如下：

kernel_name<<<4,8>>>(argument list);

由于数据在全局内存中是线性存储的，因此可以用变量blockIdx.x和threadId.x来进行以下操作。

• 在网格中标识一个唯一的线程
• 建立线程和数据元素之间的映射关系

如果把所有32个元素放到一个块里，那么只会得到一个块：

kernel_name<<<1,32>>>(argument list);

如果每个块只含有一个元素，那么会有32个块：

kernel_name<<<32,1>>>(argument list);

核函数的调用与主机线程是异步的。核函数调用结束后，控制权立刻返回给主机端。你可以调用以下函数来强制主机端程序等待所有的核函数执行结束：

cudaError_t cudaDeviceSynchronize(void);

这是一个显示的方法，对应的也有隐式方法，隐式方法就是不明确说明主机要等待设备端，而是设备端不执行完，主机没办法进行，比如内存拷贝函数：

cudaError_t cudaMemcpy(void* dst,const void * src,size_t count,cudaMemcpyKind kind);

6. 编写核函数

  核函数用__global__符号声明，在devie（GPU）上执行，在host（CPU）上调用，返回类型必须时void，不支持可变参数，不支持静态变量，不支持函数指针，核函数相对于CPU是异步的，在核函数执行完之前就会返回，这样CPU可以不用等待核函数的完成，继续执行后续代码。在host端核函数的调用方式为：

__global__ void kernel_name(argument list);

  我们之前讲过CPU-GPU的异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词来区分host和device上的函数，主要的三个函数类型限定词如下：

限定符	执行	调用	备注
__ global__	设备端执行	可以从主机调用，也可以从计算能力3以上的设备调用	必须有一个void的返回类型
__ device__	设备端执行	设备端调用
__ host __	主机端执行	主机端调用	一般忽略不写

  __device__和__host__限定符可以一齐使用，这样函数可以同时在主机和设备端进行
编译。

6.1 CUDA核函数的限制

核函数编写有以下限制：

• 只能访问设备内存
• 必须具有void返回类型
• 不支持可变数量的参数
• 不支持静态变量
• 显示异步行为

考虑一个简单的例子：循环将两个大小为N的向量A和B相加，主机端代码：

void sumArraysOnHost(float *A, float *B, float *C, const int N) {
   
  for (int i = 0; i < N; i++)
    C[i] = A[i] + B[i];
}

设备端代码：

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
   
  int i = threadIdx.x;
  C[i] = A[i] + B[i];
}

可以看到设备端代码中的循环体消失了，内置的线程坐标变量替换了数组索引。

7. 验证核函数

  开发过程中，验证自己的代码是必须的且高效的。这里的验证方法是为了验证我们最开始写的两个数组相加的程序，通过比较两个数组在主机端和设备端分别执行相加得出的结果，来进行验证核函数即设备端的代码是否正确。

#include <stdio.h>
#include <math.h>
void checkResult