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; 
}