global void computer_wavefront(const double zercoef,double wf65,double wavefront,int row,int col,int high)
{
int i=blockIdx.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx = jcol+i;
double wavefront_temp = 0;
double wavefront_temp1[Zs];
if(i<row&&j<col)
{
#pragma unroll
for(int k=2;k<Zs;k++)
{
wavefront_temp1[k] = wf65[k*col*row+idx] * zercoef[k];
}
for(int k=2;k<Zs;k++)
{
wavefront_temp += wavefront_temp1[k];
}
wavefront[idx] = wavefront_temp;
}
}
请问一下我这个kernel中还可以再优化吗？我总觉得每个线程做累加的这部分
for(int k=2;k<Zs;k++) { wavefront_temp += wavefront_temp1[k]; }
可以再优化一下，比如吧这个累加展开，好像没什么效果，希望版主给点建议哈。
LZ您好：
您kernel中的两个循环循环次数什么的完全一样，您先是生成了一组数据存放在wavefront_temp1中，然后读出来累加到wavefront_temp，我觉得您完全可以把两个循环合并，每次计算出来以后直接累加给wavefront_temp。
zercoef[k]是所有线程都是用同样的内容，您不妨考虑使用shared memory缓冲。
大致如此，祝您编码顺利~
我在之前的帖子回复过，楼主做了很多无用功，浪费力气。
例如这个，完全可以直接写成：
for (…)
{
wavefront_temp += wf65[kcolrow+idx] * zercoef[k];
}
（当然，如果你拆分在2个和上累加也行。最后再将这2个和累计成一个和，可以多少提高增加IPC的概率）。
以及，这个和，根据ICE和我的讨论，从数学说无法化简了。
但是您可以考虑将这个数组给安排一下，例如你的zercoef， 可以考虑放入constant memory(如果大小可以的话）。这个也许会提高你的效率点。
至于你询问的乘加，可能真心无法化简了。
版主你好，我原来是一个循环，后来我考虑到要用#pragma unroll 我的理解如果累计的话要迭代，不利于指令并行，所以分开了。不知道这样理解对不对？
还有就是，这里是用share memory 还是 constant memory？怎么用呢？
显然不对。。。。
你这是相当于在做无用功！人家说你每天在办公室里干活太少，于是你就干了一份，撕毁，然后重新干一遍，这样就看起来好了？
不过编译器应该能识别出你这是无用功，然后给你将2个循环体合并的。
我建议楼主还是先了解下CUDA的基本概念，而不是上来就黑写一通。
这也是对您负责。
#pragma unrolld不就是将循环展开吗？如果是一个迭代的for循环是不能展开的吧，所以我将可以展开的东西提到了外面。这是我的想法。
还有你举的这个例子不太理解，我的代码中那一部分是在做无用功，我只是把计算工作量拆分成了两部分，不是重复的两部分工作。
前文说过，
（1）首先是编译器可能自动替你合并。这样导致你的手工拆分是无意义的（等于不拆分）。
（2）而就算不替你合并，你这样也是无意义的。
直接：目标[index] = 值[index] 这样是效率最高的。
而你的拆分形式：
临时目标[index] = 值[index];
和
目标[index] = 临时目标[index];
虽然在（例如编译器将2循环合并，但循环体不组合的情况下），看上去每步的:
目标[t] = 临时目标[t]; 可以和 临时目标[t-1] = 值[t-1];合并同时执行，从而提高了IPC, 但这只是虚假的繁荣。因为他多干了无用功。
这就如同你在办公室里，为了显示你更忙碌，不使用打印机打印文档，而用手抄文档一样。虽然你会更加劳碌，但效率反而更低。
请您考虑这点。如果还是不能接受。我将不能再给出其他建议。
system
10
LZ您好：
目测您似乎没有理解玫瑰斑竹的意思，我来稍微再解释一下：
1：关于迭代能否unroll的问题，您的程序中每个线程自己用自己的数据进行迭代，此时只有自己的数据对自己之前的数据的依赖性，这个完全可以保证的。
换句话说，依据您的写法，您这个循环是可以展开的，无压力，甚至编译器也有可能自动展开。
2：您把一个循环拆成了两部分，假定编译器不为您合并，那么您第一个循环需要保存Zs个数据，然后下一个循环再读取出来，予以累加，一共增加了2*Zs次存储访问，（应该是local memory的访问）。
而如果按照玫瑰斑竹的建议，每次wf65zerceof[]的时候，直接累加到wavefront_temp变量中，则可以避免这2Zs次存储访问，而是直接给保存在寄存器中的wavefront_temp变量，寄存器是最快的，也是延迟最短的，其延迟可以通过线程切换而轻易掩盖，无需担心，因此您直接写就是最好的实现，正如玫瑰斑竹为您指出的那样。
3：您的代码中zercoef这个数组目测是不变的一个系数数组，您现在的写法是每次使用的时候都从global memory中读取，如果访存不是瓶颈则基本无妨，如果访存成为瓶颈，需要优化访存的话，您可以考虑使用其他的高速cache。
如果考虑使用constant cache，那么大致如下写：
全局 定义 constant double zercoef[Zs];
main()中 cudaMemcpyToSymbol()对zercorf赋值，然后kernel内可以直接使用这个数组，无需通过参数引入，是device端的全局数组，同时会被constant cache缓冲，在一次kernel执行中不能改变。
如果考虑使用shard memory，那么大致如下：
在kernel内
shared double zercoef[Zs];
每个block安排一组线程对其赋值，并__syncthreads();
block内部的线程使用自己block的shared memory中缓冲的这份zercoef。
以上两种方法都可以减少从global memory读取数据，在访存密集的情况下可以尝试。
大致如此，祝您好运~
system
11
ice 你好，谢谢的详细解答，其实在我的应用是一个迭代算法，
（1）zercoef是一个长度为65的向量，它所占的字节数为65sizeof（double）= 520B，在每一轮迭代过程中zercoef都要更新，但是在每轮迭代中多个kernel都要用到同一个zercoef。
（2）wf65是一个25625665的一个三维矩阵，你可以把它想象成一个长=256 宽=256 高=65的立方体，这个立方体在我的迭代开始到结束都是固定的，共占用字节数为25625665sizeof（double）= 32MB。
像我这样的应用，我觉得我的应用中对wf65的访问占用时间太多了，zercoef相对要少很多，你觉得这两个变量的访问怎么安排比较合理呢？谢谢啦！
system
12
建议楼主看看编程指导中的SHFL intrinsic相关的例程，对你很有用！
system
13
LZ您好：
看到您最新的补充说明，答复如下：
1：如果您的“每轮迭代”指的是host端
for()
{ kernel 1;
kernel 2;
…
kernel n;
}
以及每次host端的循环中，您的zercoef都是不变的，那么前面给您的对zercoef的处理建议都是可行的，您可以对zercoef使用高速的constant cache或者shared memory进行缓冲。
以及，您这里说的“迭代”和您1#代码中分拆循环的理由中的“迭代”完全无关。您kernel中并无体现出修改zercoef的地方，您kernel内的循环是每个线程独立跑的，完全可以unroll。
system
14
2：根据您1#给出的代码，您对wf65的使用是合并访问的，一般来说初步就可以了。
如果您的算法确实是访存密集型的，那么您可以考虑以下的几点：
a) 您说您的一次迭代中，会有多个kernel，而如果每个kernel都需要读取wf65，那么您可以考虑下，在您的算法框架下，能否将多个kenrel合并为一个kernel，这样可以减少对wf65全局读取的次数。
b) 以及，如果可能，wf65的数据可否在GPU端就地生成，而不用访存读取，一般来说GPU的计算资源是富余的，付出一些计算代价是可以接收的。
c) 一般来说32MB这么多的数据想直接缓冲到GPU的高速cache中是不行的。
大致如此，祝您晚安~