请问我在用归约比较一个block中元素的大小，如果我的block设置成一维的也就是dimThread（256,1,1），这样处理简单
就是
tid=threadidx.x;
int i=blockdim/2;
while(i !=0)
{
if(tid <i)
{
s[tid]=s[tid]>s[tid+i]?s[tid]:s[tid+i];
__syncthreads();
i/=2;
}
}
if(tid=0)
max=s[tid];
block是一维的比较好理解，可是如果二维的也就是dimblock（16,16,1）这样的归约该怎么写啊，求指教
LZ您好：
1：您可以审视一下您的程序，看是否必须使用2维形状的block。如果无需这样使用，那么您大可换回您熟悉的一维block形状下的规约操作。
2：如果必须使用上述形状的block，那么规约无非有两个因素，其一是控制好每次的访存，其二是控制好每次干活的线程。您可以考虑下在二维block的情况下，这两点怎么写。
大致建议如上，祝您编码顺利~
2维的你可以用乘法，变换成1维的，然后对您的数组进行规约的。一个道理的。
举个例子说，如果你的kernel是2维的block, 用2维的形状（例如，方便某些算法）进行了计算后，得到一组数据，需要累加求和，
此时可以立刻进行编号变换，组合成1维的，然后对此数组进行规约的。
（以及，反过来，如果一个1维的block, 也可以中途拆分成2维的，然后方便某些计算，您觉得呢？）
谢谢版主的回答
也就是说我的int i改成int i =(blockDim.x*bloclDim.y)/2，这样吧
谢谢您详细的回答
LZ您好：
您可以自行测试下您的设想。
以及您可能还需要参照横扫斑竹的建议，对您的线程编号一维化，再做处理。
大致如此，祝您编码顺利~
直接int i =(blockDim.x*bloclDim.y)/2是无法将您的2维编号的线程给1维化的，但这样可以：
int i = threadIdx.x + blockDim.x * threadIdx.y;
（您原来的int i = blockDim.x * blockDim.y / 2;不可以, 实际上您那样总是得到一个常数的，显然错误的。请您想象为何这个可以)
这一般也叫"flatten"化。
谢谢楼主的回答
我的这个i是归约时的那个偏移，
tid=threadidx.x;
int i=blockdim/2;
while(i !=0)
{
if(tid <i)
{
s[tid]=s[tid]>s[tid+i]?s[tid]:s[tid+i];
__syncthreads();
i/=2;
}
}
每次经i/=2;就是每次都除以2的那个
那个tid才是每个线程的线程号
起始就应该是个常数吧
版主看看是应该这样吗
LZ您好：
我直接告诉您做法好了：
1：您可以保留您的int i =(blockDim.x*bloclDim.y)/2写法，并将这个作为每次判断一维化线程编号的依据。
2：int tid = threadIdx.x + blockDim.x * threadIdx.y;这样来得知每个线程的一维化的线程编号，也就是横扫斑竹告诉您的“flatten化”。
3：以一维化的tid和1：给出的i做判断，确定干活的线程。
4：循环中比较的时候的访存方式不变，不因为比较线程的时候对线程编号进行了一维化处理，而改变这里的实现。比如您原先是每个block对应一个16*16的二维图像块，在全局的图像数据中寻址的时候，可能按照二维索引寻址比较方便，那么这里不因为你将线程编号扁平化用于判断参与规约的线程就将内部的二维寻址改写为一维的，您可以沿用之前的二维写法。
大致如此，供您参考。
祝您好运~
谢谢版主，我照您的说法去编写
LZ您好：
其实最重要的并非是以某个具体的方法去实现，而是您要理解这里面的道理和关系，从而灵活运用，因应万变。版主的说法也只是一家之言，也可能说错，而您要以您自己的实践去检验真理。
这才是钻研和研修之道，愿共勉。
祝您编码顺利~
谢谢版主的鼓励
我按照您的意思编完了，版主看看
global
void NormKernel( float M_src,int width,int height,floatN_max,floatN_min)
{
[b] int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockIdx.y*blockDim.y;
int tid= y*width+x;[/b]
__shared__ float *M_max=new float[width*height];
__shared__ float *M_min=new float[width*height];
M_max[tid]=M_src[tid];
__syncthreads();
int i=(blockDim.x*blockDim.y)/2;
while(i !=0)
{
if(tid <i+blockIdx.x)
{
M_max[tid]=M_max[tid]>M_max[tid+i]?M_max[tid]:M_max[tid+i];
M_min[tid]=M_max[tid]<M_max[tid+i]?M_max[tid]:M_max[tid+i];
__syncthreads();
i/=2;
}
}
if(tid=0)
{
N_max[blockIdx.x+gridDim.xblockIdx.y]=M_max[tid];
N_min[blockIdx.x+gridDim.xblockIdx.y]=M_min[tid];
}
}
版主帮我看看，可能有问题的地方我做了特殊标记
LZ您好：
您在掌握了所有您的算法实现意图的时候，为何不愿意自己想清楚CUDA实现思路并自行理顺一下实现？
您在拥有所有相关信息的时候，为何不愿意自己动手编译并调试一下您的代码？无论出现什么问题，解决问题的过程都将对您的水平提升起到明显的作用。
而是每一步亦步亦趋地跟着斑竹的说法，然后让斑竹为您人肉编译，人肉DEBUG，人肉执行。
这并不是一个良好而高效的方法。
您的代码经过我初步目测，已看出多个有问题或者可能有问题的地方，但在您自己动手编译和调试之前，我拒绝继续直接向您指出这些问题。
祝您编译和调试愉快~