У меня была программа 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(); и стоимость доступа к разделяемой памяти?
