Как избежать несвязанных доступов в матричном умножении ядра CUDA? - PullRequest
Новая
       6

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

0 голосов
/ 02 января 2019

Я изучаю CUDA с книгой ' Программирование массово параллельных процессоров '.Практическая проблема из главы 5 смущает меня:

Для умножения плиточной матрицы вне возможного диапазона значений для BLOCK_SIZE, для каких значений BLOCK_SIZE ядро ​​полностью избежит не слитых обращений к глобальной памяти?(вам нужно только рассмотреть квадратные блоки)

Насколько я понимаю, BLOCK_SIZE мало что делает для объединения памяти.До тех пор, пока потоки внутри единой деформации обращаются к последовательным элементам, у нас будет объединенный доступ.Я не мог понять, где ядро ​​имеет неразделенный доступ к глобальной памяти.Любые намеки от вас, ребята?

Вот исходные коды ядра: 

#define COMMON_WIDTH 512
#define ROW_LEFT 500 
#define COL_RIGHT 250
#define K 1000
#define TILE_WIDTH 32
__device__ int D_ROW_LEFT = ROW_LEFT;
__device__ int D_COL_RIGHT = COL_RIGHT;
__device__ int D_K = K;
.....
__global__
void MatrixMatrixMultTiled(float *matrixLeft, float *matrixRight, float *output){
    __shared__  float sMatrixLeft[TILE_WIDTH][TILE_WIDTH];
    __shared__  float sMatrixRight[TILE_WIDTH][TILE_WIDTH];  
   int bx = blockIdx.x; int by = blockIdx.y;
   int tx = threadIdx.x; int ty = threadIdx.y;
   int col = bx * TILE_WIDTH + tx;
   int row = by * TILE_WIDTH + ty;
   float value = 0;
   for (int i = 0; i < ceil(D_K/(float)TILE_WIDTH); ++i){
       if (row < D_ROW_LEFT && row * D_K + i * TILE_WIDTH  +tx < D_K){
        sMatrixLeft[ty][tx]  = matrixLeft[row * D_K + i * TILE_WIDTH  +tx];
       }
       if (col < D_COL_RIGHT && (ty + i * TILE_WIDTH) * D_COL_RIGHT  + col < D_K ){
        sMatrixRight[ty][tx] = matrixRight[(ty + i * TILE_WIDTH) * D_COL_RIGHT  + col];
       }
       __syncthreads();
       for (int j = 0; j < TILE_WIDTH; j++){
           value += sMatrixLeft[ty][j] * sMatrixRight[j][tx]; 
       }
       __syncthreads();
   }
   if (row < D_ROW_LEFT && col < D_COL_RIGHT ){
        output[row * D_COL_RIGHT + col] = value;
       }
}

1 Ответ

0 голосов
/ 02 января 2019

Ваш вопрос является неполным, так как код, который вы разместили, не имеет никакого отношения к BLOCK_SIZE, и это, по крайней мере, очень важно для вопроса, поставленного в книге.В более общем плане вопросы, которые ставят ядро ​​без конфигурации запуска, часто являются неполными, поскольку конфигурация запуска часто имеет отношение как к правильности, так и к поведению ядра.

Я не перечитывал эту частькниги прямо сейчас.Однако я предполагаю, что конфигурация запуска ядра включает в себя измерение блока, которое выглядит примерно так: (эта информация отсутствует в вашем вопросе, но должна была быть включена, на мой взгляд, для разумного вопроса) 

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
dim3 dimGrid(...,...);

И я предполагаю, что запуск ядра задается чем-то вроде: 

MatrixMatrixMultTiled<<<dimGrid, dimBlock>>>(...);

Ваше утверждение: «Пока потоки в рамках единой деформации обращаются к последовательным элементам, у нас будет объединенный доступ».это разумное рабочее определение.Покажем, что это нарушается для некоторых вариантов BLOCK_SIZE, учитывая вышеизложенные допущения, чтобы покрыть пробелы в вашем неполном вопросе.

Объединенный доступ - это термин, который применяется только к глобальному доступу к памяти.Поэтому мы будем игнорировать доступ к общей памяти.В этом обсуждении мы также будем игнорировать доступ к переменным __device__, таким как D_ROW_LEFT.(Доступ к этим переменным, по-видимому, равномерный . Мы можем поспорить о том, представляет ли это объединенный доступ. Я бы сказал, что это действительно объединенный доступ, но нам не нужно распаковывать его здесь.) Поэтому мыосталось только 3 "точки доступа": 

matrixLeft[row * D_K + i * TILE_WIDTH  +tx];
matrixRight[(ty + i * TILE_WIDTH) * D_COL_RIGHT  + col];
output[row * D_COL_RIGHT + col]

Теперь, чтобы выбрать пример, давайте предположим, что BLOCK_SIZE равно 16. Будет ли любая из вышеперечисленных точек доступа нарушать потоки вашего оператора подряд в рамках единого доступа деформации подрядэлементы "?

Начнем с блока (0,0).Следовательно, row равно threadIdx.y и col равно threadIdx.x.Давайте рассмотрим первый перекос в этом блоке.Поэтому первые 16 потоков в этой деформации будут иметь значение threadIdx.y, равное 0, а их значения threadIdx.x будут увеличиваться с 0..15.Аналогично, вторые 16 потоков в этой деформации будут иметь значение threadIdx.y, равное 1, и их значения threadIdx.x будут увеличиваться с 0..15.

Теперь давайте вычислим фактический индекс, сгенерированный для первого доступа.точка выше, через варп.Давайте предположим, что мы находимся на первой итерации цикла, поэтому i равно нулю.Поэтому это: 

matrixLeft[row * D_K + i * TILE_WIDTH  +tx];

сводится к: 

matrixLeft[threadIdx.y * D_K + threadIdx.x];

D_K - это просто копия устройства переменной K, равной 1000. Теперь давайте оценим уменьшенный индексВыражение выше через нашу выбранную деформацию (0) в нашем выбранном блоке (0,0): 

warp lane:    0  1  2  3  4  5  6  .. 15     16   17   18 .. 31
threadIdx.x   0  1  2  3  4  5  6     15      0    1    2    15
threadIdx.y   0  0  0  0  0  0  0      0      1    1    1     1
index:        0  1  2  3  4  5  6     15   1000 1001 1002  1015

Следовательно, сгенерированный индексный паттерн здесь показывает разрыв между 16-м и 17-м потоком в основе и доступшаблон не соответствует вашему ранее указанному условию:

"потоки в последовательных элементах доступа одной основы"

, и у нас нет объединенного доступа в этом случае (по крайней мере, для float количеств).

...