Abstract: 本文继续上文介绍CUDA编程模型关于核函数以及错误处理部分
Keywords: CUDA核函数，CUDA错误处理
开篇废话
今天的废话就是人的性格一旦形成，那么就会成为最大的指向标，或者说一个人的性格思维方式能够决定这个人的全部生命轨迹，比如有人真的爱学习（比如我，嘻嘻嘻）有人真的不爱学习，没有优劣，只是两种生活态度，因为学习这个事你学一辈子也学不完人类智慧的九牛一毛，而不学习可以有更多的时间进行社会实践，融入社会，荣华富贵，享受生命。这是两种性格，没有好坏，毕竟每个人评价生活的标尺不一。努力追求自己想要的，不要嘲笑别人所追求的，能一起走的就一起走，不能一起走的就各自安好，这就是目前我理解的声明的真谛。
继续CUDA编程模型的后半部分，关于核函数以及错误处理。
- 核函数
- 启动核函数
- 编写核函数
- 验证核函数
- 错误处理
核函数概述
核函数就是在CUDA模型上诸多线程中运行的那段串行代码，这段代码在设备上运行，用NVCC编译，产生的机器码是GPU的机器码，所以我们写CUDA程序就是写核函数，第一步我们要确保核函数能正确的运行产生正切的结果，第二优化CUDA程序的部分，无论是优化算法，还是调整内存结构，线程结构都是要调整核函数内的代码，来完成这些优化的。
我们一直把我们的CPU当做一个控制者，运行核函数，要从CPU发起，那么我们开始学习如何启动一个核函数
启动核函数
启动核函数，通过的以下的ANSI C 扩展出的CUDA C指令：
kernel_name<<<grid,block>>>(argument list);
其标准C的原型就是C语言函数调用
function_name(argument list);
这个三个尖括号’<<<grid,block>>>'内是对设备代码执行的线程结构的配置（或者简称为对内核进行配置），也就是我们上一篇中提到的线程结构中的网格，块。回忆一下上文，我们通过CUDA C内置的数据类型dim3类型的变量来配置grid和block（上文提到过：在设备端访问grid和block属性的数据类型是uint3不能修改的常类型结构，这里反复强调一下）。
通过指定grid和block的维度，我们可以配置：
- 内核中线程的数目
- 内核中使用的线程布局
我们可以使用dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化：
kernel_name<<<4,8>>>(argument list);
上面这条指令的线程布局是：
我们的核函数是同时复制到多个线程执行的，上文我们说过一个对应问题，多个计算执行在一个数据，肯定是浪费时间，所以为了让多线程按照我们的意愿对应到不同的数据，就要给线程一个唯一的标识，由于设备内存是线性的（基本市面上的内存硬件都是线性形式存储数据的）我们观察上图，可以用threadIdx.x 和blockIdx.x 来组合获得对应的线程的唯一标识（后面我们会看到，threadIdx和blockIdx能组合出很多不一样的效果）
接下来我们就是修改代码的时间了，改变核函数的配置，产生运行出结果一样，但效率不同的代码：
- 一个块：
kernel_name<<<1,32>>>(argument list);
- 32个块
kernel_name<<<32,1>>>(argument list);
上述代码如果没有特殊结构在核函数中，执行结果应该一致，但是有些效率会一直比较低。
上面这些是启动部分，当主机启动了核函数，控制权马上回到主机，而不是主机等待设备完成核函数的运行，这一点我们上一篇文章也有提到过（就是等待hello world输出的那段代码后面要加一句）
想要主机等待设备端执行可以用下面这个指令：
cudaError_t cudaDeviceSynchronize(void);
这是一个显示的方法，对应的也有隐式方法，隐式方法就是不明确说明主机要等待设备端，而是设备端不执行完，主机没办法进行，比如内存拷贝函数：
cudaError_t cudaMemcpy(void* dst,const void * src,
size_t count,cudaMemcpyKind kind);
这个函数上文已经介绍过了，当核函数启动后的下一条指令就是从设备复制数据回主机端，那么主机端必须要等待设备端计算完成。
所有CUDA核函数的启动都是异步的，这点与C语言是完全不同的
posted on
