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