是否可以在CUDA設備上定義全局內存數組？

我tr分配一個cuda全局內存數組。我的內核總結如下：是否可以在CUDA設備上定義全局內存數組？

__device__ float R_d = 0; 
__global__ void perform_summation(float* A, int N){ 
    int idx = blockDim.x*blockIdx.x+threadIdx.x; 
    extern __shared__ float sharedArray []; 
    float result[]; //THIS IS THE THING i TRIED TO CREATE 
    if(idx < N){ 
     sharedArray[threadIdx.x] = A[idx]; 
     // }else{ 
     //  sharedArray[threadIdx.x] = 0 ; 
     // } 

     for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { 
      __syncthreads(); 
      if(threadIdx.x % (2*stride) == 0){ 
       sharedArray[threadIdx.x]+=sharedArray[threadIdx.x+stride]; 
      } 
     } 
    } 
    if(idx % blockDim.x == 0){ 
//  R_d += sharedArray[threadIdx.x]; 
     result[blockIdx.x] = sharedArray[threadIdx.x]; 
    } 
    for (int i = 0; i < gridDim.x; ++i) { 
     R_d += result[i]; 
    } 
}

作爲總結y內核需要一個數組並通過映射縮減方法找到元素的總和。每個塊將相關元素放入共享內存中，然後彙總內部的所有數據，而不是將結果放到我嘗試創建的全局數組中。在最後，我將總結全局數組的所有數字以找到最後一個答案。

作爲第一種方法，我沒有使用全局數組來收集每個塊的結果，我只是將塊的結果總和到變量R_d中，但它不起作用，並且只顯示來自最後一個塊的值爲結果。我想因爲我沒有同步。塊之間的最後一個塊將覆蓋最後的所有值。這裏是我在第一次嘗試在內核的末尾完成的工作

f(idx < N){ 
     sharedArray[threadIdx.x] = A[idx]; 
     // }else{ 
     //  sharedArray[threadIdx.x] = 0 ; 
     // } 

     for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { 
      __syncthreads(); 
      if(threadIdx.x % (2*stride) == 0){ 
       sharedArray[threadIdx.x]+=sharedArray[threadIdx.x+stride]; 
      } 
     } 

     if(threadIdx.x == 0){ 
      R_d += sharedArray[threadIdx.x]; 
     } 
    }

所以我實際上有兩個問題。如何在設備內存中爲我提出的第一個解決方案定義一個全局內存陣列，並且第二個解決方案是否有任何解決方案只使用R_d變量？

來源

2013-04-14 erogol

R_d甚至沒有在任何地方定義。如果你要發佈代碼，至少要確保它能夠編譯。就目前而言，這個代碼只是無稽之談。 – talonmies

你是對的我只是更新了上面的代碼 – erogol

您可以通過cudaMalloc:

cudaMalloc((void **)&ptr, size);

分配在全局設備存儲器陣列，但你不想做那裏面的內核，則調用內核和指針傳遞給內核之前做到這一點。

至於減少，看看these nVidia slides，它解釋得很好。基本上，它取決於你使用了多少塊和線程。可以說有幾個街區。因此，在共享存儲器定義一個數組：

__shared__ float cache[BLOCK_THREADS];

共享存儲器分配給每個塊的，所以我們在每個塊中在cache的值相加至所述第一元件。

__syncthreads(); 

int step = (BLOCK_THREADS >> 1); //the same result as BLOCK_THREADS/2 
while(step > 0) { 
    if (threadInBlock < step) { 
     cache[threadInBlock] += cache[threadInBlock + step]; 
    } 
    __syncthreads(); 
    step = (step >> 1); 
}

因此，這將每個塊中的所有元素總和爲cache[0]。現在我們可以再次使用約簡，或者我們可以用原子操作將每個塊中的所有和相加。如果塊的數量明顯少於每個塊的線程數量，那麼這將是可以的。

__syncthreads(); 
if (threadInBlock == 0) { 
    atomicAdd(result, cache[0]); 
}

請注意result是一個指向全局內存中單個值的指針。還要注意，只有當BLOCK_THREADS是2的冪 - 這是很常見的，因爲每個塊的線程數應該是32的倍數（與經線對齊）。

來源

2013-04-14 12:55:42

是否可以在CUDA設備上定義全局內存數組？

回答

相關問題