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

一个都看不懂啊!什么时候也得接触一下PTX了。

PTX级别的优化,性能提高应该是按照50%-1倍左右的性能为提高;对于lz的学习方式,适合实践性的,要是理论性的,可能从算法方面入手更好,好的并行算法,直接放到parallel的系统上就可以优化很多。
就像matrix的乘法运算,高效的乘法运算,其实就跟ibm当年的mpi的集群的算法一样了。
细节扣起来比较麻烦:)一般最后的优化步骤才会考虑去做ptx的优化。
要是用70%的时间得到30%的性能提升个人觉得是不太划算的。
当然写ptx一类的汇编 出神入化了,也可以哈·~呵呵

[ 本帖最后由 OpenHero 于 2010-3-15 13:18 编辑 ]


[[i] 本帖最后由 cyrosly 于 2010-3-15 23:50 编辑 [/i]]

对,基本同意LS的观点,我也是以算法为主,这里这是将在算法确定的情况下,不过当算法已经最优(或者至少你开发人员已经找到中的最优),当遇到local memory使用量太大,而CUDA C不能显式强制制定使用register还是local memory,这种情况下使用PTX还是不错的选择。说点多余的,至于理论性的,就说当初接触CUDA又是硬件架构又是GPU汇编什么的,时间久了给我的感觉就是仅仅多了解了些专业术语名词和脑中多了点更深层次的逻辑和硬件执行模型的理解(可能有些人也会因此有些许的优越感),但实际开发中我看很多人对硬件架构和内部实现逻辑了解的很清楚,实际开发中却不能很好的利用这些,只能说太过钻牛角尖,研究些理论化太深而又几乎没什么实际应用价值的术语,完全没有必要,因为如果彻底的往深层剖析,他们那点理解也只是表面的,于是逐渐从一个CUDA开发人员变成了一个半生不熟的GPU软件和硬件架构人员,而且还是撇脚的。
所以我是很赞同以算法为主的,以及建议新人不要走上类似上面说的那种情况,除非立志成为一位硬件架构师。

[ 本帖最后由 cyrosly 于 2010-3-16 00:04 编辑 ]

学习了

讲的不错,学习了

完全看不懂,好像是汇编代码

汇编和cal差不多?

汇编完全看不懂了