performance - CUDA，相同工作的更多线程 = 尽管占用率更高，但运行时间更长，为什么？

Question

我遇到了一个奇怪的问题，即通过增加线程数量来增加占用率会降低性能。

我创建了以下程序来说明问题:

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cutil.h>

__global__ void less_threads(float * d_out) {
    int num_inliers;
    for (int j=0;j<800;++j) {
        //Do 12 computations
        num_inliers += j*(j+1);
        num_inliers += j*(j+2);
        num_inliers += j*(j+3);
        num_inliers += j*(j+4);
        num_inliers += j*(j+5);
        num_inliers += j*(j+6);
        num_inliers += j*(j+7);
        num_inliers += j*(j+8);
        num_inliers += j*(j+9);
        num_inliers += j*(j+10);
        num_inliers += j*(j+11);
        num_inliers += j*(j+12);
    }

    if (threadIdx.x == -1)
        d_out[threadIdx.x] = num_inliers;
}

__global__ void more_threads(float *d_out) {
    int num_inliers;
    for (int j=0;j<800;++j) {
        // Do 4 computations
        num_inliers += j*(j+1);
        num_inliers += j*(j+2);
        num_inliers += j*(j+3);
        num_inliers += j*(j+4);
    }
    if (threadIdx.x == -1)
        d_out[threadIdx.x] = num_inliers;
}


int main(int argc, char* argv[])
{
    float *d_out = NULL;
    cudaMalloc((void**)&d_out,sizeof(float)*25000);

    more_threads<<<780,128>>>(d_out);
    less_threads<<<780,32>>>(d_out);


    return 0;
}

PTX输出为:

    .entry _Z12less_threadsPf (
        .param .u32 __cudaparm__Z12less_threadsPf_d_out)
    {
    .reg .u32 %r<35>;
    .reg .f32 %f<3>;
    .reg .pred %p<4>;
    .loc    17  6   0
 //   2  #include <stdlib.h>
 //   3  #include <cuda_runtime.h>
 //   4  #include <cutil.h>
 //   5  
 //   6  __global__ void less_threads(float * d_out) {
$LBB1__Z12less_threadsPf:
    mov.s32     %r1, 0;
    mov.s32     %r2, 0;
    mov.s32     %r3, 0;
    mov.s32     %r4, 0;
    mov.s32     %r5, 0;
    mov.s32     %r6, 0;
    mov.s32     %r7, 0;
    mov.s32     %r8, 0;
    mov.s32     %r9, 0;
    mov.s32     %r10, 0;
    mov.s32     %r11, 0;
    mov.s32     %r12, %r13;
    mov.s32     %r14, 0;
$Lt_0_2562:
 //<loop> Loop body line 6, nesting depth: 1, iterations: 800
    .loc    17  10  0
 //   7     int num_inliers;
 //   8     for (int j=0;j<800;++j) {
 //   9         //Do 12 computations
 //  10         num_inliers += j*(j+1);
    mul.lo.s32  %r15, %r14, %r14;
    add.s32     %r16, %r12, %r14;
    add.s32     %r12, %r15, %r16;
    .loc    17  11  0
 //  11         num_inliers += j*(j+2);
    add.s32     %r17, %r15, %r12;
    add.s32     %r12, %r1, %r17;
    .loc    17  12  0
 //  12         num_inliers += j*(j+3);
    add.s32     %r18, %r15, %r12;
    add.s32     %r12, %r2, %r18;
    .loc    17  13  0
 //  13         num_inliers += j*(j+4);
    add.s32     %r19, %r15, %r12;
    add.s32     %r12, %r3, %r19;
    .loc    17  14  0
 //  14         num_inliers += j*(j+5);
    add.s32     %r20, %r15, %r12;
    add.s32     %r12, %r4, %r20;
    .loc    17  15  0
 //  15         num_inliers += j*(j+6);
    add.s32     %r21, %r15, %r12;
    add.s32     %r12, %r5, %r21;
    .loc    17  16  0
 //  16         num_inliers += j*(j+7);
    add.s32     %r22, %r15, %r12;
    add.s32     %r12, %r6, %r22;
    .loc    17  17  0
 //  17         num_inliers += j*(j+8);
    add.s32     %r23, %r15, %r12;
    add.s32     %r12, %r7, %r23;
    .loc    17  18  0
 //  18         num_inliers += j*(j+9);
    add.s32     %r24, %r15, %r12;
    add.s32     %r12, %r8, %r24;
    .loc    17  19  0
 //  19         num_inliers += j*(j+10);
    add.s32     %r25, %r15, %r12;
    add.s32     %r12, %r9, %r25;
    .loc    17  20  0
 //  20         num_inliers += j*(j+11);
    add.s32     %r26, %r15, %r12;
    add.s32     %r12, %r10, %r26;
    .loc    17  21  0
 //  21         num_inliers += j*(j+12);
    add.s32     %r27, %r15, %r12;
    add.s32     %r12, %r11, %r27;
    add.s32     %r14, %r14, 1;
    add.s32     %r11, %r11, 12;
    add.s32     %r10, %r10, 11;
    add.s32     %r9, %r9, 10;
    add.s32     %r8, %r8, 9;
    add.s32     %r7, %r7, 8;
    add.s32     %r6, %r6, 7;
    add.s32     %r5, %r5, 6;
    add.s32     %r4, %r4, 5;
    add.s32     %r3, %r3, 4;
    add.s32     %r2, %r2, 3;
    add.s32     %r1, %r1, 2;
    mov.u32     %r28, 1600;
    setp.ne.s32     %p1, %r1, %r28;
    @%p1 bra    $Lt_0_2562;
    cvt.u32.u16     %r29, %tid.x;
    mov.u32     %r30, -1;
    setp.ne.u32     %p2, %r29, %r30;
    @%p2 bra    $Lt_0_3074;
    .loc    17  25  0
 //  22     }
 //  23  
 //  24     if (threadIdx.x == -1)
 //  25         d_out[threadIdx.x] = num_inliers;
    cvt.rn.f32.s32  %f1, %r12;
    ld.param.u32    %r31, [__cudaparm__Z12less_threadsPf_d_out];
    mul24.lo.u32    %r32, %r29, 4;
    add.u32     %r33, %r31, %r32;
    st.global.f32   [%r33+0], %f1;
$Lt_0_3074:
    .loc    17  26  0
 //  26  }
    exit;
$LDWend__Z12less_threadsPf:
    } // _Z12less_threadsPf

    .entry _Z12more_threadsPf (
        .param .u32 __cudaparm__Z12more_threadsPf_d_out)
    {
    .reg .u32 %r<19>;
    .reg .f32 %f<3>;
    .reg .pred %p<4>;
    .loc    17  28  0
 //  27  
 //  28  __global__ void more_threads(float *d_out) {
$LBB1__Z12more_threadsPf:
    mov.s32     %r1, 0;
    mov.s32     %r2, 0;
    mov.s32     %r3, 0;
    mov.s32     %r4, %r5;
    mov.s32     %r6, 0;
$Lt_1_2562:
 //<loop> Loop body line 28, nesting depth: 1, iterations: 800
    .loc    17  32  0
 //  29     int num_inliers;
 //  30     for (int j=0;j<800;++j) {
 //  31         // Do 4 computations
 //  32         num_inliers += j*(j+1);
    mul.lo.s32  %r7, %r6, %r6;
    add.s32     %r8, %r4, %r6;
    add.s32     %r4, %r7, %r8;
    .loc    17  33  0
 //  33         num_inliers += j*(j+2);
    add.s32     %r9, %r7, %r4;
    add.s32     %r4, %r1, %r9;
    .loc    17  34  0
 //  34         num_inliers += j*(j+3);
    add.s32     %r10, %r7, %r4;
    add.s32     %r4, %r2, %r10;
    .loc    17  35  0
 //  35         num_inliers += j*(j+4);
    add.s32     %r11, %r7, %r4;
    add.s32     %r4, %r3, %r11;
    add.s32     %r6, %r6, 1;
    add.s32     %r3, %r3, 4;
    add.s32     %r2, %r2, 3;
    add.s32     %r1, %r1, 2;
    mov.u32     %r12, 1600;
    setp.ne.s32     %p1, %r1, %r12;
    @%p1 bra    $Lt_1_2562;
    cvt.u32.u16     %r13, %tid.x;
    mov.u32     %r14, -1;
    setp.ne.u32     %p2, %r13, %r14;
    @%p2 bra    $Lt_1_3074;
    .loc    17  38  0
 //  36     }
 //  37     if (threadIdx.x == -1)
 //  38         d_out[threadIdx.x] = num_inliers;
    cvt.rn.f32.s32  %f1, %r4;
    ld.param.u32    %r15, [__cudaparm__Z12more_threadsPf_d_out];
    mul24.lo.u32    %r16, %r13, 4;
    add.u32     %r17, %r15, %r16;
    st.global.f32   [%r17+0], %f1;
$Lt_1_3074:
    .loc    17  39  0
 //  39  }
    exit;
$LDWend__Z12more_threadsPf:
    } // _Z12more_threadsPf

请注意，两个内核总共应该完成相同数量的工作，(如果 threadIdx.x == -1 是阻止编译器优化所有内容并留下空内核的技巧)。工作应该与 more_threads 使用 4 倍多的线程但每个线程做的工作少 4 倍相同。

Profiler 结果的重要结果如下:

more_threads:GPU 运行时间 = 1474 us，reg per thread = 6，occupancy=1，branch=83746，divergent_branch = 26，instructions = 584065，gst request=1084552

less_threads:GPU 运行时间 = 921 us，reg per thread = 14，occupancy=0.25，branch=20956，divergent_branch = 26，instructions = 312663，gst request=677381

正如我之前所说，使用更多线程的内核运行时间更长，这可能是由于指令数量增加所致。

为什么有更多的说明？

为什么会有分支，考虑到没有条件代码，更不用说发散分支了？

在没有全局内存访问的情况下，为什么会有任何 gst 请求？

这是怎么回事!

谢谢

更新

添加了 PTX 代码并修复了 CUDA C，因此它应该可以编译

Answer 1

这两个函数做的工作量不同。

more_threads<<<780, 128>>>():

• 780 block
• 每 block 128 个线程
• 每个循环 4 倍数
• 每个循环添加 8 个
• 780*128*800*(4+8) = 958,464,000 次失败

less_threads<<<780, 32>>>():

• 780 block
• 每个 block 32 个线程
• 每个循环 12 个 mul
• 每次循环添加 24 个
• 780*32*800*(12+24) = 718,848,000 次失败

因此，more_threads 比 less threads 做更多的工作，这就是指令数量增加而 more_threads 变慢的原因。修复 more_threads ，在循环内只做 3 次计算:780*128*800*(3+6) = 718,848,000。