【问题标题】:CUDA bug with threadIdx?threadIdx 的 CUDA 错误?
【发布时间】:2013-04-12 15:57:27
【问题描述】:

我想我遇到了一个 CUDA 错误。有人可以确认/评论代码(见下文)。

代码(附件)将根据“BUG”定义产生不同的结果。 BUG=0 结果为 8(正确),而 BUG=1 结果为 4(错误)。代码的区别只在这里:

#if BUG
unsigned int na=threadIdx.x, nb=threadIdx.y, nc=threadIdx.z;
#else
unsigned int na=0, nb=0, nc=0;
#endif

我只提交了一个线程,所以在这两种情况下都是 na==nb==nc==0 并且我还用语句检查了这一点:

assert( na==0 && nb==0 && nc==0 );
printf("INITIAL VALUES: %u %u %u\n",na,nb,nc);

这是我的编译和运行:

nvcc -arch=sm_21 -DBUG=0 -o bug0 bug.cu
nvcc -arch=sm_21 -DBUG=1 -o bug1 bug.cu
./bug0
./bug1

nvcc:NVIDIA (R) Cuda 编译器驱动程序 版权所有 (c) 2005-2012 NVIDIA 公司建于 Fri_Sep_21_17:28:58_PDT_2012 Cuda 编译 工具,版本 5.0,V0.2.1221

nvcc 使用 g++-4.6 运行

最后是测试代码:

/* Compilation & run
   nvcc -arch=sm_21 -DBUG=0 -o bug0 bug.cu
   nvcc -arch=sm_21 -DBUG=1 -o bug1 bug.cu
   ./bug0
   ./bug1
 */

#include <stdio.h>
#include <assert.h>

__global__
void b(unsigned int *res)
{
#if BUG
    unsigned int na=threadIdx.x, nb=threadIdx.y, nc=threadIdx.z;
#else
    unsigned int na=0, nb=0, nc=0;
#endif

    assert( na==0 && nb==0 && nc==0 );
    printf("INITIAL VALUES: %u %u %u\n",na,nb,nc);

    unsigned int &iter=*res, na_max=2, nb_max=2, nc_max=2;
    iter=0;
    while(true)
    {
        printf("a-iter=%u     %u %u %u\n",iter,na,nb,nc);

        if( na>=na_max )
        {
            na  = 0;
            nb += blockDim.y;

            printf("b-iter=%u     %u %u %u\n",iter,na,nb,nc);

            if( nb>=nb_max )
            {
                printf("c-iter=%u     %u %u %u\n",iter,na,nb,nc);
                nb  = 0;
                nc += blockDim.z;
                if( nc>=nc_max )
                    break;  // end of loop
            }
            else
                printf("c-else\n");
        }
        else
            printf("b-else\n");

        printf("result    %u %u %u\n",na,nb,nc);
        iter++;

        na += blockDim.x;
    }
}

int main(void)
{
    unsigned int res, *d_res;
    cudaMalloc(&d_res,sizeof(unsigned int));
    b<<<1,1>>>(d_res);
    cudaMemcpy(&res, d_res, sizeof(unsigned int), cudaMemcpyDeviceToHost);
    cudaFree(d_res);

    printf("There are %u combinations (correct is 8)\n",res);

    return 0;
}

【问题讨论】:

    标签: cuda nvcc


    【解决方案1】:

    这似乎是一个汇编程序错误。如果我采用您的示例的简化版本:

    template<int bug>
    __global__
    void b(unsigned int *res)
    {
        unsigned int na, nb, nc;
        switch(bug) {
            case 1:
            na=threadIdx.x;
            nb=threadIdx.y;
            nc=threadIdx.z;
            break;
    
            default:
            na = nb = nc = 0;
            break;
        }
    
        unsigned int &iter=*res, na_max=2, nb_max=2, nc_max=2;
        iter=0;
        while(true)
        {
            if( na>=na_max )
            {
                na  = 0;
                nb += blockDim.y;
    
                if( nb>=nb_max )
                {
                    nb  = 0;
                    nc += blockDim.z;
                    if( nc>=nc_max ) break;
                }
            }
    
            iter++;
            na += blockDim.x;
        }
    }
    

    并实例化两个版本,发出的 PTX 似乎是相同的,除了在带有bug=1 的版本中使用了tid.{xyz}(在右侧):

    .visible .entry _Z1bILi0EEvPj(                         .visible .entry _Z1bILi1EEvPj(
            .param .u64 _Z1bILi0EEvPj_param_0                      .param .u64 _Z1bILi1EEvPj_param_0
            )                                                      )
    {                                                      {
        .reg .pred  %p<4>;                                     .reg .pred   %p<4>;
        .reg .s32   %r<28>;                                    .reg .s32    %r<28>;
        .reg .s64   %rd<3>;                                    .reg .s64    %rd<3>;
    
    
        ld.param.u64    %rd2, [_Z1bILi0EEvPj_param_0];         ld.param.u64     %rd2, [_Z1bILi1EEvPj_param_0];
        cvta.to.global.u64  %rd1, %rd2;                        cvta.to.global.u64   %rd1, %rd2;
        mov.u32     %r26, 0;                                   .loc 2 11 1
        .loc 2 22 1                                                mov.u32  %r27, %tid.x;
            st.global.u32   [%rd1], %r26;                      .loc 2 12 1
        .loc 2 33 1                                                mov.u32  %r25, %tid.y;
            mov.u32     %r1, %ntid.z;                          .loc 2 13 1
        .loc 2 28 1                                                mov.u32  %r26, %tid.z;
            mov.u32     %r2, %ntid.y;                          mov.u32  %r24, 0;
        .loc 2 39 1                                            .loc 2 22 1
            mov.u32     %r3, %ntid.x;                              st.global.u32    [%rd1], %r24;
        mov.u32     %r27, %r26;                                .loc 2 33 1
        mov.u32     %r25, %r26;                                    mov.u32  %r4, %ntid.z;
        mov.u32     %r24, %r26;                                .loc 2 28 1
                                                                   mov.u32  %r5, %ntid.y;
    BB0_1:                                                     .loc 2 39 1
        .loc 2 25 1                                                mov.u32  %r6, %ntid.x;
            setp.lt.u32     %p1, %r27, 2;                  
        @%p1 bra    BB0_4;                                 BB1_1:
                                                               .loc 2 25 1
        .loc 2 28 1                                                setp.lt.u32  %p1, %r27, 2;
            add.s32     %r25, %r2, %r25;                       @%p1 bra     BB1_4;
        .loc 2 30 1                                        
            setp.lt.u32     %p2, %r25, 2;                      .loc 2 28 1
        mov.u32     %r27, 0;                                       add.s32  %r25, %r5, %r25;
        .loc 2 30 1                                            .loc 2 30 1
            @%p2 bra    BB0_4;                                     setp.lt.u32  %p2, %r25, 2;
                                                               mov.u32  %r27, 0;
        .loc 2 33 1                                            .loc 2 30 1
            add.s32     %r26, %r1, %r26;                           @%p2 bra     BB1_4;
        .loc 2 34 1                                        
            setp.gt.u32     %p3, %r26, 1;                      .loc 2 33 1
        mov.u32     %r27, 0;                                       add.s32  %r26, %r4, %r26;
        mov.u32     %r25, %r27;                                .loc 2 34 1
        .loc 2 34 1                                                setp.gt.u32  %p3, %r26, 1;
            @%p3 bra    BB0_5;                                 mov.u32  %r27, 0;
                                                               mov.u32  %r25, %r27;
    BB0_4:                                                     .loc 2 34 1
        .loc 2 38 1                                                @%p3 bra     BB1_5;
            add.s32     %r24, %r24, 1;                     
        st.global.u32   [%rd1], %r24;                      BB1_4:
        .loc 2 39 1                                            .loc 2 38 1
            add.s32     %r27, %r3, %r27;                           add.s32  %r24, %r24, 1;
        bra.uni     BB0_1;                                     st.global.u32    [%rd1], %r24;
                                                               .loc 2 39 1
    BB0_5:                                                         add.s32  %r27, %r6, %r27;
        .loc 2 41 2                                            bra.uni  BB1_1;
            ret;                                           
    }                                                      BB1_5:
                                                               .loc 2 41 2
                                                                   ret;
                                                           }
    

    然而,汇编器的输出是另一回事(同样是左侧的 bug=0 和右侧的 bug=1):

        /*0008*/    MOV R1, c [0x0] [0x44];                MOV R1, c [0x0] [0x44];              
        /*0010*/    MOV R6, c [0x0] [0x140];               MOV R6, c [0x0] [0x140];
        /*0018*/    MOV R7, c [0x0] [0x144];               MOV R7, c [0x0] [0x144];
        /*0020*/    S2R R0, SR_Tid_X;                      MOV R0, RZ;
        /*0028*/    MOV R4, RZ;                            MOV R2, RZ;
        /*0030*/    S2R R3, SR_Tid_Z;                      MOV R3, RZ;
        /*0038*/    ST.E [R6], RZ;                         MOV R4, RZ;
        /*0048*/    S2R R2, SR_Tid_Y;                      ST.E [R6], RZ;
        /*0050*/    ISETP.LT.U32.AND P0, pt, R0, 0x2, pt;  ISETP.LT.U32.AND P0, pt, R2, 0x2, pt;
        /*0058*/    SSY 0xd0;                              @P0 BRA 0xb0;
        /*0060*/    @P0 BRA 0xc0;                          IADD R3, R3, c [0x0] [0x2c];
        /*0068*/    IADD R2, R2, c [0x0] [0x2c];           MOV R2, RZ;
        /*0070*/    MOV R0, RZ;                            ISETP.LT.U32.AND P0, pt, R3, 0x2, pt;
        /*0078*/    ISETP.LT.U32.AND P0, pt, R2, 0x2, pt;  @P0 BRA 0xb0;
        /*0088*/    SSY 0xa0;                              IADD R0, R0, c [0x0] [0x30];
        /*0090*/    @P0 BRA 0xc0;                          MOV R2, RZ;
        /*0098*/    IADD.S R3, R3, c [0x0] [0x30];         ISETP.GT.U32.AND P0, pt, R0, 0x1, pt;
        /*00a0*/    ISETP.GT.U32.AND P0, pt, R3, 0x1, pt;  MOV R3, RZ;
        /*00a8*/    MOV R0, RZ;                            @P0 EXIT;
        /*00b0*/    MOV R2, RZ;                            IADD R4, R4, 0x1;
        /*00b8*/    @P0 EXIT;                              IADD R2, R2, c [0x0] [0x28];
        /*00c8*/    IADD.S R4, R4, 0x1;                    ST.E [R6], R4;
        /*00d0*/    ST.E [R6], R4;                         BRA 0x50;
        /*00d8*/    IADD R0, R0, c [0x0] [0x28];           BRA 0xd8;
        /*00e0*/    BRA 0x50;                              NOP CC.T;
        /*00e8*/    BRA 0xe8;                              NOP CC.T;
        /*00f0*/    NOP CC.T;                              NOP CC.T;
        /*00f8*/    NOP CC.T;                              NOP CC.T;     
    

    右边的代码缺少两条SSY 指令,运行它会导致内核处于无限循环中,这与某种 SIMT 正确性问题一致,例如未检测到的分支分歧或同步障碍周围的分歧。真正有趣的是,在单个块中仅运行单个线程时会挂起

    如果我是你,我建议在 NVIDIA 注册开发者网站上提交错误报告。

    【讨论】:

    • 感谢您的检查!我已经提交了(Issue Reference #130418-000028),但是我没有看到 NVidia 团队的任何反应,所以我决定将它发送到这里。 StackOverflow 可以更快地识别问题!
    猜你喜欢
    • 2012-06-09
    • 2012-10-22
    • 2015-09-12
    • 1970-01-01
    • 2013-10-26
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多