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结果中的重要结果如下L：

more_threads：GPU 运行时 = 1474 us，每个线程的注册 = 6，占用 = 1，分支 = 83746，divergent_branch = 26，指令 = 584065，gst 请求 = 1084552

less_threads：GPU 运行时 = 921 us，每线程注册 = 14，占用 = 0.25，分支 = 20956，分歧分支 = 26，指令 = 312663，gst 请求 = 677381

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

为什么会有更多说明？

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

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

这是怎么回事！

谢谢

更新

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

Answer 1

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

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

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

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

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

所以，more_threads 比 less 线程做更多的工作，这就是指令数量增加和 more_threads 变慢的原因。要修复 more_threads，只需在循环内执行 3 次计算：780*128*800*(3+6) = 718,848,000。

Answer 2

由于您的代码只有算术指令，因此您不需要很高的占用率来隐藏算术单元的延迟。事实上，即使您确实有内存指令，只要您的读/写效率很高，您也可以在大约 50% 的占用率下最大限度地提高性能。有关入住率和性能的更多信息，请参阅录制的 Advanced CUDA C 演示文稿。

在您的情况下，鉴于您的内核不需要高占用率来使算术单元饱和，因此使用较少的较大块将比使用更多较小的块获得更好的性能，因为启动块需要成本。不过总体而言，与实际运行代码的时间相比，启动块的成本可以忽略不计。

为什么会有更多说明？

请记住，计数器不是按块计数（也称为 CTA），而是按 SM（流式多处理器）或每个 TPC（纹理处理集群）计数，这是一组两个或三个 SM，具体取决于您的设备。指令计数是每个 SM。

期望less_threads 内核的指令更少是公平的，但是每个块启动的warp 数量是其四倍，这意味着每个SM 将执行代码的次数大约是四倍。考虑到较短的内核代码，您的测量似乎并没有不合理。

为什么会有分支？

其实你确实有条件代码：

for (int j=0;j<800;++j)

这有一个条件，但是一个 warp 中的所有线程确实都在执行相同的路径，因此它不是发散的。我的猜测是管理代码中的某个地方存在分歧，如果您担心，可以查看 PTX 代码来分析这一点。与执行的指令数相比，26 非常低，因此这不会影响您的性能。

为什么会有任何 gst 请求？

在您的代码中：

if (threadIdx.x == -1)
  d_out[blockIdx.x*blockDim.x+threadIdx.x] = num_inliers;

这将由加载/存储单元处理，因此即使它不会导致实际事务也将被计算在内。 gst_32/gst_64/gst_128 计数器指示实际内存传输（您的设备具有 1.2 或 1.3 的计算能力，旧设备具有不同的计数器集）。

Answer 3

1. 两个函数的代码行数不同，所以指令数不同

2. for 循环是使用分支实现的。最后一行代码总是发散的

3. 全局存储请求与全局分数不同。操作已设置，但从未提交。