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

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已经很久没有了，今天傍晚突然心血来潮，决定再重温下曾经的感觉，不过还不错，写的还算顺利。总之，当你将难得学会后，再学容易的，就会有种居高临下的感觉，不但对“它将要做什么，如何做的，有更深的了解和体会“并因此而变的更加容易。 以后又时间还会继续这个话题，内容也将更加详细，限于工作的问题，今天点到即止，如有兴趣讨论相关问题可加我QQ:295553381o(^!^)o

来自 “ ITPUB博客 ” ，链接：http://blog.itpub.net/23056049/viewspace-629406/，如需转载，请注明出处，否则将追究法律责任。

转载于:http://blog.itpub.net/23056049/viewspace-629406/