编写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)