Интерпретация подробного вывода 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 КБ для аргументов ядра, вы будете в порядке.