Графические процессоры NVIDIA с возможностью вычислений 1.0 - 7.x будут выдавать инструкции для деформации по порядку. Регистры специального назначения clock и clock64 могут использоваться для измерения временных интервалов кода путем считывания регистра до и после последовательности инструкций.
Это может быть полезно для оценки количества циклов, которые потребовались для выдачи последовательности инструкций для одной основы.
СЛУЧАЙ 1: Задержка выдачи инструкций
Показания часов64 вставляются до и после последовательности инструкций. В приведенном ниже случае clock64 считывает перенос одной глобальной загрузки. Этот стиль оценивает задержку выдачи команды глобальной инструкции загрузки. Деформация может быть остановлена между началом и концом CS2R, увеличивая продолжительность. Причины могут быть следующие: - not_selected - планировщик деформации выбрал более высокий приоритет warp - no_instruction - LDG находился в новой строке кэша команд и деформация останавливается до тех пор, пока строка кэша не будет выбрана - mio_throttle - инструкция LDG не может быть выдана, поскольку Очередь инструкций для загрузочного хранилища заполнена. - lg_throttle - Инструкция LDG не может быть выдана, поскольку очередь инструкций для модуля хранилища нагрузки достигла локального / глобального водяного знака.
Для повышения точности рекомендуется измерить последовательность инструкций, а не одну инструкция.
PTX
asm volatile("mov.u64 %0, %%clock64;" : "=l"(start) :: "memory");
asm volatile("ld.global.ca.u32 data, [%0];"::"l"(po):"memory");
asm volatile("mov.u64 %0, %%clock64;" : "=l"(stop) :: "memory");
SASS (SM_70)
/*0420*/ CS2R R2, SR_CLOCKLO ;
/*0430*/ LDG.E.64.STRONG.CTA R4, [R4] ;
/*0440*/ CS2R R6, SR_CLOCKLO ;
СЛУЧАЙ 2: Задержка выполнения инструкции
A Clock64 read вставляется перед последовательностью инструкций. Набор инструкций, которые гарантируют завершение последовательности инструкций и считывание часов64, вставляется после последовательности инструкций. В приведенном ниже случае целое добавление вставляется перед последним чтением, которое зависит от значения из глобальной загрузки. Этот метод может использоваться для оценки продолжительности выполнения глобальной загрузки.
PTX
asm volatile("mov.u64 %0, %%clock64;" : "=l"(start) :: "memory");
asm volatile("ld.global.ca.u32 data, [%0];"::"l"(po):"memory");
asm volatile("add.u32 %0, data, %0;":"+l"(sink)::"memory");
asm volatile("mov.u64 %0, %%clock64;" : "=l"(stop) :: "memory");
SASS (SM_70)
/*0420*/ CS2R R2, SR_CLOCKLO ;
/*0430*/ LDG.E.64.STRONG.CTA R4, [R4] ;
/*0440*/ IADD R4, R4, 1 ;
/*0450*/ CS2R R6, SR_CLOCKLO ;
ДИАГРАММА
Период измерения для случая 1 и случая 2 показан на диаграмме в форме волны. На диаграмме показаны инструкции CS2R и IADD, выполняемые за 4 цикла. В инструкциях CS2R указано время третьего цикла.
Для случая 1 измеренное время может составлять всего 2 цикла. Для случая 1 измеренное время включает нагрузку из глобальной памяти. Если нагрузка попадает в кэш L1, то это время составляет 20-50 циклов, иначе время, вероятно, превышает 200 циклов.
ПРЕДУПРЕЖДЕНИЕ
На практике это Тип выдачи команды или задержки выполнения инструкции очень сложно реализовать. Эти методы могут быть использованы для написания микро-тестов или больших последовательностей кода. В случае микропроцессоров очень важно понимать и потенциально изолировать другие факторы, такие как планирование деформации, пропуски кэша команд, постоянные пропуски кэша и т. Д. c.
Компилятор не обрабатывает чтение тактов / clock64 как забор инструкции. Компилятор может переместить чтение в неожиданное место. Рекомендуется всегда проверять сгенерированный код SASS.
Compute Capability 6.0 и выше поддерживает вытеснение на уровне команд. Прерывание уровня инструкций приведет к неожиданным результатам.