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

и grid и blocks (blockDim.x равен ширине матрицы, т.е. 64):

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

Я не могу этого достичь. Может ли кто-нибудь указать, что с этим не так? Должен ли я использовать пошаговый индекс или смещение?

Answer 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). Далее у вас нет синхронизации, и вы используете эту переменную во втором назначении. Это будет иметь неопределенный результат, потому что потоки из первой деформации читают из этой переменной, а потоки из второй деформации все еще пытаются что-то там записать. Вы должны изменить 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;
}

Answer 2

VAR не должен был быть объявлен как общий, потому что в текущей форме все потоки перебирают данные друг друга при загрузке из глобальной памяти: VAR = data_old[glob_index];,

У вас также есть доступ за пределами доступа при доступе VAR2[threadIdx.x + 1]таким образом, ваше ядро ​​никогда не завершает работу (в зависимости от вычислительных возможностей устройства - устройства 1.x не проверяли доступ к общей памяти так строго).

Вы могли обнаружить последнее, проверив коды ошибок всех вызовов функций CUDA на наличие ошибок.