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

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

2 голосов
/ 06 июля 2011

У моего ядра версия 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) на количество выданных инструкций и количество выполненных инструкций? Спасибо

Ответы [ 2 ]

2 голосов
/ 06 июля 2011

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

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

2 голосов
/ 06 июля 2011

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

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

...