问一个是否需要规约的问题

有一个下列结构的计算问题:
for( i=0; i<n1; i++ )
{
for( k=0; k<n2; k++ )
{
a[i] += b[k];
}
}
其中n1在100-1000之间,n2在500-10000之间。请有经验的版主评估一下是把里面的循环做成核函数规约计算快呢?还是将外面的循环做成核函数,核函数里面用for循环累加快呢?还是有其他的好办法?

楼主如果我是您,我将会选择启动100-1000个block,每个block计算500-10000次加法。可以可以有效的提高速度。

而每个block,您可以选择128这种常见的数据规模,这样128个线程每个线程大约累加8次(128 * 8 约等于1000),然后128个值再在block内部规约(需要log2(128) = 7次规约)。

这样往往可以达到您的卡的访存峰值。(贵算法是访存受限的)

感谢您的深夜来访。
——————————————————————————————————
代为修正了一处笔误。

以及,感谢刚才ICE提醒,忘记了您的多个block中所需要的b都是一样的:

b的大小在500-10000个元素,也就是如果是float的话,在2KB-40KB之间,此时强烈建议您ultilize贵卡的read only cache. 方式如下:
(1)如果您是计算能力3.0的话,请您使用texutre或者surface访问。
(2)如果您的计算能力是3.5的话,请您使用__ldg()访问。
(3)如果您的计算能力是2.0的话,您不需要做任何,已经有48K的L1了。目测效果会很好。

感谢您的深夜来访。

以及,可能还有其他方式更好,但我上文建议的途径简单易用,性能也可能不错。在实现时间上很划算。建议您立刻就做它。

感谢您的深夜来访。

以及,稍微补充一下2#的玫瑰斑竹:

不知道LZ这个二重循环仅仅是随手写的用来示意的还是原本就要这样实现,您的算法中,每次都是将b数组的所有元素全部加一遍,然后累加到每个a[i]元素上。

以及,这个是一个累加的过程,所以最后的结果还和a[i]的初值有一定关系,不过LZ既然可以考虑规约实现,那么应该在算法意图上是可以容忍因为浮点计算顺序不同带来的一些差别,那么直接将b相加并存起来,然后给每个a[i]加上这个和即可,原先的双重循环做了很多重复的工作。

大致如此,提醒LZ注意。

感谢深夜来访,祝您晚安。

是的。ICE这个建议真心NB. 是目前最好的建议了。我刚才忘记考虑这点,深表歉意!
学习ICE好榜样,天天加班不打烊,论坛建设不忘本,服务大家斗志强,1234!

是的。ICE这个建议真心NB. 是目前最好的建议了。我刚才忘记考虑这点,深表歉意!
学习ICE好榜样,天天加班不打烊,论坛建设不忘本,服务大家斗志强,1234!

谢谢啊!妙!实在是妙!感谢玫瑰版主深夜回答.

[
不好意思,只是个示意,b[k]应该为b【i】[k]。谢谢版主们深夜。。。[/i]

没有关系。感谢您的深夜来访。请您立刻实践。

您可以直接参考2#玫瑰斑竹的建议,实现即可。

祝您好运~

再弱弱地问一下:如果n1很大,在5000-1千万之间,而n2很小,在10-60之间,又该采取什么策略?

那相当于有大量的数,每个数都需要累加上少量的几个数,那么建议您这样:
(1)外面的循环改成1个线程负责(而不是一个block)。在这个线程内,累加上几个少量的数,保存回去。
(2)使用shared memory进行协作,提供1个线程访问这几个少量的数时候的效率。

其中第二点,如果这几个少量数很少很少,例如,每个数都需要加上2个数的时候,直接读取也行。(相当于每个线程间隔1个元素读取,虽然不是很合并,但也凑合)。

如果这个少量继续退化成1个数,那么就直接变成了两个向量的加法了,那就直接加吧。新建一个CUDA项目,直接编译即可(因为代码已经有了)。

感谢来访。

LZ您好:

这种情况您可以选择一个warp完成一个内循环的工作,依然采用规约实现,和之前的一个block实现一个内循环类似。

同时一个block还是选择典型线程数量大小,一个block计算多个内循环。

这样可以避免block太小,导致occupancy下降。

其他情况不变,以及,这个还是访存密集型的。

祝您好运~