CUDA -- 并行计算入门

知道了CUDA编程基础，我们就来个简单的实战：利用CUDA编程实现两个向量的加法。在实现之前，先简单介绍一下CUDA编程中内存管理API。首先是在device上分配内存的cudaMalloc函数。

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

这个函数和C语言中的malloc类似，但是在device上申请一定字节大小的显存，其中devPtr是指向所分配内存的指针。同时要释放分配的内存使用cudaFree函数，这和C语言中的free函数对应。另外一个重要的函数是负责host和device之间数据通信的cudaMemcpy函数：

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

其中src指向数据源，而dst是目标区域，count是复制的字节数，其中kind控制复制的方向：cudaMemcpyHostToHost,  cudaMemcpyHostToDevice,  cudaMemcpyDeviceToHost 及 cudaMemcpyDeviceToDevice，如cudaMemcpyHostToDevice将host上数据拷贝到device上。

需要指出的是cudaMemcpy是阻塞式的API，也就是CPU端代码在调用该API时，只有当该API完成拷贝之后，CPU才能继续处理后面的任务。这有一个好处就是保证了计算结果已经完全从GPU端拷贝到了CPU。同时CUDA也提供了非阻塞拷贝的APIcudaMemcpyAsync， 非阻塞拷贝也称为异步拷贝，指的是该API在拷贝完成之前就返回，使得CPU可以继续处理后续的代码。异步拷贝API使得CPU与GPU之间的数据拷贝与CPU计算的并发称为可能。如果该API与CUDA中流（Stream）相结合使用，也可以实现数据的拷贝与GPU计算进行并发执行，这一点会在流与并发这一部分进行介绍。

在host和device之间异步拷贝数据的一个简单例子如下：

#include "cuda.h" 
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
 
__device__ int devData;
__host__ __device__ int run_on_cpu_or_gpu() {
    return 1;
}
 
__global__ void run_on_gpu() {
    printf("run_on_cpu_or_gpu GPU: %d
", run_on_cpu_or_gpu());
}
 
int main() {
    int val = run_on_cpu_or_gpu();
    cudaMemcpyToSymbol(devData, &val, sizeof(int));
    printf("run_on_cpu_or_gpu CPU: %d
", run_on_cpu_or_gpu());
    cudaMemcpyFromSymbol(&val, devData, sizeof(int));
    run_on_gpu<<<1, 1>>>();
    cudaDeviceReset();
    return 0;
}

现在我们来实现一个向量加法的实例，这里grid和block都设计为1-dim，首先定义kernel如下：

// 两个向量加法kernel，grid和block均为一维
__global__ void add(float* x, float * y, float* z, int n)
{
    // 获取全局索引
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    // 步长
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
    {
        z[i] = x[i] + y[i];
    }
}

然后按照CUDA程序的执行流程继续编写代码:

• 1. 分配host内存，并进行数据初始化；
• 2. 分配device内存，并从host将数据拷贝到device上；
• 3. 调用CUDA的核函数在device上完成指定的运算；
• 4. 将device上的运算结果拷贝到host上；
• 5. 释放device和host上分配的内存。

代码如下：

int main()
{
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    // 申请host内存
    float *x, *y, *z;
    x = (float*)malloc(nBytes);
    y = (float*)malloc(nBytes);
    z = (float*)malloc(nBytes);

    // 初始化数据
    for (int i = 0; i < N; ++i)
    {
        x[i] = 10.0;
        y[i] = 20.0;
    }

    // 申请device内存
    float *d_x, *d_y, *d_z;
    cudaMalloc((void**)&d_x, nBytes);
    cudaMalloc((void**)&d_y, nBytes);
    cudaMalloc((void**)&d_z, nBytes);

    // 将host数据拷贝到device
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
    // 定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    // 执行kernel
    add << < gridSize, blockSize >> >(d_x, d_y, d_z, N);

    // 将device得到的结果拷贝到host
    cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);

    // 检查执行结果
    float maxError = 0.0;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    std::cout << "最大误差: " << maxError << std::endl;

    // 释放device内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    // 释放host内存
    free(x);
    free(y);
    free(z);
    
    return 0;
}

在这里可以附一个完整的利用CUDA 并行化思想来对数组进行求和和CPU求和的对比程序：

// 相关 CUDA 库
#include "cuda_runtime.h"
#include "cuda.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <cstdlib>

using namespace std;

const int N = 100;

// 块数
const int BLOCK_data = 3; 
// 各块中的线程数
const int THREAD_data = 10; 

// CUDA初始化函数
bool InitCUDA()
{
    int deviceCount; 

    // 获取显示设备数
    cudaGetDeviceCount (&deviceCount);

    if (deviceCount == 0) 
    {
        cout << "找不到设备" << endl;
        return EXIT_FAILURE;
    }

    int i;
    for (i=0; i<deviceCount; i++)
    {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 获取设备属性
        {
            if (prop.major>=1) //cuda计算能力
            {
                break;
            }
        }
    }

    if (i==deviceCount)
    {
        cout << "找不到支持 CUDA 计算的设备" << endl;
        return EXIT_FAILURE;
    }

    cudaSetDevice(i); // 选定使用的显示设备

    return EXIT_SUCCESS;
}

// 此函数在主机端调用，设备端执行。
__global__ 
static void Sum (int *data,int *result)
{
    // 取得线程号
    const int tid = threadIdx.x; 
    // 获得块号
    const int bid = blockIdx.x; 
    
    int sum = 0;

    // 有点像网格计算的思路
    for (int i=bid*THREAD_data+tid; i<N; i+=BLOCK_data*THREAD_data)
    {
        sum += data[i];
    }
    
    // result 数组存放各个线程的计算结果
    result[bid*THREAD_data+tid] = sum; 
}

int main ()
{
    // 初始化 CUDA 编译环境
    if (InitCUDA()) {
        return EXIT_FAILURE;
    }
    cout << "成功建立 CUDA 计算环境" << endl << endl;

    // 建立，初始化，打印测试数组
    int *data = new int [N];
    cout << "测试矩阵: " << endl;
    for (int i=0; i<N; i++)
    {
        data[i] = rand()%10;
        cout << data[i] << " ";
        if ((i+1)%10 == 0) cout << endl;
    }
    cout << endl;

    int *gpudata, *result; 
    
    // 在显存中为计算对象开辟空间
    cudaMalloc ((void**)&gpudata, sizeof(int)*N); 
    // 在显存中为结果对象开辟空间
    cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data*THREAD_data);
    
    // 将数组数据传输进显存
    cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice); 
    // 调用 kernel 函数 - 此函数可以根据显存地址以及自身的块号，线程号处理数据。
    Sum<<<BLOCK_data,THREAD_data,0>>> (gpudata,result);
    
    // 在内存中为计算对象开辟空间
    int *sumArray = new int[THREAD_data*BLOCK_data];
    // 从显存获取处理的结果
    cudaMemcpy (sumArray, result, sizeof(int)*THREAD_data*BLOCK_data, cudaMemcpyDeviceToHost);
    
    // 释放显存
    cudaFree (gpudata); 
    cudaFree (result);

    // 计算 GPU 每个线程计算出来和的总和
    int final_sum=0;
    for (int i=0; i<THREAD_data*BLOCK_data; i++)
    {
        final_sum += sumArray[i];
    }

    cout << "GPU 求和结果为: " << final_sum << endl;

    // 使用 CPU 对矩阵进行求和并将结果对照
    final_sum = 0;
    for (int i=0; i<N; i++)
    {
        final_sum += data[i];
    }
    cout << "CPU 求和结果为: " << final_sum << endl;

    getchar();

    return 0;
}

cuda从入门到精通（四）之并行计算入门_xiangxianghehe的博客-CSDN博客_cuda并行计算