2

我的内核有这样的 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）对发出指令的数量和执行的指令数量有影响吗？谢谢

4

2 回答 2

2

PTX 只是已编译代码的中间表示。这不是 GPU 实际执行的。还有一个组装步骤会发出 GPU 运行的代码，这可以在编译时发生，也可以在驱动程序中使用 JIT 编译。结果，您的指令计数以及您从中推断出的任何内容都是无效的。

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

于 2011-07-06T08:23:11.083 回答
2

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

您可以使用cuobjdump --sassCUDA 4.0 提供的工具从 .cubin 文件中提取 GPU 汇编代码，使其更具可读性。

于 2011-07-06T08:25:06.857 回答