Комментировать

Хм.. в самом загруженном ядре

Хм.. в самом загруженном ядре увидел:
.reg .f32 %f<2>;
.reg .f64 %fd<17>;
.reg .pred %p<5>;
.reg .s32 %r<46>;
.reg .s64 %rl<2>;

что-то как-то слишком много - компилятор совсем не стесняется.
Я правильно понимаю, что это требует больше 80 регистров? И это значит, что на Compute Capability 2.0, не получиться запустить 512 потоков в блоке? У меня сейчас это ядро запускается с 512 потоками в блоке, и результат получается адекватный.
Видимо придётся переписать код, в ассемблерном стиле - сразу объявить переменные для всех промежуточных результатов, и каждую операцию писать в новой строке.. бред какой-то..
И что значит:
"Since PTX supports virtual registers, it is quite common for a compiler frontend to generate
a large number of register names."
Что за виртуальные регистры? У одного регистра разные имена(тогда всё ясно, но как тогда для OpenCL посчитать количество реально используемых регистров)?