Kopieren einer Struktur in den Gerätespeicher CUDA

Question

Ich bin neu in CUDA und gehe durch die CUDA Toolkit-Dokumentation. Dort fand ich ein Beispiel, bei dem die Matrixmultiplikation Shared Memory verwendet. Hier werden beim Kopieren der Matrix-Struktur vom Host-Speicher in den Gerätespeicher nur die Datenelemente kopiert. Was ich nicht verstehen kann ist, wie andere Variablen in den Gerätespeicher kopiert werden.

Matrix-Struktur ist wie folgt

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

Dann ist hier das Codebeispiel, wo die Datenübertragung geschieht

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); 
} 

Hier ist, was ich nicht verstehe, ist, wie die Breite, schreiten und Höhe kopiert wird Gerätespeicher. Denn hier ist cudaMalloc und cudaMemcpy nur für die Elemente. Gibt es etwas, was ich verpasst habe, dies zu verstehen?

Die Kernal-Codes

__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; 
} 

Matrixmultiplikation kernal Code

__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

Für diejenigen, die über sich fragen, der Beispielcode wir sprechen, ist hier auf Nvidias CUDA-Toolkit-Dokumentation, die in dem Speicher Thema geteilt: CUDA C Programming guide : Shared memory

Also, warum funktioniert dieses Beispiel? Ja, nur das Array "elements" wird auf der Geräteseite gesendet, indem die Funktionen cudaMalloc und cudaMemcpy verwendet werden. Ja, die Matrix-Dimensionen werden innerhalb des Kernels auf der Geräteseite verwendet, ohne explizit mit cudaMemcpy in den Gerätespeicher kopiert zu werden.

Sie müssen Arrays und Parameter nicht auf die gleiche Weise betrachten. Lassen Sie mich erklären, wie diese Werte an den Kernel gesendet werden.

1. Wir erklären die Matrix auf CPU-Seite, alle Mitglieder sind nicht initialisierten
2. wir die Dimensionen zuweisen, ist der Zeiger noch nicht initialisierten
3. Wir zuzuteilen und kopieren Speicher auf Geräteseite mit API-Funktionen, der Zeiger initialisiert wird aber es zielt auf Gerätespeicher und kann nicht wie ein normales Host-Array verwendet werden
4. Wir geben die Matrix als Parameter für den Kernel. NICHT nach Zeiger, sondern nach Wert.

Und das ist der Trick. Die vollständige Struktur-Instanz wird als Parameter übergeben, und es enthält:

1. drei ganzen Zahlen, die Dimensionen der Matrix
2. Ein Zeiger auf das Array, das die Matrixdaten

ganze Zahlen geben enthält als Parameter im Kernel-Start ist offensichtlich möglich und es funktioniert gut. Es ist auch möglich, den Zeiger auf ein Array zu geben: Der Zeiger wird kopiert, was bedeutet, dass wir einen weiteren Zeiger erstellen, der auf dieselbe Zone im Speicher zeigt. Wenn das Array, auf das wir ausgerichtet waren, auf dem Host-Speicher war, würde dies zu Fehlern führen, aber da es auf der Geräteseite mit API-Funktionen initialisiert wurde, funktioniert es gut.