CUDA: —ptxas-options = -v разделяемая память и cudaFuncAttributes.sharedSizeBytes не совпадают
Я пытаюсь использовать reqRegs а также sharedSizeBytes в cudaFuncAttributes структура для динамической оптимизации размера блока ядра во время выполнения.
Моя текущая реализация снимает текст stdout с nvcc --ptxas-options=-v обнаружить регистр и использование общей памяти ядра. Этот метод немного хакерский и зависит от точного формата выходного текста из --ptxas-options=-v, который может измениться без предупреждения.
Моя проблема в том, что я вижу несоответствие между значением общей памяти ‘smm’, сообщенным в --ptxas-options=-v вывод и sharedSizeBytes в cudaFuncAttributes struct, которая беспокоит меня, что либо оценка общей памяти, которую я использовал до сих пор, неверна, либо что sharedSizeBytes переменная ненадежна, то есть я не могу использовать ее для оптимизации размера блока во время выполнения.
Вот вывод nvcc --ptxas-options=-v для одного такого ядра …
ptxas info : Used 14 registers, 2088 bytes smem, 48 bytes cmem[1]
… по сравнению со значением cudaFuncAttributes.sharedSizeBytes = 296 во время выполнения для точно такого же ядра. Кто-нибудь знает, что здесь может происходить?
Вот еще один пример с другим ядром:
ptxas info : Used 18 registers, 2132 bytes smem, 48 bytes cmem[1]
где cudaFuncAttributes.sharedSizeBytes = 340 во время выполнения.
Благодарю.
Решение
Спасибо Роберт и Марко за ваши ответы. Они помогли мне исключить несколько случаев.
Оказывается, что несоответствие в использовании общей памяти, о которой сообщают, произошло из-за объема общей памяти, использованной после первой тестовой компиляции (сообщенный --ptxas-options=-v) отличается от объема разделяемой памяти, используемой конечной программой с измененным размером блока (сообщается cudaFuncAttributes.sharedSizeBytes). (РЕДАКТИРОВАТЬ для ясности)
Разница в разделяемой памяти была вызвана тем, что распределение массива разделяемой памяти зависело от размера блока; например:
__shared__ float myArray[BLOCK_SIZE];
Вышеупомянутое утверждение использует другой объем разделяемой памяти в программе с размером блока 256, чем тот же исходный код, скомпилированный с оптимизированным размером блока 192. Сейчас это кажется очевидным, но в оптимизированном коде CUDA следует обратить внимание поколение.
Другие решения