2013-09-01 80 views
1

我刚问过自己是否用例如threadIdx.x对性能有影响。内核上传到设备后,这些变量是否会保持不变?复杂的线程索引计算是否会影响性能?

我想导航到一个巨大的数组,其索引依赖于threadIdx.x,threadIdx.y和threadIdx.z。我需要像

array[threadIdx.y % 2 + ...] 
+0

每次计算都会产生影响,并且有公共信息每次计算和每次读取或写入所需的周期数。关于像threadIdx这样的内置变量的访问,据我所知,它们可以被视为常量。如果您执行重复计算,只需将这些计算结果保存在寄存器中或信任编译器,并检查您的内核的汇编和gpu资源使用情况(如果您有足够的寄存器,则每个sm的并发块数不会减少)。我更喜欢通过信任编译器来编写干净的优化代码。 – djmj

回答

0

如果有人感兴趣,我评估了相应的PTX代码。

(1)复杂的线程ID计算会影响性能。 “threadIdx.x”等不是常量。 (2)“threadIdx.y%2”有效地实现并且对应于“threadIdx.y & 0x00000001”(Cuda Toolkit 5.5)。

2

模运算我认为

array[threadIdx.y % 2 + ...] 

仅仅是一个例子。

一般来说,%操作可能会很慢。一个有用的技巧,以加快指数计算中指出,

foo%n==foo&(n-1) if n is a power of 2 

所以,也许编译器上面的例子会使这个优化你,但如果你有foo%n,上面的技巧是值得一用是。

3

在索引计算中有一个加法和一个模数。

从CUDA编程指南:operator+的吞吐量非常高(对于3.5计算型GPU为160)。

operator%需要几十个操作,吞吐量类似于operator+

在你的情况下,你正在使用operator%文字常量,编译器很可能会优化它。另外你的常量是两个数字(2)的乘方,所以编译器将用operator&(与operator+相同的吞吐量)替换它。

对应用程序进行配置以避免浪费时间优化算术运算而不获得任何性能,这一点非常重要。 通常情况下,算术运算完全被内存加载和存储操作隐藏,在这种情况下,您需要专注于优化内存吞吐量。

相关问题