構造体をデバイスメモリにコピーするCUDA

Question

私はCUDAを初めて使用し、CUDAツールキットのドキュメントを参照しています。そこでは、行列の乗算が共有メモリを使用している例が見つかりました。マトリクス構造をホストメモリからデバイスメモリにコピーするとき、データ要素のみがコピーされます。私が理解できないことは、他の変数がデバイスメモリにどのようにコピーされるかです。ここではその後

typedef struct { 
    int width; 
    int height; 
    int stride; 
    float* elements; 
} Matrix; 

を次のデータ転送は、私は幅、ストライドと高さがコピーされているかを理解していないものをここで

void MatMul(const Matrix A, const Matrix B, Matrix C) 
{ 
    // Load A and B to device memory 
    Matrix d_A; 
    d_A.width = d_A.stride = A.width; d_A.height = A.height; 
    size_t size = A.width * A.height * sizeof(float); 
    cudaMalloc(&d_A.elements, size); 
    cudaMemcpy(d_A.elements, A.elements, size, 
       cudaMemcpyHostToDevice); 
    Matrix d_B; 
    d_B.width = d_B.stride = B.width; d_B.height = B.height; 
    size = B.width * B.height * sizeof(float); 
    cudaMalloc(&d_B.elements, size); 
    cudaMemcpy(d_B.elements, B.elements, size, 
    cudaMemcpyHostToDevice); 

    // Allocate C in device memory 
    Matrix d_C; 
    d_C.width = d_C.stride = C.width; d_C.height = C.height; 
    size = C.width * C.height * sizeof(float); 
    cudaMalloc(&d_C.elements, size); 

    // Invoke kernel 
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
    dim3 dimGrid(B.width/dimBlock.x, A.height/dimBlock.y); 
    MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 

    // Read C from device memory 
    cudaMemcpy(C.elements, d_C.elements, size, 
       cudaMemcpyDeviceToHost); 

    // Free device memory 
    cudaFree(d_A.elements); 
    cudaFree(d_B.elements); 
    cudaFree(d_C.elements); 
} 

起こるのコードサンプルであるとして

マトリックス構造でありますデバイスメモリ。ここでは、cudaMallocとcudaMemcpyは要素のためだけです。私がこれを理解するのに欠けているものはありますか？気になる人のために

カーネルコード

__device__ float GetElement(const Matrix A, int row, int col) 
{ 
    return A.elements[row * A.stride + col]; 
} 

// Set a matrix element 
__device__ void SetElement(Matrix A, int row, int col, 
          float value) 
{ 
    A.elements[row * A.stride + col] = value; 
} 

// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is 
// located col sub-matrices to the right and row sub-matrices down 
// from the upper-left corner of A 
__device__ Matrix GetSubMatrix(Matrix A, int row, int col) 
{ 
    Matrix Asub; 
    Asub.width = BLOCK_SIZE; 
    Asub.height = BLOCK_SIZE; 
    Asub.stride = A.stride; 
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row 
             + BLOCK_SIZE * col]; 
    return Asub; 
} 

行列の乗算カーネルコード

__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) 
{ 
    // Block row and column 
    int blockRow = blockIdx.y; 
    int blockCol = blockIdx.x; 

    // Each thread block computes one sub-matrix Csub of C 
    Matrix Csub = GetSubMatrix(C, blockRow, blockCol); 

    // Each thread computes one element of Csub 
    // by accumulating results into Cvalue 
    float Cvalue = 0; 

    // Thread row and column within Csub 
    int row = threadIdx.y; 
    int col = threadIdx.x; 

    // Loop over all the sub-matrices of A and B that are 
    // required to compute Csub 
    // Multiply each pair of sub-matrices together 
    // and accumulate the results 
    for (int m = 0; m < (A.width/BLOCK_SIZE); ++m) { 

     // Get sub-matrix Asub of A 
     Matrix Asub = GetSubMatrix(A, blockRow, m); 

     // Get sub-matrix Bsub of B 
     Matrix Bsub = GetSubMatrix(B, m, blockCol); 

     // Shared memory used to store Asub and Bsub respectively 
     __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 
     __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; 

     // Load Asub and Bsub from device memory to shared memory 
     // Each thread loads one element of each sub-matrix 
     As[row][col] = GetElement(Asub, row, col); 
     Bs[row][col] = GetElement(Bsub, row, col); 

     // Synchronize to make sure the sub-matrices are loaded 
     // before starting the computation 
     __syncthreads(); 
     // Multiply Asub and Bsub together 
     for (int e = 0; e < BLOCK_SIZE; ++e) 
      Cvalue += As[row][e] * Bs[e][col]; 

     // Synchronize to make sure that the preceding 
     // computation is done before loading two new 
     // sub-matrices of A and B in the next iteration 
     __syncthreads(); 
    } 

    // Write Csub to device memory 
    // Each thread writes one element 
    SetElement(Csub, row, col, Cvalue); 
} 

Answer 1

、我々はについて話すのサンプルコードでは、共有メモリのトピックで、NVIDIAのCUDAツールキットのドキュメントにここにあります： CUDA C Programming guide : Shared memory

このサンプルはなぜ機能しますか？ はい、cudaMallocとcudaMemcpy関数を使用して、デバイス側で "要素"配列のみが送信されます。 はい、行列の寸法は、cudaMemcpyでデバイスメモリに明示的にコピーされることなく、デバイス側のカーネル内で使用されます。

同じように配列とパラメータを考慮する必要があります。これらの値がどのようにカーネルに送られるかを説明しましょう。

1. 我々は、CPU側の行列を宣言する、すべてのメンバーは、我々は、大きさを割り当てる
2. 未初期化され、ポインタは、我々は、ポインタが初期化され、API関数でデバイス側のメモリを割り当て、コピー依然として
3. 初期化されていませんデバイスメモリをターゲットとし、通常のホスト配列のようには使用できません。
4. 行列をパラメータとしてカーネルに与えます。ポインタではなく、値で指定します。

それはトリックです。完全な構造インスタンスは、パラメータとして与えられ、それが含まれる：

1. 三つの整数、マトリックスの寸法
2. ポインタ

として整数を与えるマトリクスデータを含む配列にカーネルの起動時のパラメータは明らかに可能であり、正常に動作します。 配列へのポインタを与えることも可能です。ポインタがコピーされます。つまり、メモリ内の同じゾーンを指す別のポインタを作成します。ターゲットとする配列がホストメモリ上にある場合、エラーが発生しますが、API関数を使用してデバイス側で初期化されているため、正常に動作します。