Интерпретация выходных данных --ptxas-options=-v

Question

Я пытаюсь понять использование ресурсов для каждого из моих потоков CUDA для рукописного ядра.

Я собрал свой kernel.cu подать в kernel.o файл с nvcc -arch=sm_20 -ptxas-options=-v

и я получил следующий вывод

ptxas info    : Compiling entry function '_Z12searchkernel6octreePidiPdS1_S1_' for 'sm_20'
ptxas info    : Function properties for _Z12searchkernel6octreePidiPdS1_S1_
    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 (что я предполагаю, это постоянная память) повторяется дважды с разными цифрами? В ядре я не использую постоянную память. Означает ли это, что компилятор под капотом скажет графическому процессору использовать некоторую постоянную память?

Answer 1

• Каждый поток CUDA использует 46 регистров? Да исправить
• Нет регистров разлива в локальную память? Да исправить
• 72 байта - общая сумма памяти для кадров стека __global__ а также __device__ функции? Да исправить
• В чем разница между хранилищами разлива 0 байт и разливами 0 байт?
  • Справедливый вопрос, нагрузки могут быть больше, чем хранилища, так как вы можете пролить вычисленное значение, загрузить его один раз, выбросить его (т.е. сохранить что-то еще в этот регистр), а затем загрузить его снова (т.е. повторно использовать). Обновление: обратите внимание также, что количество разливов / хранилищ разлива основано на статическом анализе, как описано @njuffa в комментариях ниже.
• Почему информация для cmem (которая, как я предполагаю, является постоянной памятью) повторяется дважды с разными цифрами? В ядре я не использую постоянную память. Означает ли это, что компилятор под капотом скажет графическому процессору использовать некоторую постоянную память?
  • Постоянная память используется для нескольких целей, включая __constant__ переменные и аргументы ядра, используются разные "банки", которые начинают немного детализироваться, но при условии, что вы используете менее 64 КБ для __constant__ переменные и менее 4 КБ для аргументов ядра вы будете в порядке.