//归约求和运算测试程序，程序运行会卡死，只能重启计算机。求版主花几分钟帮我看看
#include <stdio.h>
#define m 6
#define n 128
global void reduct(float ad,float cach)
{
int tx=threadIdx.x+blockIdx.xblockDim.x;
while (tx<mn)
{
int i=blockDim.x/2;
while (i!=0)
{
if (tx<i)
ad[tx]+=ad[tx+i]; //对矩阵的每行进行归约求和
__syncthreads();
i/=2;
}
tx+=blockIdx.xblockDim.x;
}
__syncthreads();
for (int i=0;i<m;i++)
cach[i]=ad[i*n]; //将每行的第一个元素，即每行的和，存放在cach数组中。
}
void main()
{
float a[m][n],b[m]; //a为待处理矩阵，b为处理结果矩阵。
for (int i=0;i<m;i++)
for (int j=0;j<n;j++)
a[i][j]=1; //每个元素都是1
float ad,cach;
cudaMalloc((void)&ad,mnsizeof(float));
cudaMemcpy(ad,a,mnsizeof(float),cudaMemcpyHostToDevice);
cudaMalloc((void**)&cach,msizeof(float));
reduct<<<3,128>>>(ad,cach); //启用3个block,每个block里面有128个sp
cudaMemcpy(b,cach,m*sizeof(float),cudaMemcpyDeviceToHost);
cudaFree(ad);
cudaFree(cach);
for (int i=0;i<m;i++)
printf(“%f\t”,b[i]);
printf(“\n”);
}
LZ您好：
您规约这么写是不行的，其他逻辑问题暂且不提，您的第一个线程（tx=0）会在您的循环“while (tx<m*n)”中死循环。
请您完整修改您的规约代码。
祝您调试顺利~
我将核函数中的 if (tx<i) 改为 if (tx<(blockIdx.xn+i))，将tx+=blockIdx.xblockDim.x 改为 tx+=3*n.
这样就可以在第一次while循环中实现前三行的归约，但是第二次while循环似乎进不去。
我实在想不到如何实现归约和循环，我是初学者，您能给我一些思路吗？
LZ您好：
您可以参考programming guide中的规约例子。
以及建议您在书写代码之前，先安排好：
1：每次规约有哪些线程参与操作。
2：如何利用这些线程的编号辅助寻址。
您之前的循环写的相当混乱，不建议在此基础上小修小补。
大致如上，祝您中秋快乐~
假如我有两个block,每个block有两个sp.
第一个block里面的sp地址为0，1；那么第二个block里面的sp地址为0,1还是2，3呢？
LZ您好：
您的概念是错误的。
sp=stream processor，又称为cuda core，是硬件概念。
而block下属的是thread，是软件概念。
sp不隶属于任何block。
以及，您可以使用threads的编号来辅助寻址，threads的编号是行优先排列的，您可以一维化为全局的编号，也可以一维化为block内部的编号。
以及他们和访存的寻址并无直接关系。他们和访存的寻址关系取决于您的写法，您需要自己维护一个正确的写法。以及这种写法并非唯一，也无教条，您可以灵活应用。
大致如此，希望您能先理解这些概念，再着手调整您的代码。
祝您中秋快乐~
我参考了SDK里面的reduct 2,但是它是一次性将所有数据归约完，而我需要循环运算才能把矩阵的所有行归约完。实在找不到例子，自己调出了以下核函数，结果正确，但不知道能否更简洁？
global void reduct(float ad,float od)
{
shared float as[mn];
int tid=threadIdx.x;
int bx=blockIdx.x;
int ix=threadIdx.x+blockIdx.xblockDim.x;
while (ix<m*n){
as[threadIdx.x]=ad[ix];
for (int s=blockDim.x/2;s>0;s>>=1){
if (tid<s)
as[tid]+=as[tid+s];
__syncthreads();
}
if (tid==0)
od[bx]=as[0];
ix+=blockDim.x*gridDim.x;
bx+=gridDim.x;
__syncthreads();
}
}
我在想一个问题，我把一行数据从global mem里读到share mem里，然后进行归约求和，再对这个和除以数据个数，就得到这一行的均值。现在我想让这一行的每个元素（未归约前的元素）都减去均值，从而滤掉直流成分。
但是经过归约之后，拷贝到share mem里的序列值已经改变了。我只能从global mem里再将原序列再一次读进来，这样的话效率就很低。另外，由于我每一个元素都要减去一个均值，也就是每一个线程都要同时读取同一个单元，这算是conflict吗？
是否存在一种算法，在归约完成的同时也完了直流成分的滤波？
LZ您好：
大致看了您的代码，您这样的实现是正确的，也是常见做法。
您采取了每个block规约一行的做法，使用block内部的线程index进行线程行为控制，使用全局的线程index进行辅助寻址，这是常规的正确做法。
以及根据您的代码写法，您的线程数量要比您的需要规约的点数要少。您可以使用更多的线程。
同时需要指出的是，规约的写法有很多种，并不限于手册上展示的那种折半规约的方式。具体使用哪种方式，一般要根据算法意图和计算量予以设计。如果得到一个结果数据需要规约计算的数据很多，那么可以先让每个block进行各自线程自己累加的规约形式，当这样进行到需要规约的数据数量等同于线程数量的时候，再进行折半规约。
折半规约每更新一次，干活的线程数量就减少一半，某些情况下并不是十分高效。
LZ您好：
1：如果您的规约结果是整个grid规约后得到的一个全局结果，那么根据这里面的数据依赖性，您必须整个grid结束才能实现这种同步，并得到唯一的结果。这种情况下，您只能重启kernel，继续向下运行。
如果您的规约结果是一个block进行规约得到的结果，那么您可以在block规约结束之后，__syncthreads()，然后整个block的线程都可以安全读取这个规约的结果，然后自己减去直流分量。
因为这个block内部的规约结果需要整个block的线程都读取，那么您需要把数据存放在shared memory（推荐）或者global memory中。
整个warp的线程都读取同一个数据，如果数据在shared memory中，这会引发一次广播，而不会有bank conflict，您可以参考programming guide的说明。如果数据在global memory中，这会引起一次global memory 的访问，以及会浪费访存带宽，因为丢弃了大量数据。
2：如果shared memory容量够用，您可以将规约生成的中间数据写在shared memory中的另外的地方，不破坏原始数据，以便之后直接使用。当然这个具体实现需要根据您的具体问题权衡。
3：我不懂您的信号处理的算法，因而无法回答您是否有算法可以直接完成直流分量滤波。
如果直接按照上边的实现过程考虑的话，我觉得既然您使用规约的结果作为直流分量，那么那么总得得到这个结果，才能计算滤波。如果每步或者每几步都对原始数据进行部分修正，那么总体的计算量可能得不偿失。因为每次修正需要将所有的数据都修正一遍，这样应该不如计算好最后的修正量，只修正一次。
大致如此，预祝您国庆节快乐~
回答得非常仔细，给了我很大帮助，非常感谢您。
不客气的，欢迎您常来论坛。
祝您好运~