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