0

我正在尝试将展平的 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);

以及grid和blocks（blockDim.x等于矩阵宽度，即64）：

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

我无法实现它。有人可以指出这有什么问题吗？我应该使用跨步索引还是偏移量？

4

2 回答 2

1

VAR不应将其声明为共享，因为在当前形式中，当您从全局内存加载时，所有线程都会在彼此的数据上乱涂乱画：VAR = data_old[glob_index];.

您在访问时也有越界访问VAR2[threadIdx.x + 1]，因此您的内核永远不会完成（取决于设备的计算能力 - 1.x 设备没有严格检查共享内存访问）。

您可以通过检查对 CUDA 函数的所有调用的返回代码是否有错误来检测后者。

于 2012-11-30T11:05:52.820 回答
1

共享变量由单个块中的所有线程共享。这意味着您没有共享变量的 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 在您的块中始终为零。

并避免使用 uint。他们有（或曾经有）一些性能问题。

实际上，对于这样简单的任务，您不需要使用共享内存

__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;
}
于 2012-11-30T11:17:50.763 回答