1

我有兴趣获得在 GeForce GTX 550 Ti 上执行 1 个双精度 FLOP 所需的纳秒数。

为了做到这一点，我采用了这种方法：我发现卡的单精度峰值性能为 691.2 GFLOPS，这意味着双精度峰值性能将是它的 1/8，即 86.4 GFLOPS。然后为了获得每个核心的 FLOPS，我将 86.4 GFLOPS 除以核心数量 192，得到每个核心 0.45 GFLOPS。0.45 GFLOPS 意味着每个内核每纳秒 0.45 FLOPS。如果我遵循正确的方法，那么我想知道每个内核运行了多少线程来获取这些 GFLOPS 数字，我在哪里可以找到这些信息？

此外，下面显示的我的小测试仅由一个线程执行 236000232 个周期。为了找到执行 1 次循环迭代所需的时间（以纳秒为单位），我做了 236000232/10^6 = 236 个周期。该卡的着色器时钟为 1800Mhz，这意味着执行一次循环需要 236/1.8 = 131 纳秒。这个数字比上面的数字大得多（每个内核 0.45 纳秒）。我确信我在这里遗漏了一些东西，因为数字非常不同。请帮助我理解它背后的数学原理。

    __global__ void bench_single(float *data)
{
    int i;
    double x = 1.;
    clock_t start, end;
    start = clock();
    for(i=0; i<1000000; i++)
    {
        x = x * 2.388415813 + 1.253314137;
    }
    end = clock();
    printf("End and start %d - %d\n", end, start);
    printf("Finished in %d cycles\n", end-start);
}

谢谢，

4

2 回答 2

4

计算能力 2.1 设备的双精度吞吐量为每个周期 4 次操作（如果执行 DFMA，则为 8 次）。这假设所有 32 个线程在分派的 warp 中都处于活动状态。

4 ops/cycle/SM * 4 SMs * 1800 MHz * 2 ops/DFMA = 56 GFLOPS 双倍

该计算假定经线中的所有线程都处于活动状态。

您问题中的代码包含两个可以融合到 DFMA 中的相关操作。使用 cuobjdump -sass 检查程序集。如果您在同一个 SM 上启动多个 warp，则测试会变成对相关指令吞吐量而不是延迟的衡量。

于 2013-02-02T04:46:16.657 回答
3

您应该意识到内核设计存在问题，这意味着您使用此代码进行的任何测量都与双精度指令吞吐量绝对没有关系。

因为包含所有双精度操作的计算循环的结果没有用于内存写入，所以编译器优化会将其删除。CUDA 5 编译器为您的内核发出以下 PTX：

.visible .entry _Z12bench_singlePf(
    .param .u32 _Z12bench_singlePf_param_0
)
{
    .local .align 8 .b8     __local_depot0[8];
    .reg .b32   %SP;
    .reg .b32   %SPL;
    .reg .s32   %r<16>;


    mov.u32     %SPL, __local_depot0;
    cvta.local.u32  %SP, %SPL;
    add.u32     %r3, %SP, 0;
    .loc 2 13 1
    cvta.to.local.u32   %r4, %r3;
    // inline asm
    mov.u32     %r1, %clock;
    // inline asm
    // inline asm
    mov.u32     %r2, %clock;
    // inline asm
    st.local.v2.u32     [%r4], {%r2, %r1};
    cvta.const.u32  %r5, $str;
    // Callseq Start 0
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r5;
    .param .b32 param1;
    st.param.b32    [param1+0], %r3;
    .param .b32 retval0;
    .loc 2 13 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r6, [retval0+0];
    }
    // Callseq End 0
    .loc 2 14 1
    sub.s32     %r7, %r2, %r1;
    cvta.const.u32  %r8, $str1;
    st.local.u32    [%r4], %r7;
    // Callseq Start 1
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r8;
    .param .b32 param1;
    st.param.b32    [param1+0], %r3;
    .param .b32 retval0;
    .loc 2 14 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r9, [retval0+0];
    }
    // Callseq End 1
    .loc 2 15 2
    ret;
}

两条时钟加载指令是相邻的，唯一的其他代码是调用printf. 该 PTX 中没有计算循环。

您应该重新设计内核，以便编译器无法推断出循环结果未使用并对其进行优化。一种方法是这样的：

__global__ 
void bench_single(float *data, int flag=0)
{
    int i;
    double x = 1.;
    clock_t start, end;
    start = clock();
    for(i=0; i<1000000; i++) {
        x = x * 2.388415813 + 1.253314137;
    }
    end = clock();
    printf("End and start %d - %d\n", end, start);
    printf("Finished in %d cycles\n", end-start);

    if (flag) {
        data[blockIdx.x] = x;
    }
}

内核末尾的条件写入阻止了编译器优化循环，所以现在编译器发出这个 PTX：

.visible .entry _Z12bench_singlePfi(
    .param .u32 _Z12bench_singlePfi_param_0,
    .param .u32 _Z12bench_singlePfi_param_1
)
{
    .local .align 8 .b8     __local_depot0[8];
    .reg .b32   %SP;
    .reg .b32   %SPL;
    .reg .pred  %p<3>;
    .reg .f32   %f<2>;
    .reg .s32   %r<28>;
    .reg .f64   %fd<44>;


    mov.u32     %SPL, __local_depot0;
    cvta.local.u32  %SP, %SPL;
    ld.param.u32    %r6, [_Z12bench_singlePfi_param_0];
    ld.param.u32    %r7, [_Z12bench_singlePfi_param_1];
    add.u32     %r10, %SP, 0;
    .loc 2 13 1
    cvta.to.local.u32   %r1, %r10;
    // inline asm
    mov.u32     %r8, %clock;
    // inline asm
    mov.f64     %fd43, 0d3FF0000000000000;
    mov.u32     %r27, 1000000;

BB0_1:
    .loc 2 10 1
    fma.rn.f64  %fd4, %fd43, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd5, %fd4, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd6, %fd5, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd7, %fd6, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd8, %fd7, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd9, %fd8, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd10, %fd9, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd11, %fd10, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd12, %fd11, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd13, %fd12, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd14, %fd13, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd15, %fd14, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd16, %fd15, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd17, %fd16, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd18, %fd17, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd19, %fd18, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd20, %fd19, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd21, %fd20, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd22, %fd21, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd23, %fd22, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd24, %fd23, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd25, %fd24, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd26, %fd25, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd27, %fd26, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd28, %fd27, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd29, %fd28, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd30, %fd29, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd31, %fd30, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd32, %fd31, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd33, %fd32, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd34, %fd33, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd35, %fd34, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd36, %fd35, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd37, %fd36, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd38, %fd37, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd39, %fd38, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd40, %fd39, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd41, %fd40, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd42, %fd41, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd43, %fd42, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    .loc 2 9 1
    add.s32     %r27, %r27, -40;
    setp.ne.s32     %p1, %r27, 0;
    @%p1 bra    BB0_1;

    cvta.to.global.u32  %r5, %r6;
    // inline asm
    mov.u32     %r11, %clock;
    // inline asm
    .loc 2 13 1
    st.local.v2.u32     [%r1], {%r11, %r8};
    cvta.const.u32  %r12, $str;
    // Callseq Start 0
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r12;
    .param .b32 param1;
    st.param.b32    [param1+0], %r10;
    .param .b32 retval0;
    .loc 2 13 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r14, [retval0+0];
    }
    // Callseq End 0
    .loc 2 14 1
    sub.s32     %r15, %r11, %r8;
    cvta.const.u32  %r16, $str1;
    st.local.u32    [%r1], %r15;
    // Callseq Start 1
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r16;
    .param .b32 param1;
    st.param.b32    [param1+0], %r10;
    .param .b32 retval0;
    .loc 2 14 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r17, [retval0+0];
    }
    // Callseq End 1
    .loc 2 16 1
    setp.eq.s32     %p2, %r7, 0;
    @%p2 bra    BB0_4;

    .loc 2 17 1
    cvt.rn.f32.f64  %f1, %fd43;
    mov.u32     %r18, %ctaid.x;
    shl.b32     %r19, %r18, 2;
    add.s32     %r20, %r5, %r19;
    st.global.f32   [%r20], %f1;

BB0_4:
    .loc 2 19 2
    ret;
}

请注意，现在有一个很好的浮点乘加指令流，来自编译器部分展开循环的地方。

正如 Greg Smith 指出的那样，除非您在给定的 SM 上运行足够的 warp 以克服指令调度延迟，否则您不应该期望获得指令吞吐量的真正度量。这可能意味着您将要运行至少一个大块。另请注意， printf 调用将对吞吐量产生很大的负面影响。如果每个块只有一个线程写出其结果，或者（更好）将其存储到全局内存中，您将获得更具代表性的数字。运行大量块，您将获得一些可以平均的测量值。作为最后的检查，您还应该使用反汇编目标代码cudaobjdump确保汇编器不会在时钟读取指令的位置附近移动，否则您所依赖的内核时序会中断。旧版本的汇编程序有指令重新排序的习惯，这可能会破坏插入内核 C 代码或 PTX 的一系列时钟读取的功能。

于 2013-02-02T07:52:35.160 回答