小弟初学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(xx+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+Ml)][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位项目就链接前者。即可。