CUDA還原

Question

我目前使用下面的還原功能使用CUDA來概括所有元素的數組中的多個陣列：

__global__ void reduceSum(int *input, int *input2, int *input3, int *outdata, int size){ 
    extern __shared__ int sdata[]; 

    unsigned int tID = threadIdx.x; 
    unsigned int i = tID + blockIdx.x * (blockDim.x * 2); 
    sdata[tID] = input[i] + input[i + blockDim.x]; 
    __syncthreads(); 

    for (unsigned int stride = blockDim.x/2; stride > 32; stride >>= 1) 
    { 
     if (tID < stride) 
     { 
      sdata[tID] += sdata[tID + stride]; 
     } 
     __syncthreads(); 
    } 

    if (tID < 32){ warpReduce(sdata, tID); } 

    if (tID == 0) 
    { 
     outdata[blockIdx.x] = sdata[0]; 
    } 
} 

然而，正如你可以從函數的參數看，我會喜歡能夠在一個簡化函數內求和三個單獨的數組。現在很明顯，一個簡單的方法是啓動內核三次，每次傳遞一個不同的數組，這當然會起作用。我只是把它作爲一個測試內核來寫，但真正的內核最終會得到一個結構數組，而且我需要爲每個結構的所有X，Y和Z值執行一個加法，這就是爲什麼我需要在一個內核中總結它們。

我已經initalised併爲所有三個陣列

int test[1000]; 
    std::fill_n(test, 1000, 1); 
    int *d_test; 

    int test2[1000]; 
    std::fill_n(test2, 1000, 2); 
    int *d_test2; 

    int test3[1000]; 
    std::fill_n(test3, 1000, 3); 
    int *d_test3; 

    cudaMalloc((void**)&d_test, 1000 * sizeof(int)); 
    cudaMalloc((void**)&d_test2, 1000 * sizeof(int)); 
    cudaMalloc((void**)&d_test3, 1000 * sizeof(int)); 

我不確定什麼網格和塊尺寸，我應該使用這種內核的，我不完全知道如何修改減少迴路分配的內存將數據放置，因爲我想它，即 輸出數組：

Block 1 Result|Block 2 Result|Block 3 Result|Block 4 Result|Block 5 Result|Block 6 Result| 

     Test Array 1 Sums    Test Array 2 Sums   Test Array 3 Sums   

我希望是有道理的。還是有更好的方法只有一個約簡函數，但能夠返回Struct.X，Struct.Y或struct.Z的總和？

這裏的結構：

template <typename T> 
struct planet { 
    T x, y, z; 
    T vx, vy, vz; 
    T mass; 
}; 

我需要添加了所有的VX和存儲，所有的VY並儲存起來，所有的VZ和存儲。

Answer 1

或者還有更好的方法只有一個還原函數，但能夠返回Struct.X，Struct.Y或struct.Z的總和嗎？

通常加速計算的主要焦點是速度。 GPU代碼的速度（性能）通常在很大程度上取決於數據存儲和訪問模式。因此，儘管如您在您的問題中指出的那樣，我們可以通過多種方式實現解決方案，讓我們專注於應該相對較快的事情。

像這樣的減少算術/操作強度不大，所以我們對性能的關注主要圍繞數據存儲以有效訪問。當訪問全局內存時，GPU通常會以大塊的形式進行操作 - 32字節或128字節的塊。爲了有效利用內存子系統，我們希望在每個請求中使用所請求的所有32或128個字節。

但你的結構隱含的數據存儲模式：

template <typename T> 
struct planet { 
    T x, y, z; 
    T vx, vy, vz; 
    T mass; 
}; 

相當多的規則了這一點。對於此問題，您關心的是vx,vy和vz。那些3項應是給定結構（元件）內連續的，但在這些結構的陣列，他們將通過必要的存儲對於其它結構的物品分開，至少：

planet0:  T x 
       T y 
       T z    --------------- 
       T vx  <--   ^
       T vy  <--   | 
       T vz  <--  32-byte read 
       T mass     | 
planet1:  T x      | 
       T y      v 
       T z    --------------- 
       T vx  <-- 
       T vy  <-- 
       T vz  <-- 
       T mass 
planet2:  T x 
       T y 
       T z 
       T vx  <-- 
       T vy  <-- 
       T vz  <-- 
       T mass 

（爲求例如，假設是Tfloat）

此指出陣列結構（AOS）存儲格式在GPU的一個關鍵缺點。由於GPU的訪問粒度（32字節），因此從連續結構訪問相同的元素是不合適的。在這種情況下，性能通常的建議是將AOS存儲向SOA轉換（陣列結構）：

template <typename T> 
struct planets { 
    T x[N], y[N], z[N]; 
    T vx[N], vy[N], vz[N]; 
    T mass[N]; 
}; 

以上只是一個可能的例子，可能不是你真正會用什麼作爲結構將成爲目的不大，因爲我們只有一個行星結構N。關鍵是，現在當我訪問vx連續行星，各個vx元件都在相鄰的存儲器，所以32字節的讀出給我32字節值得vx數據的，沒有浪費的或不使用的元件。

利用這樣的變換，還原問題再次變得相對簡單，從代碼的組織的觀點來看。您可以基本上使用與單個數組縮減代碼相同的內容，可以連續調用3次，也可以直接對內核代碼進行擴展，以獨立處理所有3個數組。 A「3合1」的內核可能是這個樣子：

template <typename T> 
__global__ void reduceSum(T *input_vx, T *input_vy, T *input_vz, T *outdata_vx, T *outdata_vy, T *outdata_vz, int size){ 
    extern __shared__ T sdata[]; 

    const int VX = 0; 
    const int VY = blockDim.x; 
    const int VZ = 2*blockDim.x; 

    unsigned int tID = threadIdx.x; 
    unsigned int i = tID + blockIdx.x * (blockDim.x * 2); 
    sdata[tID+VX] = input_vx[i] + input_vx[i + blockDim.x]; 
    sdata[tID+VY] = input_vy[i] + input_vy[i + blockDim.x]; 
    sdata[tID+VZ] = input_vz[i] + input_vz[i + blockDim.x]; 
    __syncthreads(); 

    for (unsigned int stride = blockDim.x/2; stride > 32; stride >>= 1) 
    { 
     if (tID < stride) 
     { 
      sdata[tID+VX] += sdata[tID+VX + stride]; 
      sdata[tID+VY] += sdata[tID+VY + stride]; 
      sdata[tID+VZ] += sdata[tID+VZ + stride]; 
     } 
     __syncthreads(); 
    } 

    if (tID < 32){ warpReduce(sdata+VX, tID); } 
    if (tID < 32){ warpReduce(sdata+VY, tID); } 
    if (tID < 32){ warpReduce(sdata+VZ, tID); } 

    if (tID == 0) 
    { 
     outdata_vx[blockIdx.x] = sdata[VX]; 
     outdata_vy[blockIdx.x] = sdata[VY]; 
     outdata_vz[blockIdx.x] = sdata[VZ]; 
    } 
} 

（編碼的瀏覽器 - 沒有測試 - 只是你的表現出什麼「基準內核」的擴展名）

的高於AoS - > SoA數據轉換可能會在代碼中的其他地方具有性能優勢。由於所提出的內核將一次處理3個數組，網格和塊尺寸應完全相同，你會在單陣列情況下用什麼供您參考內核。共享內存存儲將需要增加（三倍）每塊。

Answer 2

羅伯特Crovella給了一個很好的答案，突出AOS的重要性 - >往往提高了對GPU性能SoA的佈局轉換，我只是想提出一箇中間地帶，可能是更方便。 CUDA語言僅爲您描述的目的提供了幾種矢量類型（請參閱this section of the CUDA programming guide）。

例如，CUDA定義INT3，存儲3點的整數數據類型。

struct int3 
{ 
    int x; int y; int z; 
}; 

類似類型的花車，字符，雙打等存在什麼是好的關於這些數據類型是，他們可以用一條指令，這可能會給你一個小的性能提升加載。有關這方面的討論，請參閱this NVIDIA blog post。對於這種情況，它也是一種更「自然」的數據類型，它可能會使代碼的其他部分更易於使用。例如，您可以定義：

struct planets { 
    float3 position[N]; 
    float3 velocity[N]; 
    int mass[N]; 
}; 

使用此數據類型的還原內核可能看起來像這樣（改編自Robert's）。

__inline__ __device__ void SumInt3(int3 const & input1, int3 const & input2, int3 & result) 
{ 
    result.x = input1.x + input2.x; 
    result.y = input1.y + input2.y; 
    result.z = input1.z + input2.z; 
} 

__inline__ __device__ void WarpReduceInt3(int3 const & input, int3 & output, unsigned int const tID) 
{ 
    output.x = WarpReduce(input.x, tID); 
    output.y = WarpReduce(input.y, tID); 
    output.z = WarpReduce(input.z, tID);  
} 

__global__ void reduceSum(int3 * inputData, int3 * output, int size){ 
    extern __shared__ int3 sdata[]; 

    int3 temp; 

    unsigned int tID = threadIdx.x; 
    unsigned int i = tID + blockIdx.x * (blockDim.x * 2); 

    // Load and sum two integer triplets, store the answer in temp. 
    SumInt3(input[i], input[i + blockDim.x], temp); 

    // Write the temporary answer to shared memory. 
    sData[tID] = temp; 

    __syncthreads(); 

    for (unsigned int stride = blockDim.x/2; stride > 32; stride >>= 1) 
    { 
     if (tID < stride) 
     { 
      SumInt3(sdata[tID], sdata[tID + stride], temp); 
      sData[tID] = temp; 
     } 
     __syncthreads(); 
    } 

    // Sum the intermediate results accross a warp. 
    // No need to write the answer to shared memory, 
    // as only the contribution from tID == 0 will matter. 
    if (tID < 32) 
    { 
     WarpReduceInt3(sdata[tID], tID, temp); 
    } 

    if (tID == 0) 
    { 
     output[blockIdx.x] = temp; 
    } 
}