Как загрузить несколько чисел Float4 из памяти в регистры с помощью встроенной сборки GCN в AMD HIP? ⇐ C++

[Гость]

Мотивация
Я провожу микротесты графических процессоров AMD, чтобы понять их характеристики производительности и повысить производительность ядра. Теперь я подозреваю, что различные результаты распределения регистров и планирования инструкций могут повлиять на эффективную пропускную способность памяти. Я заметил, что компилятор пытается чередовать инструкции памяти и инструкции вычисления, а также пытается сохранить регистры, загружая новые значения сразу после завершения предыдущей арифметической инструкции. Я обнаружил, что в некоторых случаях может быть заметная разница в производительности. Если я намеренно добавляю некоторые неоптимизируемые операции, такие как запись фиктивного значения в LDS, чтобы помешать компиляторам выполнять такое чередование, иногда это повышает производительность. Я подозреваю, что причина в том, что это изменяет количество одновременных запросов к памяти, выполняемых в данный момент, что приводит к уменьшению используемой пропускной способности памяти.

Поэтому я решил использовать встроенную сборку при настройке AMD HIP, чтобы лучше контролировать микротесты.
Попытка 1
Следующая программа HIP пытается использовать встроенную ассемблерную сборку для загрузки 4 чисел с плавающей запятой в float4 tmp11, tmp12, tmp13, tmp14

#include #include __global__ void kernel( float* __restrict массив, плавать4* наружу, uint32_t идентификатор ) { float* a_ptr = &array[idx]; float4 tmp11, tmp12, tmp13, tmp14; #ifdef __HIP_PLATFORM_AMD__ asm нестабильный( "global_load_dwordx4 %0, %1, выкл\n\t" : "=v" (tmp11) : "в" (a_ptr) ); asm нестабильный( "global_load_dwordx4 %0, %1, выключено, смещение: 16\n\t" : "=v" (tmp12) : "в" (a_ptr) ); asm нестабильный( "global_load_dwordx4 %0, %1, выключено, смещение: 32\n\t" : "=v" (tmp13) : "в" (a_ptr) ); asm нестабильный( "global_load_dwordx4 %0, %1, выключено, смещение: 48\n\t" "s_waitcnt vmcnt(0)" : "=v" (tmp14) : "в" (a_ptr) ); #endif } int main (недействительный) { } Компилируя исходный код с помощью hipcc -S main.cpp -o main.S -O3 и проверяя main.S, я обнаружил, что сгенерированная сборка неверна. Все значения загружаются в одни и те же регистры.

;;#ASMSTART global_load_dwordx4 v[0:3], v[4:5], выкл. ;;#АСМЕНД ;;#АСМСТАРТ global_load_dwordx4 v[0:3], v[4:5], выкл., смещение:16 ;;#АСМЕНД ;;#АСМСТАРТ global_load_dwordx4 v[0:3], v[4:5], выкл., смещение:32 ;;#АСМЕНД ;;#АСМСТАРТ global_load_dwordx4 v[0:3], v[4:5], выкл., смещение:48 s_waitcnt vmcnt (0) ;;#АСМЕНД Попытка 2 Компилятор, похоже, определил, что загрузки не имеют никакого эффекта, и он может повторно использовать одни и те же регистры для всех ассемблерных инструкций, что имеет смысл. Если я проделаю с переменными некоторые арифметические операции, значения действительно будут загружены в разные регистры, вероятно, из-за разного распределения регистров.

Поэтому моя следующая попытка — загрузить несколько значений в несколько регистров с использованием нескольких инструкций в одном операторе встроенного ассемблера. Теперь компилятор знает, что результат должен поступать в разные регистры.

#include #include __global__ void kernel( float* __restrict массив, плавать4* наружу, uint32_t идентификатор ) { float* a_ptr = &array[idx]; float4 tmp11, tmp12, tmp13, tmp14; #ifdef __HIP_PLATFORM_AMD__ asm нестабильный( "global_load_dwordx4 %0, %4, выкл\n\t" "global_load_dwordx4 %1, %4, смещение выключения: 16\n\t" "global_load_dwordx4 %2, %4, смещение выключения: 32\n\t" "global_load_dwordx4 %3, %4, смещение выключения: 48\n\t" "s_waitcnt vmcnt(0)" : «=v» (tmp11), «=v» (tmp12), «=v» (tmp13), «=v» (tmp14) : "в" (a_ptr) ); #endif } int main (недействительный) { } К сожалению, сгенерированная сборка по-прежнему неверна.

;;#ASMSTART global_load_dwordx4 v[0:3], v[0:1], выкл. global_load_dwordx4 v[4:7], v[0:1], смещение выключения:16 global_load_dwordx4 v[8:11], v[0:1], смещение выключения:32 global_load_dwordx4 v[12:15], v[0:1], смещение выключения: 48 s_waitcnt vmcnt (0) ;;#АСМЕНД Первая инструкция загрузки блокирует регистры v[0:1], поэтому все последующие загрузки не будут работать должным образом.
Вопрос
Как загрузить несколько чисел с плавающей точкой 4 из памяти в регистры с помощью встроенной сборки GCN в AMD HIP, используя правильный синтаксис встроенной сборки?