1

我想我遇到了一个 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 Corporation 基于 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;
}
4

1 回答 1

3

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

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 似乎是相同的,除了tid.{xyz}在版本中使用bug=1(在右侧):

.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 注册开发者网站上提交错误报告。

于 2013-04-19T15:02:43.220 回答