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\n", 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\n", 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並行計算