Регистры и разделяемая память в зависимости от возможностей компиляции?
Привет, когда я собираю с nvcc -arch=sm_13 Я получил:
ptxas info : Used 29 registers, 28+16 bytes smem, 7200 bytes cmem[0], 8 bytes cmem[1]
когда я использую nvcc -arch=sm_20 Я получил:
ptxas info : Used 34 registers, 60 bytes cmem[0], 7200 bytes cmem[2], 4 bytes cmem[16]
Я думал, что все параметры ядра передаются в общую память, но для sm_20 это не так...?! Возможно, они также передаются в регистры? Глава моей функции выглядит следующим образом:
__global__ void func(double *, double , double, int)
Пока спасибо!
2 ответа
В устройствах с вычислительной возможностью 2.x аргументы ядра хранятся в постоянной памяти. Разница в регистре, вероятно, связана с различиями в коде, сгенерированном для функций математической библиотеки между версиями. Есть ли такие вещи, как трансцендентные функции или sqrt в ядре?
Как утверждает @talonmies, различия в разделяемой памяти связаны с тем, что устройства SM 2.x передают аргументы ядра через постоянную, а не разделяемую память.
Однако одно из основных отличий в использовании регистров в устройствах SM 2.x заключается в том, что, хотя устройства SM 1.x имеют выделенные регистры адресов для команд загрузки и хранения, SM 2.x использует регистры общего назначения для адресов. Это ведет к увеличению давления в регистре SM 2.x. К счастью, файл регистра также в 2 раза больше на GF100 (SM 2.0) по сравнению с GT200 (SM 1.3).