Какова взаимосвязь между размерным характером потоков и размерами самих данных в CUDA?

Question

Я с самого начала читал, что использование 2D-блока потоков - это самый простой способ работы с 2D-набором данных. Я пытаюсь реализовать следующие матричные операции в последовательности:

1. Поменять элементы на нечетные и четные позиции каждой строки в матрице

   1 2         2 1
   3 4 becomes 4 3
   

2. Отражение элементов матрицы по главной диагонали

   2 1         2 4
   4 3 becomes 1 3
   

Для реализации я написал следующее ядро:

__global__ void swap_and_reflect(float *d_input, float *d_output, int M, int N)
{
    int j = threadIdx.x;
    int i = threadIdx.y;
    for(int t=0;t<M*N;t++)
      d_output[t] = d_input[t];
    float temp = 0.0;
    if (j%2 == 0){
            temp = d_output[j];
            d_output[j] = d_output[j+1];
            d_output[j+1] = temp;         
    }
    __syncthreads(); // Wait for swap to complete
    if (i!=j){
        temp = d_output[i];
        d_output[i] = d_output[j];
        d_output[j] = temp;       
    }
}

Отражение не случилось, как и ожидалось. Но в этот момент я часто путаюсь с 2D-структурой исполняющих потоков с 2D-структурой самой матрицы.

Не могли бы вы исправить мое понимание многомерного расположения потоков и того, как оно соотносится с размерностью самих данных? Я полагаю, что это причина того, почему у меня неправильная часть отражения.

Любые указатели / ресурсы, которые могли бы помочь мне визуализировать / понять это правильно, были бы огромной помощью.

Спасибо за чтение.

Answer 1

Индексы потоков располагаются в вашем гипотетическом блоке 4x4 в парах (x, y) как

(0,0)  (0,1)
(1,0)  (1,1)

, а порядок:

thread ID       (x,y) pair
---------       ----------
0               (0,0)
1               (1,0)
2               (0,1)
3               (1,1)

Вам необходимо выбрать упорядочение вашего массива в памяти , а затем соответственно изменить ядро, например:

if (i!=j){
    temp = d_output[i+2*j];
    d_output[i+2*j] = d_output[j+2*i];
    d_output[j+2*i] = temp;       
}