技术开发 频道

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

  【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/

 

0