Что такое bank conflicts в shared memory и как их избежать?

Краткий тезис

Bank conflicts — это ситуация в CUDA shared memory, когда несколько потоков одного warp одновременно обращаются к разным адресам, принадлежащим одному и тому же bank (32 banks по 4 байта). Это приводит к сериализации доступа и падению пропускной способности памяти. В контексте LLM bank conflicts часто возникают при транспонировании матриц в attention (например, при вычислении softmax или перестановке QK^T). Основные методы устранения: padding (добавление отступов), изменение stride (шага доступа) и использование warp shuffle инструкций (__shfl_down_sync).

1. Термин: Shared memory и banks

Shared memory — это быстрая (но ограниченная) память на чипе GPU, разделяемая между потоками одного блока. Она организована в 32 banks (для архитектур с compute capability 2.0+). Каждый bank может обслуживать один адрес за такт. Размер bank — 4 байта (32 бита). Если два потока обращаются к разным 4-байтовым словам, которые лежат в одном bank, возникает bank conflict.

Warp — группа из 32 потоков, которые выполняются одновременно (SIMT). Все потоки warp'а исполняют одну инструкцию, но могут обращаться к разным данным. Если в warp'е возникает bank conflict, доступы к shared memory сериализуются: банк обрабатывает запросы последовательно, что снижает эффективную пропускную способность.

2. Что такое bank conflict (определение и пример)

Bank conflict — это ситуация, когда два или более потока из одного warp обращаются к разным адресам, которые отображаются на один и тот же bank. Отображение адреса на bank определяется младшими битами: bank_index = (byte_address / 4) % 32.

Пример:

• Поток 0 читает адрес 0 (bank 0).
• Поток 1 читает адрес 128 (32 слова × 4 байта = 128 байт; 128/4 = 32, 32 % 32 = 0) — тоже bank 0.
• Возникает 2-way bank conflict (конфликт между двумя потоками).

Если все 32 потока обращаются к одному bank — это 32-way conflict, максимальная сериализация (в 32 раза медленнее).

Исключение: если все потоки читают один и тот же адрес (broadcast), конфликта нет — значение передаётся всем потокам за один такт.

3. Почему bank conflicts — проблема

Каждый bank может выполнить только один доступ за такт. При конфликте доступы к одному bank'у выполняются последовательно. Это снижает пропускную способность shared memory пропорционально числу конфликтующих потоков.

Формула эффективной пропускной способности:

Effective bandwidth = (peak bandwidth) / (max conflict degree)

где conflict degree — максимальное количество потоков, обратившихся к одному bank'у.

Для LLM это критично, так как attention — узкое место инференса. Bank conflicts в shared memory могут замедлить вычисление softmax и матричного умножения (при использовании tiling).

4. Как bank conflicts возникают в attention (LLM)

В FlashAttention и других оптимизированных реализациях attention используется shared memory для хранения блоков Q, K, V. Типичные операции, вызывающие bank conflicts:

1. Транспонирование K (при вычислении QK^T). Если матрица K хранится в row-major, а читается по столбцам, соседние потоки могут обращаться к разным строкам, но с одинаковым смещением внутри строки → конфликт.
2. Softmax — редукция по строкам (max, sum). При использовании shared memory для хранения промежуточных значений, если каждый поток обрабатывает свой элемент строки, а строка длиннее 32 элементов, могут возникнуть конфликты.
3. Перестановка данных после умножения (например, при записи результата обратно в глобальную память).

Пример кода с конфликтом (псевдо-CUDA):

__shared__ float smem[32][32];
int tx = threadIdx.x;
int ty = threadIdx.y;
// Чтение по столбцам: каждый поток (tx, ty) читает smem[ty][tx]
float val = smem[ty][tx]; // bank index = (ty*32 + tx) % 32 = tx % 32 — нет конфликта, если tx разные.
// Но если матрица 33x33 и мы используем smem[33][33], то bank index = (ty*33 + tx) % 32 — уже не trivially conflict-free.

5. Способы избежать bank conflicts

5.1 Padding (добавление отступов)

Добавление фиктивных столбцов к массиву в shared memory, чтобы сдвинуть отображение на banks. Например, для матрицы 32×32 объявляем shared float smem[32][32+1]. Тогда доступ smem[ty][tx] даёт bank index = (ty*33 + tx) % 32, что для разных ty даёт разные остатки, и конфликты исчезают.

Плюсы: простота, эффективно для небольших матриц. Минусы: тратится дополнительная память (1 столбец на каждые 32).

5.2 Изменение stride (шага доступа)

Если доступ к shared memory организован с шагом, кратным 32, то все потоки попадают в один bank. Чтобы избежать, нужно выбирать stride, не кратный 32. Например, при транспонировании можно использовать smem[ty][tx] с padding, как выше.

5.3 Warp shuffle instructions

Вместо shared memory можно использовать регистры и инструкции warp shuffle (__shfl_down_sync, __shfl_xor_sync). Они позволяют обмениваться данными между потоками warp'а без обращения к shared memory, полностью устраняя bank conflicts.

Пример редукции через shuffle:

float val = input[threadIdx.x];
for (int offset = 16; offset > 0; offset >>= 1) {
    val += __shfl_down_sync(0xffffffff, val, offset);
}
if (threadIdx.x == 0) output = val;

Плюсы: максимальная скорость, нет shared memory. Минусы: работает только в пределах одного warp (32 потока), сложнее для произвольных паттернов доступа.

5.4 Использование vectorized loads (float4, int4)

Загрузка 16 байт за раз (4 float) уменьшает количество обращений к shared memory, но не устраняет конфликты, если адреса внутри вектора попадают в разные banks. Однако это может снизить число конфликтующих транзакций.

6. Сравнение методов

7. Инструменты профилирования

Для обнаружения bank conflicts используйте NVIDIA Nsight Compute (ncu). В разделе Memory Workload Analysis показывается L1/Shared Memory Bank Conflicts. Метрика l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum — количество конфликтов при чтении.

Пример команды:

ncu --set full -o profile ./my_kernel

8. Влияние на производительность LLM (attention)

В современных реализациях attention (FlashAttention, xformers) bank conflicts минимизируются за счёт:

• Tiling с padding (размер блока 32×32 + padding).
• Использования warp shuffle для softmax (редукция).
• Оптимизированных layout'ов (например, row-major с padding).

Устранение bank conflicts может дать прирост производительности до 20–30% на ядрах attention, что напрямую снижает latency RAG-системы.

9. Пет-проект для закрепления

Задача: Реализовать простое ядро транспонирования матрицы 32×32 в shared memory с bank conflicts и без, измерить время выполнения.

Инструменты: CUDA C++, NVIDIA Nsight Compute, Python (для вызова через PyCUDA или CuPy).

Шаги:

1. Написать ядро transpose_conflict, которое копирует матрицу из глобальной памяти в shared memory (без padding), затем транспонирует и записывает обратно.
2. Написать ядро transpose_noconflict с padding (smem[32][33]).
3. Запустить оба ядра на матрице 1024×1024 (разбить на блоки 32×32).
4. Сравнить время выполнения (CUDA events) и количество bank conflicts (Nsight Compute).
5. Дополнительно: реализовать версию с warp shuffle для редукции (например, для softmax).

Ожидаемый результат: Ядро без конфликтов работает в 1.5–2 раза быстрее, Nsight показывает 0 bank conflicts.

10. Связь с другими вопросами

11. Навигация

• Предыдущий: 703
• Следующий: 705
• Индекс: 00. Индекс разборов

Навигация

• Предыдущий: 703
• Следующий: 705
• Индекс: 00. Индекс разборов