CUDA：在共享內存上移位數組

Question

我正在嘗試將平鋪的2D矩陣加載到共享內存中，將數據沿x移動，寫回到全局內存也隨y移動。輸入數據因此沿着x和y移動。我有：

__global__ void test_shift(float *data_old, float *data_new) 

{ 

uint glob_index = threadIdx.x + blockIdx.y*blockDim.x; 

__shared__ float VAR; 
__shared__ float VAR2[NUM_THREADS]; 

// load from global to shared 

VAR = data_old[glob_index]; 

// do some stuff on VAR 

if (threadIdx.x < NUM_THREADS - 1) 
{ 
VAR2[threadIdx.x + 1] = VAR; // shift (+1) along x 
} 

__syncthreads(); 

// write to global memory 

if (threadIdx.y < ny - 1) 
{ 
glob_index = threadIdx.x + (blockIdx.y + 1)*blockDim.x; // redefine glob_index to shift along y (+1) 
data_new[glob_index] = VAR2[threadIdx.x]; 
} 

到內核的呼叫：

test_shift <<< grid, block >>> (data_old, data_new); 

和網格和塊（blockDim.x等於矩陣寬度，即64）：

dim3 block(NUM_THREADS, 1); 
dim3 grid(1, ny); 

我無法實現它。有人能指出這有什麼問題嗎？我應該使用一個跨度索引還是偏移量？

Answer 1

VAR不應該被聲明爲共享，因爲在當前表單中，當您從全局內存加載時，所有線程都會在對方的數據上塗寫：VAR = data_old[glob_index];。

當你訪問VAR2[threadIdx.x + 1]時，你也有一個越界訪問，所以你的內核永遠不會完成（取決於設備的計算能力 - 1.x設備沒有嚴格檢查共享內存訪問）。

您可能通過檢查所有對CUDA函數的調用的返回碼來檢測錯誤。

Answer 2

共享變量很好地由單個塊中的所有線程共享。這意味着你沒有blockDim.y完成共享變量，但是每個塊只有一個單獨的壓縮。

uint glob_index = threadIdx.x + blockIdx.y*blockDim.x; 

__shared__ float VAR; 
__shared__ float VAR2[NUM_THREADS]; 
VAR = data_old[glob_index]; 

if (threadIdx.x < NUM_THREADS - 1) 
{ 
    VAR2[threadIdx.x + 1] = VAR; // shift (+1) along x 
} 

這指示塊中的所有線程將數據寫入單個變量（VAR）。接下來你沒有同步，你在第二個任務中使用這個變量。這將有未定義的結果，因爲來自第一個warp的線程正在讀取這個變量，第二個warp中的線程仍然試圖在那裏寫入東西。 您應該將VAR更改爲本地，或者爲塊中的所有線程創建一個共享內存變量數組。

if (threadIdx.y < ny - 1) 
{ 
    glob_index = threadIdx.x + (blockIdx.y + 1)*blockDim.x; 
    data_new[glob_index] = VAR2[threadIdx.x]; 
} 

在VAR2 [0]中，你仍然有一些垃圾（你從來沒有寫過）。 threadIdx.y在您的塊中始終爲零。

並避免使用提示。他們有（或曾經有過）一些性能問題。

其實，對於這樣簡單的任務，你不需要使用共享內存

__global__ void test_shift(float *data_old, float *data_new) 
{ 

int glob_index = threadIdx.x + blockIdx.y*blockDim.x; 

float VAR; 

// load from global to local 
VAR = data_old[glob_index]; 

int glob_index_new; 
// calculate only if we are going to output something 
if ((blockIdx.y < gridDim.y - 1) && (threadIdx.x < blockDim.x - 1)) 
{ 
    glob_index_new = threadIdx.x + 1 + (blockIdx.y + 1)*blockDim.x; 

    // do some stuff on VAR 
} else // just write 0.0 to remove garbage 
{ 
    glob_index_new = ((blockIdx.y == gridDim.y - 1) && (threadIdx.x == blockDim.x - 1)) ? 0 : ((blockIdx.y == gridDim.y - 1) ? threadIdx.x : (blockIdx.y)*blockDim.x); 
    VAR = 0.0; 
} 

// write to global memory 

data_new[glob_index_new] = VAR; 
}