Fermi GPU (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，每个SM中有32个标量处理器，所以上面的代码将在一个SM中执行10次。正如我预期的那样，执行的指令数应该是 10*13 = 130。但是经过我的分析，结果是：发出的指令=130，执行的指令=100。 1. 为什么发出的指令数与执行的指令数不同？没有分支，所以它们不应该是相等的吗？ 2. 为什么执行的指令数比预期少？ ptx版本中的所有指令至少应该执行吗？ 3. 高速缓存未命中（L1 和 L2）对发出的指令数和执行的指令数有影响吗？ 谢谢

Answer 1

PTX只是编译代码的中间表示。这并不是 GPU 实际执行的内容。还有一个进一步的汇编步骤，它发出 GPU 运行的代码，这可以在编译时发生，也可以在驱动程序中使用 JIT 编译。因此，您的指令很重要，您从中推断出的任何内容都是无效的。

NVIDIA 推出了一个名为 cuobjdump 的工具，它可以反汇编为 Fermi 卡生成的汇编器输出，并显示在 GPU 上运行的实际机器代码

Answer 2

请记住，PTX 并不完全是 GPU 上执行的内容。 PTX 只是一种中间表示。真正的代码位于 .cubin 文件中。
这就是为什么基于 ptx 源代码进行如此准确的计算毫无意义。

您可以使用 CUDA 4.0 提供的 cuobjdump --sass 工具将 .cubin 文件中的 GPU 汇编代码提取为更具可读性的内容。