Перемножение матриц: наивное и тайловое

1. Двумерная задача — двумерный грид

Произведение C = A·B квадратных матриц N×N: элемент C_row,col — скалярное произведение строки row матрицы A и столбца col матрицы B (по N умножений и сложений на элемент, всего 2N³ операций). Элементов N² — и мы уже знаем, что делать с большим числом одинаковых независимых вычислений: по треду на элемент. Только теперь у элемента два индекса, и грид удобно сделать двумерным (dim3 из C02):

#define N 1024          // размер матрицы
#define TILE 16          // размер блока 16×16 = 256 тредов

dim3 block(TILE, TILE);
dim3 grid(N / TILE, N / TILE);        // 64×64 блока; N кратно TILE для простоты
matMul<<<grid, block>>>(A, B, C, N);

Индексация — та же формула C04, применённая дважды, по оси на измерение:

int row = blockIdx.y * blockDim.y + threadIdx.y;   // строка элемента C
int col = blockIdx.x * blockDim.x + threadIdx.x;   // столбец элемента C

Матрица хранится в памяти одномерно, построчно: элемент (row, col) лежит по смещению row * N + col. Наивное ядро:

__global__ void matMulNaive(const float* A, const float* B, float* C, int n)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    float sum = 0.0f;
    for (int k = 0; k < n; k++)
        sum += A[row * n + k] * B[k * n + col];   // строка A · столбец B

    C[row * n + col] = sum;
}

Ядро верное и уже быстрее CPU. Но посчитаем байты, как учил C06: каждый тред читает N элементов строки A и N элементов столбца B — 2N чтений глобальной памяти ради 2N операций сложения-умножения. Одна арифметическая операция на каждое чтение DRAM — для задачи, где данных N² и работы N³, это расточительность: соседние треды читают одни и те же строки и столбцы, каждый сам по себе.

2. Тайлы: один раз прочитал — шестнадцать раз использовал

Присмотримся к блоку 16×16 тредов, вычисляющему квадрат 16×16 матрицы C. Всем его тредам нужны одни и те же 16 строк A и 16 столбцов B. Идея тайлинга: пусть блок читает эти данные из глобальной памяти совместно и по одному разу — в разделяемую память (C05), а уже из неё каждый тред берёт свои сомножители сколько угодно раз.

Целиком 16 строк в разделяемую память не влезут (16 · 1024 float = 64 КБ — больше лимита блока), поэтому работа нарезается на фазы. В фазе p блок загружает два тайла 16×16: кусок строк A с колонками [16p … 16p+15] и кусок столбцов B со строками [16p … 16p+15], — перемножает их и накапливает вклад в сумму. Фаз всего N / 16, и после последней в аккумуляторе каждого треда — готовый элемент C:

__global__ void matMulTiled(const float* A, const float* B, float* C, int n)
{
    __shared__ float tA[TILE][TILE];   // тайл полосы A
    __shared__ float tB[TILE][TILE];   // тайл полосы B

    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE + ty;
    int col = blockIdx.x * TILE + tx;

    float sum = 0.0f;

    for (int p = 0; p < n / TILE; p++) {
        // каждый тред приносит по одному элементу каждого тайла:
        // 256 тредов загружают 2·256 значений — совместно и без дублей
        tA[ty][tx] = A[row * n + (p * TILE + tx)];
        tB[ty][tx] = B[(p * TILE + ty) * n + col];

        __syncthreads();               // тайлы загружены целиком — можно считать

        for (int k = 0; k < TILE; k++)
            sum += tA[ty][k] * tB[k][tx];   // 16 умножений по быстрой памяти

        __syncthreads();               // досчитали — только теперь можно грузить следующий
    }

    C[row * n + col] = sum;
}

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

Теперь арифметика выигрыша. Каждый элемент A теперь читается из глобальной памяти не 16 тредами по разу каждый, а один раз на блок — дальше он крутится в разделяемой памяти. Чтений глобальной памяти на тред: было 2N, стало 2N / TILE — в 16 раз меньше. Общее правило: тайл T×T сокращает трафик DRAM в T раз; на практике ядро ускоряется в 3–6 раз (не в 16 — потому что наивной версии частично помогали кэши, а тайловой добавились барьеры). Проверьте на своей карте секундомером из C06 — числа скажут сами.

3. Тренажёр: карта тайлов

Контрольные вопросы

• Сколько чтений глобальной памяти делает наивное ядро на один элемент C и почему это плохо?
  2N: строка A и столбец B целиком. На каждое чтение DRAM приходится одна операция умножения-сложения, то есть ядро упирается в память, а не в счёт. Хуже того, соседние треды читают одни и те же строки и столбцы независимо друг от друга — данные ездят из DRAM многократно.
• Опишите одну фазу тайлового ядра: кто что загружает и что потом считает.
  Блок 16×16 совместно загружает два тайла: каждый тред приносит по одному элементу тайла A и тайла B — 512 значений без дублей. Барьер. Затем каждый тред выполняет 16 умножений-сложений по разделяемой памяти, накапливая свой элемент C. Барьер — и следующая фаза загружает следующие тайлы полосы. Фаз N/TILE.
• Зачем в фазе два __syncthreads и что сломается без каждого из них?
  Первый — после загрузки: без него треды начнут перемножать тайл, который соседние варпы ещё не доложили, и прочитают мусор. Второй — после вычислений: без него быстрые треды начнут загружать следующий тайл поверх текущего, пока медленные ещё считают по нему, — данные перезапишутся под ногами.
• Во сколько раз тайлинг T×T сокращает трафик глобальной памяти и почему ускорение меньше?
  Ровно в T раз: элемент читается один раз на блок вместо T раз (по разу каждым тредом строки/столбца блока). Реальное ускорение — в 3–6 раз при T = 16: наивной версии частично помогали кэши L1/L2, а тайловой добавились накладные расходы барьеров.
• Почему нельзя бездумно увеличить TILE до 32?
  Блок станет 32×32 = 1024 треда — впритык к лимиту, а два тайла float 32×32 займут 8 КБ разделяемой памяти на блок: на SM поместится меньше блоков, планировщику станет нечем прятать латентность. Размер тайла — компромисс, который проверяют замером, а не интуицией.

Источники

1. CUDA C++ Programming Guide. Shared Memory: Matrix Multiplication Example // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-programming-guide/ (дата обращения: 09.07.2026).
2. CUDA C++ Best Practices Guide. Shared Memory in Matrix Multiplication // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/ (дата обращения: 09.07.2026).