目标是将c矩阵(大小为num*num，其中num不超过1024)中每一行的元素求和并赋予长度num的数组a。
发散的规约函数：
__global__ void
d_compute(float *a,float *c,int num)
{
int col=blockIdx.x*blockDim.x+threadIdx.x;
int tx=threadIdx.x;
if(col<num*num)
{
for(int i=1;i<num;i*=2)
{
if(tx%(2*i)==0) c[col]+=c[col+i];
__syncthreads(); }
if(tx==0) a[blockIdx.x]=c[col];
}
}
函数配置为d_compute<<<num,num>>>(a,c,num);
而另一种非规约的算法为
__global__ void
vectorSigma(float *a,float *c,int num)
{
int col=blockIdx.x*blockDim.x+threadIdx.x;
int n;
float sum=0;
if(col<num)
{
for(n=0;n<num;n++)
{
sum+=c[n+col*num];
}
a[col]=sum;
}
}
函数配置为vectorSigma<<<1,num>>>(a,c,num);
计算的结果是居然第二种的for循环比第一种的规约速度要快，请问这是哪里出现问题了…应该如何修正…谢谢！
楼主您好，在我回答您的问题之前，麻烦请补充代码1完整。
然后我方可进行任何评论。
我不能自动闹补出您的实际代码。
感谢来访，周末愉快。
感谢版主回复！
抱歉一楼忘了在代码1中添加int tx=threadIdx.x;
以下为1的全部代码
global void
d_compute(float a,float c,int num)
{
int col=blockIdx.xblockDim.x+threadIdx.x;
int tx=threadIdx.x;
if(col<numnum)
{
for(int i=1;i<num;i*=2)
{
if(tx%(2*i)==0) c[col]+=c[col+i];
__syncthreads();
}
if(tx==0) a[blockIdx.x]=c[col];
}
}
楼主您好，您的这个规约是在global memory中进行的，反复的读取，写入global memory, 而且是不合并的，这样会导致严重的访存问题的。
而相比下，您的代码2,虽然也是不合并的，但它至少是只读c的，无频繁读取-修改-回写。而且累加是在寄存器中的进行的(sum), 所以效率较高。
您觉得呢？
的确如版主所说！
这里仍有一点小问题。
修改代码1为共享内存并收敛发散后代码如下：
__global__ void
d_compute(float *a,float *c,int num)
{
__shared__ float s[1024];
int col=blockIdx.x*blockDim.x+threadIdx.x;
int tx=threadIdx.x;
if(col<num*num)
{
s[tx]=c[col];
__syncthreads();
for(int i=num/2;i>0;i/=2)
{
if(tx<i) s[tx]+=s[tx+i];
__syncthreads();
}
if(tx==0) a[blockIdx.x]=s[0];
}
}
此时效率较1L提升了相当多，但是仍比代码2略有差距（代码1为12ms左右，代码2为10ms左右，使用cudaEventRecord()测量）
请问这个差距能否反超？谢谢！周末快乐！
楼主您好， 我没有看懂您的新col的定义，
这个显然不应该再是col(column)了吧。而应该是每个点了。
（因为我看到您写成：col < num * num…)
这个代码极大的提升了，而且应该快的多，为何？
因为他将原来的c[1024 * id + i]的访问循环改写成了合并读取了。
原来的每个线程间，间隔了4KB啊！现在这种合并访问了，效率极大的提升了（增加的shared memory上的操作和这个比起来还是很合算的）。
所以您不应该变慢的。原来直接每个线程间隔4KB的对global memory的读取，现在直接合并了。却还慢。。这不科学。。明日我将详细和ICE讨论下此问题。看看为何会出现你报告的情况。
此外，根据您的kernel启动的方式<<<num,num>>可知， num不应该大于1024的。
那么您的num * num的float数组，最多只有4MB。
而您现在完全合并的读取，然后在shared memory规约加的新代码，却跑了12ms,
这意味着，您的代码的global memory读取，只有4MB/12ms = 0.3GB/S!
这个非常不应该的。您的新代码看起来无问题，却只有0.3GB/S的访存效率。这个真心奇怪的。
我建议楼主先：
（1）检查下您的计时是否正确（请用profiler看，不要自己计算时间，避免你出错）
（2）如果您是debug模式编译的，请改成release编译，前者的“优化”与否无意义。
期待您的回复。
根据您的代码看，您不是问题1就是问题2.
（即，要么你不是release编译的，要么你是，但是你手工计时错误）
请立刻尝试此2个建议。并反馈。
[
的确这个col是一维坐标中的下标，
其实我在5L的意思是，这个采用了共享内存的函数比起使用全局内存的代码1快了很多，但是仍然比代码2的for循环num次略微慢一点（速度比为 共享：全局：for num次=12:28:10）。
如版主所说，的确改成Release编译之后又上去了一个层次。使用profiler后观测如下显示：
[attach]3223[/attach]
共享内存的函数使用6ms,for循环num次使用10ms！
而手动计时的时间也是显示6ms，看来确定为问题二所致！
感谢版主答疑解惑！周末快乐！
system
10
楼主您好，一些常见的优化，例如您这种，i /= 2的，在正常的release下，这是一个极快的操作，
（一般是一条 >>= 1的指令）。但是在debug下，编译器可能老老实实的进行除法，处以2,因为N卡无整数除法指令，需要模拟，您的时间可能全部耗费在这个上面了。
以及，您的代码还是慢。不过也就先这样吧。
如果您要尝试，可以考虑不使用shared memory上的reduction, 而只是将shared memory作为合并访问的用途：您之前是每个线程读每列的，您现在可以，例如说，先上一些线程（至少32个），横向读取，写入到shared memory, 然后里面的线程们可以从shared memory再次告诉的读取。换句话说，用shared memory将原来的没有合并的访问给合并化，但此时就可以放弃使用shared memory到底了（无需在上面规约，不必用了就要用到低，什么都在上面干）。（或者说，你将shared memory只用来做“转置”而不是在上面拼命读写，计算到底）。
您觉得呢？