CUDA编程优化--以warp模式和思维实施-技术开发专区

CUDA编程优化--以warp模式和思维实施

作者：王玉圆编辑：王玉圆 2011-07-13 00:05 来源：专家�

　　【IT168 技术】CUDA中(如果在nVIDIA的GPU上，这些技巧同样适用于OpenCL),通常显式的让数据按照warp模式分配执行(指令在硬件层自动按照warp派发),通常可让程序性能优成倍提升。在这个系列中我们将介绍多个以warp mode执行且带来明显性能提升的例子(当然，计算规模要足够大)。废话不多说，贴代码：

.version 1.4
.target sm_10

.extern .shared .u32 smem[];

.entry kernel_reduce( .param.u32 d, .param.u32 s )
{
      .reg.u32  %r<8>;
    .reg.pred %p;

    cvt.u32.u16   %r0, %tid.x;
    shl.b32       %r1, %r0, 2;
    mov.u32       %r7, smem;
    add.u32       %r3, %r7, %r1;
        ld.param.u32  %r4, [ s ];
    add.u32       %r4, %r4, %r1;
    ld.global.u32 %r2, [ %r4 ];
    st.shared.u32 [ %r3 ], %r2;
        mov.u32       %r4, %laneid;

    setp.ge.u32   %p, %r4, 16;
    @%p bra $lable0;
    ld.shared.u32 %r5, [ %r3    ];
    ld.shared.u32 %r6, [ %r3+64 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable0:
    setp.ge.u32   %p, %r4, 8;
    @%p bra $lable1;
    ld.shared.u32 %r5, [ %r3    ];
    ld.shared.u32 %r6, [ %r3+32 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable1:
    setp.ge.u32   %p, %r4, 4;
    @%p bra $lable2;
    ld.shared.u32 %r5, [ %r3    ];
    ld.shared.u32 %r6, [ %r3+16 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable2:
    setp.ge.u32   %p, %r4, 2;
    @%p bra $lable3;
    ld.shared.u32 %r5, [ %r3   ];
    ld.shared.u32 %r6, [ %r3+8 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable3:
    setp.ge.u32   %p, %r4, 1;
    @%p bra $lable4;
    ld.shared.u32 %r5, [ %r3   ];
    ld.shared.u32 %r6, [ %r3+4 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable4:
    bar.sync 0;

    setp.ne.u32  %p, %r4, 0;
        @%p bra $lable5;
    mov.u32       %r5, %warpid;
    shl.b32       %r5, %r5, 2;
    add.u32       %r5, %r5, %r7;
    ld.shared.u32 %r6, [ %r3 ];
    st.shared.u32 [ %r5 ], %r6;

$lable5:
    bar.sync 0;

    setp.ge.u32   %p, %r0, 8;
    @%p bra $lable6;
    ld.shared.u32 %r5, [ %r3    ];
    ld.shared.u32 %r6, [ %r3+32 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable6:
    setp.ge.u32   %p, %r0, 4;
    @%p bra $lable7;
    ld.shared.u32 %r5, [ %r3    ];
    ld.shared.u32 %r6, [ %r3+16 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable7:
    setp.ge.u32   %p, %r0, 2;
    @%p bra $lable8;
    ld.shared.u32 %r5, [ %r3   ];
    ld.shared.u32 %r6, [ %r3+8 ];
    add.u32       %r5, %r5, %r6;
    st.shared.u32 [ %r3 ], %r5;

$lable8:
    setp.ne.u32   %p, %r0, 0;
    @%p bra $lable9;
    ld.shared.u32 %r4, [ %r3   ];
    ld.shared.u32 %r5, [ %r3+4 ];
    add.u32       %r4, %r4, %r5;
    ld.param.u32  %r0, [ d ];
    add.u32       %r0, %r0, %r1;
    st.global.u32 [ %r0 ], %r4;

$lable9:
    exit;
}

　　由于没涉及到浮点数的计算，因此在sm_*<13的情况下不需要使用map_f64_to_f32 指示( >=13那自然更是不用了) 这只是针对一个CTA的，但是同样的规则同样适用于大尺寸数组，这样只不过少写点索引计算的代码而已:) 有一点很不爽，代码贴上后，本来已经对齐看起来赏心悦目的代码全乱了:( 有很多网友都问过我当初是怎么学习CUDA的，由于也不是简单几句话可以说清楚的，且适用于每个人的方法不同，所以我当时的回答也就是敷衍了事。不过今天在这里，我还是说一下吧。我是从比较难的PTX程序开始学起的，甚至当我适应了PTX之后再转入CUDA "C"，编程时却有些不习惯了。再后来开发越来越多的大型项目处于开发时间的考虑就一直用C了，PTX已经很久没有了，今天傍晚突然心血来潮，决定再重温下曾经的感觉，不过还不错，写的还算顺利。总之，当你将难得学会后，再学容易的，就会有种居高临下的感觉，不但对“它将要做什么，如何做的，有更深的了解和体会“并因此而变的更加容易。

更多内容请点击：

CUDA专区：http://cuda.it168.com/

CUDA论坛：http://cudabbs.it168.com/