У меня была программа CUDA, в которой регистры ядра ограничивались Максимальное теоретическое достигнутое занятие до% 50. Поэтому я решил использовать разделяемую память вместо регистров для тех переменных, которые были постоянными между потоками блоков и были почти доступны только для чтения во время запуска ядра. Я не могу предоставить исходный код здесь; что я сделал концептуально так:Увеличение достигнутого уровня занятости не увеличивает скорость вычислений линейно
Моя первоначальная программа:
__global__ void GPU_Kernel (...) {
__shared__ int sharedData[N]; //N:maximum amount that doesn't limit maximum occupancy
int r_1 = A; //except for this first initialization, these registers don't change anymore
int r_2 = B;
...
int r_m = Y;
... //rest of kernel;
}
Я изменил выше программы:
__global__ void GPU_Kernel (...) {
__shared__ int sharedData[N-m];
__shared__ int r_1, r_2, ..., r_m;
if (threadIdx.x == 0) {
r_1 = A;
r_2 = B;
...
r_m = Y; //last of them
}
__syncthreads();
... //rest of kernel
}
Теперь нити перекосов внутри блока выполняют трансляцию читает доступ только что созданный переменные общей памяти. В то же время потоки не используют слишком много регистров для ограничения достигнутого уровня занятости.
Вторая программа имеет максимальная теоретическая достигнутая занятость равна% 100. В реальных проходах средняя достигнутая занятость для первых программ составляла ~% 48, а для второго - около 80%. Но проблема в увеличении чистой скорости составляет около 5% до 10, что намного меньше, чем я ожидал, учитывая улучшенную занятость. Почему эта корреляция не линейна?
Учитывая нижеследующее изображение из технической документации Nvidia, я думал, что, когда достигнутое занятие составляет% 50, например, половина серверов SMX (в новых архитектурах) неактивна за раз, поскольку избыточные запрашиваемые ресурсы другими ядрами чтобы они не были активными. Разве мое понимание ошибочно? Или это неполное объяснение выше явления? Или он добавлен __syncthreads(); и стоимость доступа к общей памяти?
