编写CUDA核函数验与证核函数杂谈

标签：__ 函数 res void float CUDA 证核

编写核函数 核函数也是一个函数，但是声明核函数有一个比较模板化的方法： global__ void kernel_name(argument list); 注意：声明和定义是不同的，这点CUDA与C语言是一致的 在C语言函数前没有的限定符global，CUDA C中还有一些其他在C中没有的限定符，见表10-2。 表10-2 CUDA C中一些其他在C中没有的限定符

限定符	执行	调用	备注
__global__	设备端执行	可以从主机调用也可以从计算能力3以上的设备调用	必须有一个void的返回类型
__device__	设备端执行	设备端调用
__host__	主机端执行	主机调用	可以省略

有些函数可以同时定义为 device 和 host，这种函数可以同时被设备和主机端的代码调用，主机端代码调用函数很正常，设备端调用函数与C语言一致，但是要声明成设备端代码，告诉nvcc编译成设备机器码，同时声明主机端设备端函数，那么就要告诉编译器，生成两份不同设备的机器码。 Kernel核函数编写有以下限制 1）只能访问设备内存 2）必须有void返回类型 3）不支持可变数量的参数 4）不支持静态变量 5）显示异步行为 并行程序中经常的一种现象：把串行代码并行化时对串行代码块for的操作，也就是把for并行化。
例如：
    串行： 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]; } 这两个简单的段不能执行，但是可以大致的看一下for展开并行化的样子。 验证核函数 验证核函数就是验证其正确性，下面这段代码上文出现过，但是同样包含验证核函数的方法： /* * 3_sum_arrays */ #include <cuda_runtime.h> #include <stdio.h> #include "freshman.h"     void sumArrays(float * a,float * b,float * res,const int size) {   for(int i=0;i<size;i+=4)   {     res[i]=a[i]+b[i];     res[i+1]=a[i+1]+b[i+1];     res[i+2]=a[i+2]+b[i+2];     res[i+3]=a[i+3]+b[i+3];   } } __global__ void sumArraysGPU(float*a,float*b,float*res) {   int i=threadIdx.x;   res[i]=a[i]+b[i]; } int main(int argc,char **argv) {   int dev = 0;   cudaSetDevice(dev);     int nElem=32;   printf("Vector size:%d\n",nElem);   int nByte=sizeof(float)*nElem;   float *a_h=(float*)malloc(nByte);   float *b_h=(float*)malloc(nByte);   float *res_h=(float*)malloc(nByte);   float *res_from_gpu_h=(float*)malloc(nByte);   memset(res_h,0,nByte);   memset(res_from_gpu_h,0,nByte);     float *a_d,*b_d,*res_d;   CHECK(cudaMalloc((float**)&a_d,nByte));   CHECK(cudaMalloc((float**)&b_d,nByte));   CHECK(cudaMalloc((float**)&res_d,nByte));     initialData(a_h,nElem);   initialData(b_h,nElem);     CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));   CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));     dim3 block(nElem);   dim3 grid(nElem/block.x);   sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d);   printf("Execution configuration<<<%d,%d>>>\n",block.x,grid.x);     CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));   sumArrays(a_h,b_h,res_h,nElem);     checkResult(res_h,res_from_gpu_h,nElem);   cudaFree(a_d);   cudaFree(b_d);   cudaFree(res_d);     free(a_h);   free(b_h);   free(res_h);   free(res_from_gpu_h);     return 0; } 在开发阶段，每一步都进行验证是绝对高效的，比把所有功能都写好，然后进行测试这种过程效率高很多，同样写CUDA也是这样的每个代码小块都进行测试，看起来慢，实际会提高很多效率。
CUDA小技巧，当进行调试的时候可以把核函数配置成单线程的： kernel_name<<<1,1>>>(argument list)

标签：__,函数,res,void,float,CUDA,证核
From： https://www.cnblogs.com/wujianming-110117/p/18386450