第六篇：二維陣列的傳輸 (host <-> device)

前言

       本文的目的很明確：介紹如何將二維陣列傳遞進視訊記憶體，以及如何將二維陣列從視訊記憶體傳遞迴主機端。

實現步驟

       1. 在視訊記憶體中為二維陣列開闢空間

       2. 獲取該二維陣列在視訊記憶體中的 pitch 值 (cudaMallocPitch 實現)

       3. 將二維陣列傳遞進視訊記憶體 (cudaMemcpy2D 實現)

       4. 在視訊記憶體中對該二維陣列進行處理 (目前必須按照 1 維陣列的規則進行處理)

       5. 將結果傳遞迴記憶體 (cudaMemcpy2D實現)

重要概念 - pitch

       對於記憶體的存取來說，對準偏移量為2的冪（現在一般要求2^4=16）的地址能獲取更快的速度，而如果不對齊，可能你需要的資料需要更多的存取次數才能得到。

       為了滿足這個條件，對於一個二維陣列來說（行優先row major），就希望每一行的開頭都滿足“對齊”。如果一行的長度不規整，導致下一行開頭不在指定的位置，就需要在每一行末尾進行填充（padding），從而使得每一行都對齊，這和BMP格式的畫素儲存是一個道理。

       pitch就是指每一行的位元組數 + padding的位元組數 。

       使用 cudaMemcpy2D 的目的僅是為了利用 pitch 機制提升二維陣列中元素的訪問速度。事實上二維陣列傳遞進視訊記憶體後，還是得按照一維陣列的規範去處理該二維矩陣。global 函式中目前不支援多下標訪問，好坑。。

程式碼示例

#include "cuda.h"
#include "cuda_runtime.h"

#include <iostream>

using namespace std;

// 定義測試二維陣列的行列數
const int R = 5;
const int C = 10;

int main()
{
    // 定義一個用於測試的二維陣列，其每個元素都賦值為0，並將其列印出來。
    int array2D[R][C];
    cout << "傳輸前的測試矩陣:" << endl;
    for (int i=0; i<R; i++) {
        for (int j=0; j<C; j++) {
            array2D[i][j] = 0;
            cout << array2D[i][j] << " ";
        }
        cout << endl;
    }

    // 再定義另一個同樣大小的二維陣列用於獲取從視訊記憶體傳回的結果
    int result[R][C];
    cout << "傳輸前的結果矩陣:" << endl;
    for (int i=0; i<R; i++) {
        for (int j=0; j<C; j++) {
            result[i][j] = 1;
            cout << result[i][j] << " ";
        }
        cout << endl;
    }

    // 為此二維陣列在視訊記憶體中分配記憶體
    int *d_array2D;
    cudaMalloc ((void**)&d_array2D, sizeof(int)*R*C);

    // 獲取視訊記憶體中的二維陣列的 pitch 值
    size_t d_pitch;
    cudaMallocPitch ((void**) &d_array2D, &d_pitch, sizeof(int)*C, R);

    // 將二維陣列轉移進視訊記憶體
     cudaMemcpy2D (
         d_array2D,    // 目的地址
         d_pitch,    // 目的 pitch
         array2D,    // 源地址
         sizeof(int)*C,    // 源 pitch
         sizeof(int)*C,    // 資料拷貝寬度
         R,    // 資料拷貝高度
         cudaMemcpyHostToDevice    // 資料傳遞方向
        );

     // 將二維陣列從視訊記憶體傳輸回主機端的結果矩陣中
     cudaMemcpy2D (
         result,    // 目的地址
         sizeof(int)*C,    // 目的 pitch
         d_array2D,    // 源地址
         d_pitch,    // 源 pitch
         sizeof(int)*C,    // 資料拷貝寬度
         R,    // 資料拷貝高度
         cudaMemcpyDeviceToHost    // 資料傳遞方向
        );

    // 列印傳回到結果矩陣的資料 
    cout << "從視訊記憶體獲取到測試矩陣後的結果矩陣:" << endl;
    for (int i=0; i<R; i++) {
        for (int j=0; j<C; j++) {
            cout << result[i][j] << " ";
        }
        cout << endl;
    }

    cudaFree (d_array2D);

    cin.get();

    return EXIT_SUCCESS;
}

執行測試

      

小結

       本文介紹的僅僅是二維陣列在兩端之間的傳輸！當二維陣列傳遞進了視訊記憶體，在對其操作的過程中，是需要對其進行一個一維到二維的下標操作轉換的，global 中不支援多下標訪問。之所以加入 pitch 並使用 cudaMemcpy2D 只是為了提高元素的訪問速度。

       如果需要具體處理傳遞進入的二維陣列，還要將 pitch 也作為引數傳遞進 kernel 函式，如下所示：

// 下面的 kernel 函式將二維陣列的所有位置為 2
__global__
void kernelFun (int *d_array2D, int pitch)
{
    for (int i=0; i<R; i++) {
        int *row = (int *)((char *)d_array2D+i*pitch);
        for (int j=0; j<C; j++) {
            row[j] = 2;
        }
    }

    return;
}