CUDA中減少塊大小錯誤？

Question

我正在查看此sum_reduction.cu示例和tutorial，並注意到對於某些問題大小它不起作用，例如，它適用於n = 2000的問題大小，但不適用於n = 3000。顯然它總是適用於塊大小倍數的問題大小，但教程和示例代碼都沒有這樣說。問題是，這種簡化算法僅適用於某些問題大小嗎？的例子，他們選擇N = 256K，其爲偶數時，二的冪並且還多個塊大小512

對於自容納的我粘貼（的模板版本）這裏的代碼的最重要的位：

template<typename T> 
__global__ void kernelSum(const T* __restrict__ input, T* __restrict__ per_block_results, const size_t n) { 
    extern __shared__ T sdata[]; 

    size_t tid = blockIdx.x * blockDim.x + threadIdx.x; 

    // load input into __shared__ memory 
    T x = 0.0; 
    if (tid < n) { 
     x = input[tid]; 
    } 
    sdata[threadIdx.x] = x; 
    __syncthreads(); 

    // contiguous range pattern 
    for(int offset = blockDim.x/2; offset > 0; offset >>= 1) { 
     if(threadIdx.x < offset) { 
      // add a partial sum upstream to our own 
      sdata[threadIdx.x] += sdata[threadIdx.x + offset]; 
     } 
     // wait until all threads in the block have 
     // updated their partial sums 
     __syncthreads(); 
    } 

    // thread 0 writes the final result 
    if(threadIdx.x == 0) { 
     per_block_results[blockIdx.x] = sdata[0]; 
    } 
} 

，並調用內核：

// launch one kernel to compute, per-block, a partial sum 
block_sum<double> <<<num_blocks,block_size,block_size * sizeof(double)>>>(d_input, d_partial_sums_and_total, num_elements); 

// launch a single block to compute the sum of the partial sums 
block_sum<double> <<<1,num_blocks,num_blocks * sizeof(double)>>>(d_partial_sums_and_total, d_partial_sums_and_total + num_blocks, num_blocks); 

據我瞭解，如果問題大小比塊減少這種說法T x = 0.0;確保元素置零，因此應該工作更小，但它不？

更新：我很抱歉浮動/雙重的事情是一個錯字，而準備的問題，而不是真正的問題。

Answer 1

1. 你已經發布的代碼是不相符的，因爲你的模板內核 被稱爲kernelSum但要調用一些所謂的 block_sum。

2. 而且，我不相信你的模板化內核 功能的使用也可能會被正確的，因爲寫的：

   block_sum<double> <<<num_blocks,block_size,block_size * sizeof(float)>>>(d_input, d_partial_sums_and_total, num_elements); 
        ^             ^
         | these types are required to match     | 
   

   內核模板被實例化double類型。因此，期待足夠的共享內存來存儲block_sizedouble量的基礎上，這條線：

   extern __shared__ T sdata[]; 
   

   但你只有通過所需的存儲的一半：

   block_size * sizeof(float) 
   

   我相信這是要給你意外的結果。

3. 減少書面不預計該塊 尺寸爲2的冪，由於這個循環：

   // contiguous range pattern 
   for(int offset = blockDim.x/2; offset > 0; offset >>= 1) { 
   

   這是不太可能是一個問題，第一核心的呼叫，因爲你可能選擇兩種動力，每塊的線程（block_size）數量：

   block_sum<double> <<<num_blocks,block_size,... 
   

   然而，對於第二個內核調用，這將取決於WH醚num_blocks是二的冪，這取決於你的網格計算，這還沒有顯示：

   block_sum<double> <<<1,num_blocks,... 
   

4. 最後，如果num_blocks超過設備極限的第一內核啓動會失敗。這可能發生在非常大的數據集上，但可能不適用於3000的大小，這取決於您沒有顯示的網格計算。

上述第3項是一個很難滿足任意矢量大小的要求。因此，我會建議一個備用的減少策略來處理任意大小的向量。爲此，我建議你研究CUDA reduction sample code and presentation。

這裏有一個完整的計劃，主要是基於你已經顯示的代碼，有上述問題解決，似乎爲我工作了尺寸的3000：

#include <stdio.h> 
#include <stdlib.h> 
#define DSIZE 3000 
#define nTPB 256 



template<typename T> 
__global__ void block_sum(const T* __restrict__ input, T* __restrict__ per_block_results, const size_t n) { 
    extern __shared__ T sdata[]; 

    size_t tid = blockIdx.x * blockDim.x + threadIdx.x; 

    // load input into __shared__ memory 
    T x = 0.0; 
    if (tid < n) { 
     x = input[tid]; 
    } 
    sdata[threadIdx.x] = x; 
    __syncthreads(); 

    // contiguous range pattern 
    for(int offset = blockDim.x/2; offset > 0; offset >>= 1) { 
     if(threadIdx.x < offset) { 
      // add a partial sum upstream to our own 
      sdata[threadIdx.x] += sdata[threadIdx.x + offset]; 
     } 
     // wait until all threads in the block have 
     // updated their partial sums 
     __syncthreads(); 
    } 

    // thread 0 writes the final result 
    if(threadIdx.x == 0) { 
     per_block_results[blockIdx.x] = sdata[0]; 
    } 
} 


int main(){ 

    double *d_input, *d_partial_sums_and_total, *h_input, *h_partial_sums_and_total; 

    int num_elements=DSIZE; 
    int block_size = nTPB; 
    int num_blocks = (num_elements + block_size -1)/block_size; 
    // bump num_blocks up to the next power of 2 
    int done = 0; 
    int test_val = 1; 
    while (!done){ 
    if (test_val >= num_blocks){ 
    num_blocks = test_val; 
    done = 1;} 
    else test_val *= 2; 
    if (test_val > 65535) {printf("blocks failure\n"); exit(1);} 
    } 


    h_input = (double *)malloc(num_elements * sizeof(double)); 
    h_partial_sums_and_total = (double *)malloc((num_blocks+1)*sizeof(double)); 

    cudaMalloc((void **)&d_input, num_elements * sizeof(double)); 
    cudaMalloc((void **)&d_partial_sums_and_total, (num_blocks+1)*sizeof(double)); 

    double h_result = 0.0; 
    for (int i = 0; i < num_elements; i++) { 
    h_input[i] = rand()/(double)RAND_MAX; 
    h_result += h_input[i];} 

    cudaMemcpy(d_input, h_input, num_elements*sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemset(d_partial_sums_and_total, 0, (num_blocks+1)*sizeof(double)); 

// launch one kernel to compute, per-block, a partial sum 
    block_sum<double> <<<num_blocks,block_size,block_size * sizeof(double)>>>(d_input, d_partial_sums_and_total, num_elements); 

// launch a single block to compute the sum of the partial sums 
    block_sum<double> <<<1,num_blocks,num_blocks * sizeof(double)>>>(d_partial_sums_and_total, d_partial_sums_and_total + num_blocks, num_blocks); 

    cudaMemcpy(h_partial_sums_and_total, d_partial_sums_and_total, (num_blocks+1)*sizeof(double), cudaMemcpyDeviceToHost); 

    printf("host result = %lf\n", h_result); 
    printf("device result = %lf\n", h_partial_sums_and_total[num_blocks]); 
} 

爲了簡潔/可讀性，我在上面的代碼中省去了錯誤檢查。在遇到cuda代碼時遇到困難時，您應該始終使用proper cuda error checking。

另外，未來，如果您發佈完整的代碼，以證明您正在做的事情，您將使其他人更容易幫助您，如上所述。