Регистры и разделяемая память в зависимости от вычислительных возможностей компиляции?

Привет, когда я компилирую с 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)

Спасибо!


cuda nvcc compiler-construction
person tim    schedule 17.05.2011    source источник

Ответы (2)


arrow_upward
3
arrow_downward

В устройствах с вычислительными возможностями 2.x аргументы ядер хранятся в постоянной памяти. Разница в регистрах, вероятно, связана с различиями в коде, сгенерированном для функций математической библиотеки между версиями. Есть ли в ядре такие вещи, как трансцендентные функции или sqrt?

person talonmies    schedule 17.05.2011
comment
Во-первых, спасибо за параметры, у вас есть какой-нибудь источник об этом? Далее: Нет, никаких математических функций! - person tim; 18.05.2011
comment
Для справки, это, вероятно, есть в руководстве по ptx, но все, что я могу сказать, это то, что разработчик NVIDIA сказал мне, что это так. Другое возможное различие в регистрах — это указатель. Указатели используют 1 регистр на вычислительных устройствах 1.x и 2 на Fermi, который имеет внутреннюю разрядность 64 бита. - person talonmies; 18.05.2011

arrow_upward
4
arrow_downward

Как утверждает @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).

person harrism    schedule 18.05.2011
comment
Марк, несмотря на то, что регистровый файл действительно в два раза больше на Fermi, максимальное количество регистров на поток также вдвое меньше, чем в устройствах sm 1.x. Ограничение в 64 регистра часто используется в кодах, с которыми я работаю. Я знаю, что Василий Волков также указал на влияние давления регистров на способность кода достигать высокого уровня параллелизма на уровне инструкций на картах Fermi. - person talonmies; 18.05.2011
comment
Да, Fermi увеличивает максимальное количество потоков на SM на 50% (1536/1024), удаляет регистры адресов, уменьшает вдвое максимальное количество регистров на поток, удваивает размер регистрового файла на SM и немного увеличивает задержку доступа к общей памяти. Все это приводит к повышенной чувствительности к использованию ресурсов. Дублирование регистрового файла помогает, но вы правы, все равно туго. Но это реальность массивно-параллельных процессоров. Всегда есть компромиссы, подобные этому, которые необходимо учитывать при разработке нового графического процессора. Однако в целом архитектура Fermi гораздо более эффективна, чем архитектура Tesla. - person harrism; 19.05.2011