Отслеживание использования регистра ядра cuda

Question

Я пытаюсь отследить использование регистра и натолкнулся на интересный сценарий. Рассмотрим следующий источник:

#define OL 20
#define NHS 10

__global__ void loop_test( float ** out, const float ** in,int3 gdims,int stride){

        const int idx = blockIdx.x*blockDim.x + threadIdx.x;
        const int idy = blockIdx.y*blockDim.y + threadIdx.y;
        const int idz = blockIdx.z*blockDim.z + threadIdx.z;

        const int index = stride*gdims.y*idz + idy*stride + idx;
        int i = 0,j =0;
        float sum =0.f;
        float tmp;
        float lf;
        float u2, tW;

        u2 = 1.0;
        tW = 2.0;

        float herm[NHS];

        for(j=0; j < OL; ++j){
                for(i = 0; i < NHS; ++i){
                        herm[i] += in[j][index];
                }
        }

        for(j=0; j<OL; ++j){
                for(i=0;i<NHS; ++i){
                        tmp = sum + herm[i]*in[j][index];
                        sum = tmp;
                }
                out[j][index] = sum;
                sum =0.f;
        }

}

В качестве примечания к источнику - промежуточная сумма, которую я мог сделать + =, но играл с тем, как изменение использования регистра эффектов (кажется, что нет - просто добавляет дополнительную инструкцию mov) , Кроме того, этот источник ориентирован на доступ к памяти, отображаемой в 3D-пространстве.

Подсчет регистров может показаться, что существует 22 регистра (я полагаю, что float [N] занимает N + 1 регистров - пожалуйста, поправьте меня, если я ошибаюсь) на основе объявлений.

Однако компилируется с:

nvcc -cubin -arch=sm_20 -Xptxas="-v" src/looptest.cu

Выходы:

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 25 registers, 72 bytes cmem[0]

Хорошо, так что число отличается от ожидаемого. Дополнительно, если скомпилировано с:

nvcc -cubin -arch=sm_13 -Xptxas="-v" src/looptest.cu

Использование регистра далеко меньше - 8, если быть точным (очевидно, из-за более строгого соблюдения в sm_20, чем в sm_13, по стандартам математики IEEE с плавающей запятой?):

ptxas info    : Compiling entry function '_Z9loop_testPPfPPKfS2_4int3i' for 'sm_13'
ptxas info    : Used 17 registers, 40+16 bytes smem, 8 bytes cmem[1]

В качестве последнего примечания измените макрос OL на 40, и вдруг:

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 28 registers, 72 bytes cmem[0]

В заключение я хотел бы знать, где израсходованы регистры и каковы результаты моих наблюдений за парой.

У меня недостаточно опыта сборки, чтобы пройти через cuobjdump - ответ наверняка лежит там - может быть, кто-то может рассказать мне о том, что я должен искать, или показать мне руководство о том, как подойти к сборке дамп.

Answer 1

sm_20 и sm_13 - очень разные архитектуры с очень разным дизайном набора команд (ISA).Основное отличие, которое вызывает увеличение использования регистров, которое вы видите, заключается в том, что sm_1x имеет регистры адресов специального назначения, а sm_2x и более поздние - нет.Вместо этого адреса хранятся в регистрах общего назначения, как и значения, что означает, что большинству программ требуется больше регистров для sm_2x, чем для sm_1x.

sm_20 также имеет в два раза больший размер файла регистров sm_13, чтобы компенсировать это.

Answer 2

Использование регистра не обязательно имеет тесную связь с числом переменных.

Компилятор пытается оценить выигрыш в скорости хранения переменной в регистре между двумя точками использования в коде, сравнивая потенциальный выигрыш в одном ядре со стоимостью для всех одновременно работающих ядер из-за меньшего количества доступных регистров в пуле регистров. (A Fermi SM имеет 32768 регистров). Поэтому неудивительно, если изменение вашего кода вызывает неожиданные колебания в количестве используемых регистров.

Вам действительно стоит беспокоиться только об использовании регистра, если профилировщик говорит, что ваша занятость ограничена использованием регистра. В этом случае вы можете использовать параметр --maxrregcount, чтобы уменьшить количество регистров, используемых одним ядром, чтобы увидеть, улучшает ли это общую скорость выполнения.

Чтобы уменьшить количество регистров, используемых ядром, вы можете попытаться сделать использование переменных максимально локальным. Например, если вы делаете:

set variable 1
set variable 2
use variable 1
use variable 2

Это может привести к использованию 2 регистров. Пока, если вы:

set variable 1
use variable 1
set variable 2
use variable 2

Это может привести к использованию 1 регистра.