Порядок выполнения инструкций от cuda driver

Question

Следующий фрагмент кода

asm volatile("mov.u64 %0, %%clock64;" : "=l"(start) :: "memory");
asm volatile("ld.global.ca.u64 data, [%0];"::"l"(po):"memory");
asm volatile("mov.u64 %0, %%clock64;" : "=l"(stop) :: "memory");

выглядит так в коде SASS

/*0420*/                   CS2R R2, SR_CLOCKLO ;                                           /* 0x0000000000027805 */
/*0430*/                   LDG.E.64.STRONG.CTA R4, [R4] ;                                  /* 0x0000000004047381 */
/*0440*/                   CS2R R6, SR_CLOCKLO ;                                           /* 

Я хочу быть уверен, что планировщик выдает второй CS2R после инструкция LDG и не ранее из-за какой-либо оптимизации, такой как выполнение вне порядка. Как я могу быть уверен в этом?

ОБНОВЛЕНИЕ:

Основываясь на предложении Грега, я добавил зависимую инструкцию, которая выглядит как

  asm volatile("mov.u64 %0, %%clock64;" : "=l"(start) :: "memory");
  asm volatile("ld.global.ca.u64 data, [%0];"::"l"(po):"memory");
  asm volatile("add.u64 %0, data, %0;":"+l"(sink)::"memory");
  asm volatile("mov.u64 %0, %%clock64;" : "=l"(stop) :: "memory");

, где определено uint64_t sink = 0; , Тем не менее, я вижу только один LDG между инструкциями CS2R. Я ожидал увидеть инструкцию IADD, так как я снова читаю data. Кажется, я неправильно написал инструкцию добавления asm, но больше не знаю.

Answer 1

Графические процессоры 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 и выше поддерживает вытеснение на уровне команд. Прерывание уровня инструкций приведет к неожиданным результатам.