Странные результаты для профилированных выполненных инструкций и выданных инструкций в графическом процессоре Fermi (GTX 580)

Question

Мое ядро ​​имеет версию ptx:

.version 2.2
.target sm_20, texmode_independent

.entry histogram(
        .param .u32 .ptr .global .align 4 histogram_param_0,
        .param .u32 .ptr .global .align 4 histogram_param_1
)
{
        .reg .f32       %f<2>;
        .reg .s32       %r<12>;

_histogram:
        mov.u32         %r1, %tid.x;
        mov.u32         %r2, %envreg3;
        add.s32         %r3, %r1, %r2;
        mov.u32         %r4, %ctaid.x;
        mov.u32         %r5, %ntid.x;
        mad.lo.s32      %r6, %r4, %r5, %r3;
        shl.b32         %r7, %r6, 2;
        ld.param.u32    %r8, [histogram_param_0];
        add.s32         %r9, %r8, %r7;
        ld.param.u32    %r10, [histogram_param_1];
        ld.global.f32   %f1, [%r9];
        add.s32         %r11, %r10, %r7;
        st.global.f32   [%r11], %f1;
        ret;
}

Я, как я посчитал, в моем ядре всего 13 инструкций (не считая инструкции ret). Когда я установил число рабочих элементов равным 5120, размер рабочей группы равен 64. Поскольку имеется 16 SM, в каждом из которых есть 32 скалярных процессора, поэтому приведенный выше код будет выполняться в SM 10 раз. Как я и ожидал, количество выполненных инструкций должно быть 10*13 = 130. Но после того, как я выполнил профилирование, результаты следующие: выданные инструкции =130, выполненные внедрения =100. 1. Почему количество выданных инструкций отличается от количества выполненных инструкций? Ветвей нет, не должны ли они быть равными? 2. Почему количество выполненных команд меньше ожидаемого? Должны ли все инструкции в версии ptx выполняться хотя бы? 3. Влияет ли отсутствие кеша (L1 и L2) на количество выданных инструкций и количество выполненных инструкций? Спасибо

Answer 1

PTX является лишь промежуточным представлением скомпилированного кода. Это не то, что на самом деле выполняет GPU. Существует еще один шаг сборки, который генерирует код, который запускается графическим процессором, это может произойти либо во время компиляции, либо с использованием JIT-компиляции в драйвере. В результате ваша инструкция имеет значение, и все, что вы из них делаете, является недействительным.

NVIDIA выпускает инструмент под названием cuobjdump который может разобрать выходные данные ассемблера, сгенерированные для карт Fermi, и показать фактический машинный код, запущенный на GPU

Answer 2

Имейте в виду, что PTX не совсем то, что выполняется на GPU. PTX - это просто промежуточное представление. Настоящий код находится в файлах.cubin. Вот почему делать такие точные вычисления на основе исходного кода ptx не имеет смысла.

Ты можешь использовать cuobjdump --sass инструмент, поставляемый с CUDA 4.0 для извлечения кода сборки графического процессора из файлов.cubin в нечто более удобочитаемое.