Что хотят услышать интервьюеры:

Occupancy — это доля активных warps на SM относительно максимально возможного числа, которое железо может держать одновременно. Это важный, но не единственный фактор производительности: высокая occupancy помогает скрывать latency, но сама по себе не гарантирует быстрый kernel. Оптимизация обычно сводится к балансу между числом регистров, размером shared memory и количеством потоков в блоке.

Определение:

Occupancy в CUDA — это мера того, насколько полно загружен streaming multiprocessor (SM) активными потоками выполнения. Обычно её выражают как отношение числа активных warps к максимально поддерживаемому числу warps на SM.

На occupancy влияют ограничения железа и ресурсы, которые использует kernel:

число регистров на поток;
объём shared memory на блок;
максимальное число блоков и warps на SM;
размер блока и конфигурация запуска.

Высокая occupancy помогает GPU переключаться между warps и скрывать задержки памяти или вычислений, но иногда kernel с меньшей occupancy работает быстрее, если у него меньше pressure на регистры и больше ILP.

Пример использования:

Типичная ситуация — kernel работает медленно из-за слишком большого числа регистров на поток. Тогда occupancy падает, и SM не может держать достаточно warps, чтобы скрывать latency.

__global__ void saxpy(const float* x, const float* y, float* out, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = a * x[i] + y[i];
    }
}

int main() {
    int n = 1 << 20;
    int threads = 256;
    int blocks = (n + threads - 1) / threads;

    saxpy<<<blocks, threads>>>(x, y, out, 2.0f, n);
}

Если увеличить threads до разумного значения, например 128–256, и при этом kernel не упирается в регистры или shared memory, occupancy часто становится выше. Но если компилятор начинает активно использовать регистры для более сложной логики, то чрезмерное увеличение threads может не дать выигрыша.

Пояснение кода:

Код показывает простой memory-bound kernel saxpy, где каждый поток обрабатывает один элемент массива.

threadIdx.x и blockIdx.x вычисляют глобальный индекс элемента.
Условие if (i < n) защищает от выхода за границы массива.
Каждый поток выполняет простую операцию a * x[i] + y[i].
При запуске <<<blocks, threads>>> важно подобрать размер блока так, чтобы:
- на SM помещалось достаточно warps;
- не было лишнего расхода регистров;
- shared memory не ограничивала число активных блоков.

Если бы для этого kernel использовалось много локальных переменных, сложные ветвления или большой shared memory, occupancy могла бы резко упасть. Тогда оптимизация шла бы не через “максимум потоков”, а через уменьшение использования ресурсов на поток, пересмотр структуры данных или упрощение вычислений.

Ключевые моменты:

Occupancy — это не “скорость”, а степень загрузки SM активными warps.
Высокая occupancy помогает скрывать latency, особенно в memory-bound задачах.
Основные ограничители: регистры, shared memory, число блоков и warps на SM.
Увеличение occupancy не всегда ускоряет kernel; иногда важнее ILP, коалесинг памяти и уменьшение лишней работы.
Оптимизация обычно начинается с профилирования и проверки, что именно ограничивает запуск: registers, shared memory или scheduler.
Подбирать размер блока и ресурсы нужно под конкретный kernel, а не по универсальному правилу.

Что такое occupancy и как её оптимизировать?