Не понимаю, почему добавление столбцов быстрее, чем строки в CUDA

Question

Я начал с CUDA и написал два ядра для эксперимента. Whey оба принимают 3 указателя на массив n * n (эмуляция матрицы) и n.

__global__
void th_single_row_add(float* a, float* b, float* c, int n) {
  int idx = blockDim.x * blockIdx.x * n + threadIdx.x * n;
  for (int i = 0; i < n; i ++) {
    if (idx + i >= n*n) return;
    c[idx + i] = a[idx + i] + b[idx + i];
  }
}

__global__
void th_single_col_add(float* a, float* b, float* c, int n) {
  int idx = blockDim.x * blockIdx.x + threadIdx.x;
  for (int i = 0; i < n; i ++) {
    int idx2 = idx + i * n;
    if (idx2 >= n*n) return;
    c[idx2] = a[idx2] + b[idx2];
  }
}

In th_single_row_add для каждого потока суммируют строки в n elemnts, в th_single_col_add для каждого столбца суммы потоков. Вот профиль на n = 1000 (1 000 000 элементов)

986.29us  th_single_row_add(float*, float*, float*, int)
372.96us  th_single_col_add(float*, float*, float*, int)

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

Answer 1

Потоки в CUDA не действуют по отдельности, они сгруппированы в основы из 32 потоков . Эти 32 потока выполняются в непрерывном режиме (обычно). Инструкция, выданная одному потоку, выдается всем 32 одновременно в одном и том же тактовом цикле.

Если эта инструкция является инструкцией, которая читает память (например), то может быть выполнено до 32 независимых операций чтения. требуется / требуется. Точные шаблоны адресов, необходимые для выполнения этих операций чтения, определяются кодом, который вы пишете. Если все эти адреса «смежны» в памяти, это будет эффективное чтение. Если эти адреса как-то «разбросаны» в памяти, это будет неэффективным чтением и будет медленнее.

Эта только что описанная базовая концепция называется «объединенным» доступом в CUDA. Ваш случай суммирования столбцов допускает объединенный доступ через деформацию, потому что адреса, сгенерированные каждым потоком в деформации, находятся в смежных столбцах, а местоположения соседствуют в памяти. Ваш случай суммирования строк ломает это. Адреса, генерируемые каждым потоком в деформации, не являются смежными (они являются «столбчатыми», отделенными друг от друга шириной массива) и, следовательно, не «объединяются».

Разница в производительности обусловленак этой разнице в эффективности доступа к памяти.

Вы можете больше узнать о объединении поведения в CUDA, изучив вводный подход к оптимизации CUDA, такой как здесь , особенно слайды 44-54.