多线程运算问题

加速计算 CUDA
1

下面的代码是实现两个复数相乘并进行归一化，个人对此代码无法理解
ComplexPointwiseMulAndScale<<<32,256>>>(d_signal,d_filter_signal,55.5,1.0f/55.5);

static global void ComplexPointwiseMulAndScale(Complex* a, const Complex* b, int size, float scale)
{
const int numThreads = blockDim.x * gridDim.x;
const int threadID = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = threadID; i < size; i += numThreads)
a[i] = ComplexScale(ComplexMul(a[i], b[i]), scale);
}
请问for循环中起始值int i=threadID和步长为什么这样设置,这样设置的目的是什么？

2

for (int i = threadID; i < size; i += numThreads)
这样的做法可以让相邻线程访问连续的内存空间，防止内存访问的发散，你可以参考Best Practice guide 中合并访问的一章。

你这里应该是a = ComplexScale(ComplexMul(a[i],b[i]),scale);
相邻线层在一次循环之中将访问数组a,b.相邻的元素。

3

我来稍微补充一下2# zehuanwang的回答

文中代码除了为满足合并访问之外，还是一种线程数量少于数据数量时候的常见做法。
假定您决定使用N个线程，但是需要计算的数据有M*N组，那么可以在kernel里面循环M次。
此时，先通过i=threadID来确定寻址的首地址，所有的线程都在算第一块连续的数据，访存是符合合并访问的；然后每次循环的时候整体将寻址位置向后移动一整块（也就是一个numThreads个），然后所有的线程再计算这一块，如此重复M次即可。

另外，建议您以后发代码的时候使用“代码模式”，这样可以有效地避免一些程序中的字符变成论坛发帖的转义字符，从而丢失该字符的显示，并引起其他的显示异常。

祝您编码顺利~

4

版主你好，如果我把new_size的值改为4100，线程和线程块的数量也改了，满足合并访问要求吗？
ComplexPointwiseMulAndScale<<<2050,256>>>(d_signal,d_filter_signal,4100,1.0f/4100);

static global void ComplexPointwiseMulAndScale(Complex* a, const Complex* b, int size, float scale)
{
const int numThreads = blockDim.x * gridDim.x;
const int threadID = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = threadID; i < size; i += numThreads)
a[i]= ComplexScale(ComplexMul(a[i], b[i]), scale);
}

5

LZ您好，您的代码中从未出现过“new_size”变量，所以无法回答您的问题。

一般地，修改问题规模并不会改变实现时局部的行为，所以不会影响合并访问。

关于合并访问的其他细节，请参考2#，原厂支持 zehuanwang给您建议的手册章节。

祝您编码顺利~