CUDA 简单程序的基本框架和自定义设备函数

1 cuda程序的基本框架

框架包含：

• 头文件
• 常量或者宏定义
• C++自定义函数和cuda核函数的原型声明
• main函数
• C++自定义函数核CUDA核函数的定义实现
  其中main函数中

int main（）
{
分配主机与设备代码内存
初始化主机中的数据
将某些数据从主机复制到设备
调用核函数在设备中进行计算
将某些数据从设备复制到主机
释放主机与设备内存
}

示例代码如下。

#include<math.h>
#include<stdio.h>
const double EPSILON=1.0e-15;
const double a=1.23;
const double b=2.34;
const double c=3.57;
void __global__ add(const double *x,const double *y,double *z);
void check(const double *z,const int N);
int main(void)
{
        const int N=100000000;
        const int M=sizeof(double) *N;
        double *h_x=(double*)malloc(M);
        double *h_y=(double*)malloc(M);
        double *h_z=(double*)malloc(M);
        for(int n=0;n<N;++n)
        {
           h_x[n]=a;
           h_y[n]=b;
        }
        double *d_x,*d_y,*d_z;
        cudaMalloc((void **)&d_x,M);
        cudaMalloc((void **)&d_y,M);
        cudaMalloc((void **)&d_z,M);
        cudaMemcpy(d_x,h_x,M,cudaMemcpyHostToDevice);
        cudaMemcpy(d_y,h_y,M,cudaMemcpyHostToDevice);
        const int block_size=128;
        const int grid_size=N/block_size;
        add<<<grid_size,block_size>>>(d_x,d_y,d_z);
        cudaMemcpy(h_z,d_z,M,cudaMemcpyHostToDevice);
        check(h_z,N);
        free(h_x);
        free(h_y);
        free(h_z);
        cudaFree(d_x);
        cudaFree(d_y);
        cudaFree(d_z);
        return 0;
}
void __global__ add(const double *x,const double *y,double *z)
{
        const int n = blockDim.x*blockIdx.x+threadIdx.x;
        z[n]=x[n]+y[n];
}
void check(const double *z,const int N)
{
        bool has_error=false;
        for(int n=0;n<N;++n)
        {
           if(fabs(z[n]-c)>EPSILON)
           {
                has_error =true;
           }
          printf("%s\n",has_error?"Has errors":"No errors");
        }
}

1.1 隐形的设备初始化
　　在cuda运行时API中，没有明显地初始化设备（GPU）的函数。在第一次调用一个和设备管理及版本管理及版本查询功能无关的运行时API函数时，设备将自动地初始化。

1.2 设备内存的分配与释放
　　在上述程序中，我们首先在主机中定义了三个数组并进行了初始化，通过cudaMalloc函数将它们指向了设备中的内存，而不是主机的内存。该函数是一个cuda运行时API函数，所有cuda运行时API函数都以cuda开头。
　　在c++中由malloc()函数动态分配内存，在cuda中，设备内存的动态分配可由cudaMalloc()函数实现，函数原型如下。
　　cudaError_t cudaMalloc(void **address，size_t size);

• 其中第一个参数address是待分配设备内存的指针。因为内存（地址）本身就是一个指针，所以待分配设备内存的指针就是指针的指针，即双重指针。
• 第二个参数size是待分配内存的字节数。
• 返回值是一个错误代码，如果调用成功，返回cudaSuccess,否则会返回一个代表错误的代号。

　　调用cudaMalloc（）函数时传入的第一个参数（void *）&d_x比较难懂。其中d_x是一个double类型的指针，他的地址就是指针的指针，也就是双重指针，而使用（void * * ）是一个强制类型转换操作。转换为void类型的双重指针。这种类型的转换可以不明确的写出来。所以cudaMalloc()函数的调用也可以简写为 cudaMalloc(&d_x,M);
　　用cudaMalloc()函数为什么需要一个双重指针作为变量呢？该函数的功能是改变指针d_x本身的值（将一个指针赋值给d_x）,而不是改变d_x所指内存缓冲区中的变量值。需要将d_x的地址&d_x传给函数cudaMalloc()才能达到修改指针d_x本身的值的效果。
　　总之，使用cudaMalloc函数可以为不同类型的指针变量分配设备内存。为了区分主机和设备中的变量，使用d_作为所有设备变量的前缀，使用h_作为对应主机变量的前缀。
　　正如mallloc()函数分配的主机内存需要使用free()释放一样，用cudaMalloc()函数分配的设备内存需要用cudaFree()函数释放。该函数原型为cudaError_t cudaFree(void address);
　　这里参数address就是待释放的设备内存变量（不是双重指针），返回值是一个错误代号。

1.3 主机与设备之间数据的传递
　　在分配了设备内存之后，就可以将一些数据从主机传递到设备中去，使用cudaMemcpy()方法主机中的变量数据复制到设备中相应变量d_x和d_y所指向的缓冲区中。其方法的原型是。

• 第一个参数dst是目标地址。
• 第二个参数src是源地址。
• 第三个参数count是复制数据的字节数。
• 第四个参数kind是一个枚举类型的变量，标志数据传递方向。其中udaMemcpyHostToHost表示从主机复制到主机。还有其他数据传递方向。
• 返回值是一个错误代号
• 该函数的作用是将一定字节的数据从源地址所值缓存区复制到目标地址所指缓存区。

1.4 核函数的要求

• 1.返回值是void
• 2.限定符
• 3.函数名无特殊要求，支持C++中的重载，可以用同一个函数名表示具有不同参数列表的函数。
• 4.不支持可变数据的参数列表，即参数的个数必须确定。
• 5.可以向核函数传递非指针变量，其内容对每个线程可见。
• 6.除非使用统一内存编程机制，否则传给核函数的数组（指针）必须指向设备内存。
• 7.核函数不可成为一个类的成员，通常是用一个包装函数调用核函数，而将包装函数定义为类的成员。
• 8.在计算能力3.5之前，核函数之间不能互相调用。从计算能力3.5之后，引入 了动态并行机制，在核汉书内部可以调用其他核函数，甚至可以调用自己（递归）。
• 9.无论是从主机调用，还是从设备调用，核函数都是在设备中执行。调用核函数时必须指定执行配置，即三括号和它里面的参数。

2 自定义设备函数
　　核函数可以调用不带执行配置的自定义函数，这样的自定义函数称为设备函数。他是在设备中执行，并在设备中调用的。与之相比，核函数是在设备中执行，但是在主机端被调用的。

2.1函数执行空间表示符

• 1.用__gloabal__ 修饰的函数称为核函数，一般是由主机调用，在设备中执行。
• 2.用__device__修饰的函数成为设备函数，只能被核函数或者其他设备函数调用，在设备中执行。
• 3.用 host 修饰的函数就是主机端的普通C++函数，在主机中被调用，在主机中执行，修饰符可省。
• 4.不能同时用device和global 修饰一个函数，即不能将一个函数同时定义为设备函数和核函数。
• 5.也不能同时用host核global修饰一个函数。

2.2 代码示例

double __device__ add1_device(const double x, const double y)
 {
return (x + y);
}
void __global__ add1(const double *x, const double *y, double *z,
const int N)
 {
   const int n = blockDim.x * blockIdx.x + threadIdx.x;
if (n < N)
{
    z[n] = add1_device(x[n], y[n]);
}
}