Процедура перемножения матриц на GPU SM4

Рассмотрим конкретную программу, реализующую перемножение матриц на GPU SM4. Эта программа может быть откомпилирована любыми средствами, поддерживающими библиотеки CUDA. Детали реализации алгоритма обсуждаются в комментариях.

 

// Файл matrixMul_kernel.cu

 

#ifndef _MATRIXMUL_KERNEL_H_

#define _MATRIXMUL_KERNEL_H_

 

#include < stdio.h>

#include " matrixMul.h"

 

/* Следующая процедура имеет модификатор __global__, т.е. является вычислительным ядром. Это ядро исполняет­ся одновременно на каждом мультипроцессоре и каждым потоком. Но всё-таки конкретные вычислительные процес­сы, исполняемые в различных «связках» и потоках, разли­чаются значениями таких системных переменных, как blockIdx и threadIdx. Это позволяет нам «привязать» конкретные ячейки матриц к конкретным потокам.

Параметрами процедуры являются указатели (float*) на области оперативной памяти компьютера, в которых распо­ложены матрицы A, B и C, а также целочисленные (int) значения wA и wB, задающие ширину входных матриц в единицах размера блока Block_Size */

 

__global__ void

matrixMul (float* C, float* A, float* B, int wA, int wB)

{

/* Присвоение локальным переменным значений координат текущей «связки» потоков */

int block_x = blockIdx.x;

int block_y = blockIdx.y;

 

/* Определение индексов (i, k), соответствующих конкретному потоку в составе «связки» */

int i = threadIdx.y;

int k = threadIdx.x;

 

/* Для понимания дальнейшего необходимо учесть, что в глобальной памяти графического процессора (Global memory на рис. 9.2) элементы исходных матриц A и B хранятся последовательно, в форме линейного массива, как это показано следующими формулами:

.

Как можно увидеть из рис. 9.4, элементы матрицы A, обрабатываемые одной «связкой», расположены в линейном массиве последовательно. Если нумеровать все элементы линейного массива подряд с нуля, то номер первого элемента последовательности, с которой работает заданная «связка» -

.

Вводим целочисленную константу, равную этому начальному номеру: */

 

int a_First = wA * Block_Size * block_y;

 

/* Как, опять же, видно из рис. 9.4, последователь­ность элементов матрицы A, обрабатываемая выделенной «связкой», проходит через несколько блоков размера Block_Size, которые загружаются в разделяемую память поочерёдно. Для организации цикла, внутри которого будут загружаться в память и обрабатываться эти блоки, нам потребуется номер 1-ого элемента в последнем из блоков, который вычисляется по формуле

.

Определяем соответствующую целочисленную константу: */

 

int a_Last = a_First + (wA – 1) * BLOCK_SIZE;

 

/* Для организации этого же цикла потребуется шаг, задающий изменение 1-ого номера (т.е. «левого верхнего» элемента в блоке) при переходе от текущего блока к следующему */

int a_Step = Block_Size;

 

/* Индекс первого блока матрицы B, обрабатываемого текущей связкой

int b_First = Block_Size * block_x;

 

/* Шаг итераций «сверху вниз» (см. рис. 9.3) по блокам матрицы B */

int b_Step = Block_Size * wB;

 

/* Величины шагов при переходах от одних блоков к другим в матрицах A и B проиллюстрированы на рис. 9.5.

Рис. 9.5. Переходы к следующим блокам при умножении матриц */

// Вводим переменную, в которую будет записываться
//значение искомого элемента матрицы Csub (рис. 9.4-5)

float Csub = 0;

 

// Внешний цикл по всем блокам матриц A и B

for (int a = a_First, b = b_First;

a < = a_Last;

a += a_Step, b += b_Step) {

 

/* Создаём в разделяемой памяти массив, куда из гло­бальной видеопамяти будет скопирован очередной блок матрицы A. Модификатор __shared__ как раз означает, что массив As будет создан в «быстрой» разделяемой памяти */

__shared__ float As[Block_Size][Block_Size];

 

/* Такой же массив создается для хранения блока матрицы B */

__shared__ float Bs[Block_Size][Block_Size];

 

// Копируем блоки матриц в разделяемую память:

/* каждый поток загружает 1 элемент матрицы A и 1 элемент матрицы B */

As[i][k] = A[a + wA * i + k];

BS[i][k] = B[b + wB * i + k];

/* Принцип определения номера загружаемого элемента в блоке показан на рис. 9.6 */

Рис. 9.6. Принцип расчёта линейного номера элемента с индексами (i, j) в текущем блоке: a – номер 1-го элемента в данном блоке, (i, j) – индексы вычислитель­ного потока, копирующего данный элемент; эти индексы отсчитываются не от начала матрицы, а от начала блока.

/* Синхронизируем потоки, чтобы гарантировать то, что блоки матриц загружены полностью */

 

__syncthreads();

 

/* На этом этапе все элементы обрабатываемых блоков загружены в разделяемую память. Создавать цикл для их загрузки не требуется, так как этот код будет параллельно исполняться стольким числом потоков, равным количеству элементов в блоке матрицы Block_Size× Block_Size. */

 

/* Перемножение блоков матриц A и B, загружен­ных в «быструю» память GPU. Каждый поток вычисляет 1 элемент результирующей матрицы Csub

.

Использование оператора += обеспечивает добавление рассчитываемого слагаемого ко всей сумме

. */

 

for (int j = 0; j < Block_Size; ++j)

Csub += As[i][j] * Bs[j][k];

 

/* Отметим, что при суммировании выше конкрет­ный поток работает не только с теми элементами матриц, которые были загружены в разделяемую память именно им, но и с другими элементами. Это показывает преимущество общего доступа потоков к разделяемой памяти. */

 

/* Синхронизируем потоки. Это гарантирует, что все вычисления будут закончены до того, как начнётся загрузка других блоков на следующей итерации внешнего цикла */

__syncthreads();

} // Возвращение к началу внешнего цикла

 

/* Обратное копирование вычисленного блока результирующей матрицы C в глобальную память GPU, где данные будут доступны центральному процессору. Каждый поток сохраняет 1 элемент */

int c = wB * Block_Size * block_y + Block_Size * block_x;

C[c + wB * i + k] = Csub;

}

 

#endif // #ifndef _MATRIXMUL_KERNEL_H_