Интерпретация подробного вывода ptxas, часть I
Я пытаюсь понять использование ресурсов для каждого из моих потоков CUDA для рукописного ядра.
Я скомпилировал свой файл kernel.cu
файл kernel.o
с помощью nvcc -arch=sm_20 -ptxas-options=-v
и я получил следующий вывод (пропущенный через c++filt
):
ptxas info : Compiling entry function 'searchkernel(octree, int*, double, int, double*, double*, double*)' for 'sm_20'
ptxas info : Function properties for searchkernel(octree, int*, double, int, double*, double*, double*)
72 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 46 registers, 176 bytes cmem[0], 16 bytes cmem[14]
Глядя на вывод выше, правильно ли сказать, что
- каждый поток CUDA использует 46 регистров?
- нет разливов регистров в локальную память?
У меня также есть некоторые проблемы с пониманием результатов.
-
Мое ядро вызывает множество функций __device__
. IS 72 байта - общая сумма памяти для стековых кадров функций __global__
и __device__
?
-
В чем разница между 0 byte spill stores
0 bytes spill loads
-
Почему информация для cmem
(которая, как я предполагаю, является постоянной памятью) повторяется дважды с разными цифрами? Внутри ядра я не использую постоянную память. Означает ли это, что компилятор скрытно скажет графическому процессору использовать некоторую постоянную память?
Этот вопрос "продолжается" в: Интерпретация подробного вывода ptxas, часть II
Ответы
Ответ 1
- Каждый поток CUDA использует 46 регистров? Да, правильно
- В локальную память отсутствует пролитие реестра? Да, правильно
- Является ли 72 байта суммарной суммой памяти для кадров стека функций
__global__
и __device__
? Да, правильно
- В чем разница между 0-байтовыми хранилищами разлива и 0 байтами разливов?
- Справедливый вопрос: нагрузки могут быть больше, чем магазины, так как вы можете разлить вычисленное значение, загрузить его один раз, отбросить его (т.е. сохранить что-то еще в этом регистре), а затем снова загрузить его (то есть повторно использовать). Обновление: обратите внимание также, что счет загрузки/хранения разливов основан на статическом анализе, как описано @njuffa в комментариях ниже.
- Почему информация для cmem (которую я предполагаю - постоянная память) повторяется дважды с разными цифрами? Внутри ядра я не использую постоянную память. Означает ли это, что компилятор находится под капотом и собирается сказать, что GPU использует некоторую постоянную память?
- Постоянная память используется для нескольких целей, включая переменные
__constant__
и аргументы ядра, используются разные "банки", которые начинают немного детализироваться, но пока вы используете меньше, чем 64 КБ для ваших переменных __constant__
и менее 4 КБ для аргументов ядра, вы будете в порядке.