有个问题想请问各位大大

小弟初学cuda,很多东西还不是很熟悉,导师最近给出一段C++下的用傅里叶变换计算相场的代码,先给出代码:
void Initial(fftw_complex dphi){
float x, y, z,rand01,cut;
float A=0.109; //float A=4.0
(-2.0u0+sqrt(-5epsilon-11u0u0))/15.0;
for(int l=0; l <L; l ++){
for (int m=0; m <N; m ++){
for (int n=0; n <N; n ++) {
x=(l+1-N/2)dx;
y=(m+1-M/2)dx;
z=(n+1-N/2)dx;
cut=0.5
(1.-tanh(0.1
(sqrt(x
x+yy+zz)-60)));
//dphi[n+N*(m+Ml)][0]=cut(4A(cos(qx)cos(qy)+cos(qy)cos(qz)+cos(qz)cos(qx))+u0)+(1-cut)u0;
dphi[n+N
(m+M
l)][0]=cut8Acos(qx)cos(qy)cos(qz)+u0;
rand01=float(rand())/RAND_MAX;
dphi[n+N*(m+Ml)][0]=(1.+(0.002rand01-0.001))dphi[n+N(m+Ml)][0]; //实部
dphi[n+N
(m+M*l)][1]=0;//虚部
}
}
}
}

这个是用FFTW库写的,并非 CUFFT,我想问一下,这里三重循环,要划分线程网格的话,该怎么弄?block搞成三维的吗?

楼主您好,您的循环看上去彼此无关的。

您应该可以直接展开这多重循环。
但是其中的一些循环控制变量可能超过block的表示范围,(例如是一个2048 * 2048 * 2048)的循环,您的block只能有512线程,那么您可以多少考虑拆解成:
(512,1,1) * (4, 2048, 2048)这样的。

以及,本论坛的ICE是FFT的高手。您可以等待ICE来详细回复此问题。

LZ您好:

大致看了您1#的代码,正如2#横扫斑竹所言,您的循环是互不影响的,直接展开即可实现对此代码的CUDA并行化。

至于网格的划分,这个并无太多特别指出,您也无需将block搞成3维的。您可以使用一维或者二维的block形状,用多个block拼为对应于原循环的总体形状即可。以及横扫斑竹给出您了一个例子,供您参考。

此外,对于您叙述的“这个是用FFTW库写的”,似乎代码中都是您使用常见的三角函数和其他数学函数实现的,并未见到FFTw函数,只有一个参数类型是FFTw的。

大致如此,祝您编码顺利~

谢谢斑竹的回复,作为初学者,很多东西不懂,斑竹给出的(512,1,1)是一维的block吗?然后(4,2048,2048)是什么呢?望斑竹耐心回答,谢谢啦~~

LZ您好,前者是一维的block的形状,在x方向由512个thread,因为这个参数是DIM3类型的,所以y,z方向的值是1。该block总共有51211=512个threads。
后者是grid里面block排列的形状,是三维排列的,x方向是4个block,y方向是2048个,z方向也是2048个,一共有420482048个blocks。

大致这样,您也可以参阅CUDA C Programming Guide。

祝您好运~

谢谢ice版主的回答:)
我自己想法就是把网格划分为:block(16,16,1),grid((L+16-1)/16,(M+16-1)/16),然后还有N,因为程序使用l,m,n的数值作为输入计算,n用循环控制,写出代码是这样的,望斑竹再帮我看看。

__global__ void INITIALFIVE(cufftDoubleComplex *dphi,float A,int L,int M,int N,float dx,float u0)
{
   float x, y, z,rand01,cut;
	int l,m,n;
	l=blockIdx.x*blockDim.x+threadIdx.x;
	m=blockIdx.y*blockDim.y+threadIdx.y;
	
		if(l>=0&&l<L)
		{
			if(m>=0&&m<M)
			{
				 for(n=0;n<N;n++)
				{
						x=(l+1-N/2)*dx;
						 y=(m+1-M/2)*dx;
						z=(n+1-N/2)*dx;
						cut=0.5*(1.-tanh(0.1*(sqrt(x*x+y*y+z*z)-60)));
					   //dphi[n+N*(m+M*l)][0]=cut*(4*A*(cos(q*x)*cos(q*y)+cos(q*y)*cos(q*z)+cos(q*z)*cos(q*x))+u0)+(1-cut)*u0;
					   dphi[n+N*(m+M*l)][0]=cut*8*A*cos(q*x)*cos(q*y)*cos(q*z)+u0;
					   rand01=float(rand())/RAND_MAX;
					   dphi[n+N*(m+M*l)].x=(1.+(0.002*rand01-0.001))*dphi[n+N*(m+M*l)].x; //实部
					  dphi[n+N*(m+M*l)].y=0;//虚部
				 }
			}
	   }
}

LZ您好:

1:建议您将全部循环拆开的,做法上文已有,不再赘述。

2:您选择了拆开两重循环,保留一重循环,在逻辑上是可行的。

3:您的kernel这样写的话,因为您循环处理不当,将造成GPU访存效率低下,继而影响效能。

4:以及您使用了两重判断来确定干活的threadID,完全可以写到一起,并无需和0比较。

大致这样的,对于3:,建议您了解一下 “合并访问”的相关内容。

祝您编码顺利~

为什么我看高性能编程之cuda,书上说grid只能二维,block可以三维?这里grid是三维的。。。

LZ您好,因为该书成书之时,只有计算能力1.x的硬件,只支持2维的grid。

自fermi架构开始,可以支持3维的grid。

详情请参考CUDA C Programming Guide Appendix.F,以及请参考尽可能新版的官方资料以与时俱进。

祝您好运~

谢谢斑竹的回答,学到不少东西,今天用cufft库来测试,包含了头文件cufft.h为什么链接的时候老是提示错误 3 error LNK2019: 无法解析的外部符号 _cufftPlan1d@16,该符号在函数 _main 中被引用

您还需要链接cufft.lib

请链接:C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\lib\x64\cufft.lib
或者: C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\lib\win32\cufft.lib

如果你是32位项目就链接后者,
如果你是64位项目就链接前者。即可。