CUDA程序基本框架

参考书目： CUDA编程基础与实践

作者: 樊哲勇

章节：第3章 CUDA程序基本框架

先看一段简单的数组相加的模板

#include <math.h>
#include <stdio.h>
#include <time.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)
{
    time_t start, end;
    time(&start);

    const int N = 100000000;
    const int M = sizeof(double) * N;
    double *h_x = (double*) malloc(M);//这里是一重指针是因为这是返回值,直接赋值,而cudaMalloc是修改传的参数
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);
    // double *h_x = new double[M]; //也可以使用new运算符开辟空间 delete释放
    // double *h_y = new double[M];
    // double *h_z = new double[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是要求用传双重指针的方式改变一个指针的值,而不是直接返回一个指针,该函数的返回值已经用于返回错误代号
    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, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);
    // delete[] h_x; //也可以使用new运算符开辟空间 delete释放
    // delete[] h_y;
    // delete[] h_z;
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    time(&end);
    double elapsed = difftime(end, start);
    printf("Elapsed time: %.2f seconds\n", elapsed);
    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("Did %d additons, %s\n",N, has_error ? "Has errors" : "No errors");
}

相关API介绍

内建变量blockIdx,threadIdx,gridDim,blockDim

blockIdx,threadIdx为uint3变量,是一个结构体,有x,y,z三个成员
gridDim,blockDim为dim3变量,是一个结构体,有x,y,z三个成员

满足下列关系

blockIdx.x $\in [0, gridDim.x-1]$
blockIdx.y $\in [0, gridDim.y-1]$
blockIdx.z $\in [0, gridDim.z-1]$
threadIdx.x $\in [0, blockDim.x-1]$
threadIdx.x $\in [0, blockDim.y-1]$
threadIdx.x $\in [0, blockDim.z-1]$

类似malloc()函数动态分配内存,在cuda中分配显存由cudaMalloc()实现,原型如下

cudaError_t cudaMalloc(void **address, size_t size);

我们知道如果要修改函数传入的参数需要以其指针或引用的形式传入,这里我们需要修改内存地址,其本身就是一个指针,所以传入的参数为双重指针,不同于malloc(),它是返回一个内存地址,不需要修改原有的变量,而是直接赋值

类似内存释放函数free(), cuda中显存释放为cudaFree(),原型为

cudaError_t cudaFree(void *address)

返回值都为错误代号

数据传递函数cudaMemcpy()原型

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

参数分别表示目标地址,源地址,要复制的字节数, 传递方向变量kind如下

enum __device_builtin__ cudaMemcpyKind
{
    cudaMemcpyHostToHost          =   0,      /**< Host   -> Host */
    cudaMemcpyHostToDevice        =   1,      /**< Host   -> Device */
    cudaMemcpyDeviceToHost        =   2,      /**< Device -> Host */
    cudaMemcpyDeviceToDevice      =   3,      /**< Device -> Device */
    cudaMemcpyDefault             =   4       /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
};

cudaMemcpyDefault表示自动判断方向

可以使用一个线程与多个线程进行对比

void __global__ add(const double *x, const double *y, double *z)
{
    for (int n = 0; n < 100000000; ++n)
    {
        z[n] = x[n] + y[n];
    }
}

调用核函数时只配置一个线程

add<<<1, 1>>>(d_x, d_y, d_z);

通过运行时间可以看出并行计算的优势

使用if语句对数据越界进行判断

示例为正好所有线程(grid_size*block_size)都进行计算的情况,但也有计算的元素数不等于开辟的线程数的情况也就是说选择合适的grid_size和block_size要大于等于要计算的数,至于多开辟的那些线程,没有对应的显存空间, 就要使用if进行判断不对其进行计算

void __global__ add(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] = x[n] + y[n];
    }
}

自定义设备函数

CUDA函数标识符

`__global__`	`__device__`	`__host__`
核函数,主机调用,设备执行,无返回值	设备函数,只能被核函数和设备函数调用	主机函数,普通的C++函数

以下是核函数调用设备函数的示例

double __device__ add1_device(const double x, const double y)//可以有返回值,也可以传入指针或者引用对z进行修改
{
    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]);
    }
}