CUDA：二維數組索引給出意想不到的結果

Question

我開始學習CUDA，並且我想寫一個簡單的程序，將一些數據複製到GPU，對其進行修改並將其傳回。我已經搜索並試圖找到我的錯誤。我很確定這個問題出現在我的內核中，但我並不完全確定哪裏出了問題。

這裏是我的內核：

__global__ void doStuff(float* data, float* result) 
{ 
    if (threadIdx.x < 9) // take the first 9 threads 
    { 
     int index = threadIdx.x; 
     result[index] = (float) index; 
    } 
} 

這裏是我main相關部分：

#include <stdlib.h> 
#include <stdio.h> 

int main(void) 
{ 
    /* 
     Setup 
    */ 
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0}; 

    float* data_array; 
    float* result_array; 

    size_t data_array_pitch, result_array_pitch; 
    int width_in_bytes = 3 * sizeof(float); 
    int height = 3; 

    /* 
     Initialize GPU arrays 
    */ 
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height); 
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height); 

    /* 
     Copy data to GPU 
    */ 
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice); 

    dim3 threads_per_block(16, 16); 
    dim3 num_blocks(1,1); 

    /* 
     Do stuff 
    */ 
    doStuff<<<num_blocks, threads_per_blocks>>>(data_array, result_array); 

    /* 
     Get the results 
    */ 
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost); 

    for (int i = 1; i <= 9; ++i) 
    { 
     printf("%f ", simple[i-1]); 
     if(!(i%3)) 
      printf("\n"); 
    } 

    return 0; 
} 

當我運行此我得到0.000000 1.000000 2.00000用於第一列和垃圾其他兩個。

Answer 1

我不確定如果你剛剛開始學習cuda，我會專注於二維數組。

也很好奇，如果您手動將代碼鍵入到問題中，因爲您定義了threads_per_block變量，但是在內核調用中使用threads_per_blocks。

無論如何，有幾個問題你的代碼：

1. 使用二維數組時，它幾乎總是需要通過音高 參數（以某種方式）內核。 cudaMallocPitch 在每行的末尾分配帶有額外填充的數組，因此 下一行開始於一個良好對齊的邊界。這通常會導致分配粒度爲128或256字節的 。所以你的第一行有3個有效的數據實體，其後有足夠的空間來填充 ，比如說256字節（等於你的音調變量）。所以我們必須改變內核調用和內核本身來解決這個問題。
2. 你的內核本質上是一個1D內核（例如，它不理解或使用threadIdx.y）。因此，推出2D網格並沒有意義。儘管在這種情況下它不會造成任何傷害，但它會造成冗餘，在其他代碼中可能會令人困惑和麻煩。

這裏呈現出一些變化，會給你預期的結果，基於上述意見的更新代碼：

#include <stdio.h> 


__global__ void doStuff(float* data, float* result, size_t dpitch, size_t rpitch, int width) 
{ 
    if (threadIdx.x < 9) // take the first 9 threads 
    { 
     int index = threadIdx.x; 
     result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index; 
    } 
} 

int main(void) 
{ 
    /* 
     Setup 
    */ 
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0}; 

    float* data_array; 
    float* result_array; 

    size_t data_array_pitch, result_array_pitch; 
    int height = 3; 
    int width = 3; 
    int width_in_bytes = width * sizeof(float); 

    /* 
     Initialize GPU arrays 
    */ 
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height); 
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height); 

    /* 
     Copy data to GPU 
    */ 
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice); 

    dim3 threads_per_block(16); 
    dim3 num_blocks(1,1); 

    /* 
     Do stuff 
    */ 
    doStuff<<<num_blocks, threads_per_block>>>(data_array, result_array, data_array_pitch, result_array_pitch, width); 

    /* 
     Get the results 
    */ 
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost); 

    for (int i = 1; i <= 9; ++i) 
    { 
     printf("%f ", simple[i-1]); 
     if(!(i%3)) 
      printf("\n"); 
    } 
    return 0; 
} 

您也可能會發現this question有趣的閱讀。

編輯：響應一個問題中的註釋：

result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index; 
       1    2      3 

爲了計算正確的元素索引到音調陣列我們必須：

1. 計算從（虛擬）行索引線程索引。我們通過對每個（非音調）行的寬度（在元素中，而不是字節中）進行線程索引的整數除法。
2. 將行索引乘以每個的寬度行。每個高度行的寬度由pitched參數給出，以字節爲單位。爲了將這個傾斜的字節參數轉換成傾斜的元素參數，我們除以每個元素的大小。然後通過將數量乘以步驟1中計算的行索引，我們現在已經索引到正確的行中。
3. 通過採用線程索引的餘數（模除法）除以寬度（在元素中）計算線程索引的（虛擬）列索引。一旦我們有了列索引（在元素中），我們將它添加到在步驟2中計算出的開始正確行索引中，以確定此線程將負責的元素。

對於相對直接的操作，以上是相當大的努力，這是爲什麼我建議專注於基本cuda概念而不是首先使用俯仰陣列的一個例子。例如，我會想象如何處理1和2D線程塊，以及1和2D網格，然後處理傾斜陣列。在某些情況下，傾斜陣列對於訪問二維數組（3D數組）來說是一種有用的性能增強器，但它們絕不是必須處理CUDA中的多維數組。

Answer 2

其實它也可以通過更換線

int width_in_bytes = 3 * sizeof(float); 

來完成：

int width_in_bytes = sizeof(float)*9; 

，因爲這是告訴cudaMemcpy2D參數多少字節從SRC複製到DST，在你要求複製3個浮點數的第一個代碼，但是你想複製的數組長度爲9，所以你需要的寬度是9個浮點數的大小。

儘管此解決方案有效，但您的代碼仍然存在一些低效率問題;例如，如果你真的想該塊的前9個線程做一些事情，在「如果」你應該以添加下列條件，（& &）

threadIdx.y==0