请教版主：
我用的是GTX650，共384个核。我给核函数分配了3个block，每个block包含128个线程。每个线程对应处理一个float型数据。把global的数组ad[128]拷到share里的as[128]，
tid=threadIdx.x;
i=threadIdx.x+blockIdx.x*blockDim.x;
as[tid]=ad;
__syncthreads();
…
ad=as[tid];
__syncthreads();
1、启动cuda debugging，程序会自动在as[tid]=ad; 处停下来，弹窗提示memory checker detected 256 access violations.我看了数据传递情况，数据是可以成功传递的，但是为什么会有access violation呢？tid和i不是对应的吗？
2、运行到最后一句ad=as[tid];也有上述提示，并且数据传递不成功。只有前64个数据成功传回，后面64个数据传不回来。我分两次才把数据完整传回来。为什么不能一次传回来呢，指令没写对？
LZ您好：
1：您访存越界。
2：这个基本上是您代码的BUG。
3：请使用“代码模式”提供您的源码，并尽量保留细节。
4：根据已有信息，（并自动脑补修正被论坛转义字符吞掉的[ i ]）可以初步推测是您ad[128]这个数组访存越界，按照您的写法，需要ad[384]。请您仔细考虑您代码的写法。
5：请您先修正越界的问题，再考虑后面copy不成功的问题。
大致如上，祝您调试顺利~
楼主您好：
请您确保您的ad[ i ]（where i = threadIdx.x+blockIdx.x*blockDim.x), 是在有效的下标范围内。因为您说过一共是ad[128], 换言之，您要确保您的i, 也就是您的线程总数<=128, 但根据您的描述，您一共有384个线程。请您检查您的ad [ i ]的下标变换(i = threadIdx.x + blockIdx.x * blockDim.x )符合您的设计原意。
以及，您有两处使用[ i ]的地方，分别是从as读取到ad和从ad读取到as, 这两处同理。
简单的说：下标越界。
解决方案：确定您的下标计算符合您的算法设计。
我明白了，版主的意思是我在分配SHARE的时候要把所有block所需要的share memoy一次性分配对吗？
我记得哪里有一句叫“调用核函数时，主机会自动为每个block生成一个副本。除了地址不一样，副本的其它内容完全一样。”我的理解是我只需分配其中一个block的share money空间，其它block的也会自动生成。是我理解错了吧。
我每天都要吃饭，难道不对吗？
既然是对的，为何程序运行出错？
你不能用一个正确的事实来反对其他和这个事实无关的事情！
回到你的问题，“每个block的shared memory是独立的副本”这个是恒成立的。但是你不能用他排除你下标越界的可能。
请三思。
你一共就ad[128],
然后你用<<<3,128>>>启动kernel, 然后用threadIdx.x + blockIdx.x * blockDim.x做下标，必然会越界。
你说呢！！
但是这并不能否定每个block有自己的一个shared memory数组的副本这个事实。
同理这也无法否定地球绕太阳转的事实，
哪怕你越界了。
请您三思。
LZ您好：
1：您：“版主的意思是我在分配SHARE的时候要把所有block所需要的share memoy一次性分配对吗？”我和横扫两位版主在上述回帖中从来没有这个意思，这是您的理解，以及这个是不正确的理解。
shared memory是block内可见的，所以在申请的时候，只需要申请一个block使用的大小即可。这部分shared memory可以用作block内部生成数据的存储，也可以用作global memory数据的手控缓存。在后面一种情况下，block申请的空间未必是总的global memory中的数据量，可能是，也可能不是。
2：在kernel里面声明的shared memory，每个block将拥有自己的副本，每个副本的大小是一样的，一般来说你需要block对自己的shared memory副本进行初始化，之后使用。
这个横扫斑竹已经详细说明过了。
3：前面已经说过多遍，是您的ad数组越界！这个是个global memory中的数组。不知您为何一直揪着shared memory不放，请您仔细重新考虑下。
在我没有弄懂之前，我还是少说两句了。
这是我模仿sample写的归约程序，如前所述，后面64位不能输出来。
两位版主请不要生气，我的思想没你们的清晰，这是很正常的。
下面是代码，有劳版主大人细心指点。
#include<stdio.h>
define M 12
define N 128
define nf N/2
define blockSize nf
define block_per_grid 3
global void reduction(float matd)
{
int tid=threadIdx.x;
int xr=threadIdx.x+blockIdx.xblockDim.x;
shared float mats[N],mats_sum[nf];
while (xr<MN)
{
mats[tid]=matd[xr];
__syncthreads();
if (tid<nf)
mats_sum[tid]=matd[xr]+matd[xr+blockDim.x/2];
__syncthreads();
if (tid<=32){
mats_sum[tid]+=mats_sum[tid+32];
mats_sum[tid]+=mats_sum[tid+16];
mats_sum[tid]+=mats_sum[tid+8];
mats_sum[tid]+=mats_sum[tid+4];
mats_sum[tid]+=mats_sum[tid+2];
mats_sum[tid]=(1.0/N)(mats_sum[tid]+mats_sum[tid+1]);
}
matd[xr]=abs(mats[tid]-mats_sum[0]);
__syncthreads();
xr+=blockDim.xgridDim.x;
}
}
int main()
{
float mat[M][N];
for (int i=0;i<M;i++)
for (int j=0;j<N;j++)
mat[i][j]=i+j;
float matd;
int d_size=MNsizeof(float);
cudaMalloc(&matd,d_size);
cudaMemcpy(matd,mat,d_size,cudaMemcpyHostToDevice);
reduction<<<block_per_grid,N>>>(matd);
cudaMemcpy(mat,matd,d_size,cudaMemcpyDeviceToHost);
cudaFree(matd);
for (int i=0;i<N;i++)
printf(“%.4f\t”,mat[0][i]);
}
[/i][/i]
没在论坛发过代码，现查的加不管用，见谅。
LZ您好：
大致看了您的代码，在自动修正了斜体转义字符的影响，以及从您迷惑性很强的#define blockSize nf（此时nf=64）但是invoke kernel的时候却使用N（N=128）作为真正的block size这个天坑中爬出来以后，大致发现了您的问题：
根据您的代码，您的mats_sum数组有64个元素，但是您在规约的时候写为“if (tid<=32){
mats_sum[tid]+=mats_sum[tid+32];”，这样将访问到mats_sum[64]也就是第65个元素，即，您访问shared memory越界了。
此外，您使用shared memory缺少相应的同步指令和volatile指定，即便修正了前面的访存问题，计算结果也将是不正确的。您可以参阅programming guide中Synchronization Instruction章节的介绍和示例代码。
请您在保证kernel可以正确运行之后，再考虑copy的问题。
以及，您的8#的问题代码和您1#中所叙述的问题有关么？这分明是两个完全不同的错误，按照您1#的说法，问题在于global memory访存越界，按照您8#给出的代码，是您shared memory访问越界。
请问您8#中“如前所述”含义何在？
快速回复主题的窗口和高级模式的窗口下，工具栏上都有一个“<>”样的工具图标，点开将代码粘贴进去即可。
//我的block里有N个线程，不是nf个。<<<3, 128>>>//在版主大人的指示下参阅了手册，改正tid<32, 增加volatile float* s_ptr//在把128个元素读入share的同时进行一次归约，此时只余下64个元素。根据手册，此时不再需要syncthreads//运行调试，前面运行正确，直到matd[xr]=abs(mats[tid]-mats_sum[0]);发现matd只有前32位得到正确结果//如果是我没有理解清楚您的意思，请批评指正，并请版主大人进一步指示 #include<stdio.h>
#define M 12
#define N 128
#define nf N/2
#define block_per_grid 3
__global__ void reduction(float *matd)
{
int tid=threadIdx.x;
int xr=threadIdx.x+blockIdx.x*blockDim.x;
__shared__ float mats[N],mats_sum[nf];
while (xr<M*N)
{
mats[tid]=matd[xr];
__syncthreads();
if (tid<nf)
mats_sum[tid]=matd[xr]+matd[xr+nf]; //读入共享存储器的同时进行进行一次归约
__syncthreads();
if (tid<32){ //此时mats_sum中只有64个数，不再使用syncthreads
volatile float* s_ptr=mats_sum; //参考了手册，并按您的指示设置改为tid<32,设置volatile
s_ptr[tid] += s_ptr[tid+32];
s_ptr[tid] += s_ptr[tid+16];
s_ptr[tid] += s_ptr[tid+8];
s_ptr[tid] += s_ptr[tid+4];
s_ptr[tid] += s_ptr[tid+2];
s_ptr[tid]=(1.0/N)*(s_ptr[tid]+s_ptr[tid+1]); //完成最后一步归约的同时求出平均值
}
matd[xr]=abs(mats[tid]-mats_sum[0]); //原序列每个元素减去平均值，此时matd中只有前32位得到正确结果
__syncthreads();
xr+=blockDim.x*gridDim.x;
}
}
int main()
{
float mat[M][N];
for (int i=0;i<M;i++)
for (int j=0;j<N;j++)
mat[i][j]=i+j;
float *matd;
int d_size=M*N*sizeof(float);
cudaMalloc(&matd,d_size);
cudaMemcpy(matd,mat,d_size,cudaMemcpyHostToDevice);
reduction<<<block_per_grid,N>>>(matd); //核函数启动格式为<<<3, 128>>>
cudaMemcpy(mat,matd,d_size,cudaMemcpyDeviceToHost);
cudaFree(matd);
for (int i=0;i<N;i++)
printf("%.4f\t",mat[0][i]); //输出第一行结果以检查是否正确
}
//还是没有理解版主说的越界问题，不知道我的思维落入了哪个陷阱
//特此贴出我写的小波分解核函数<<<3,n>>>
//我按版主的意思分析了一下，内部索引tid范围是0～127，mats[tid]没有越界
//外部索引xr的范围是0～383，正好在glbal memory数组cd[m][n]的范围内。
//mats[2n]前n个是原始序列，后n个存放处理结果
//每个block需要share的容量是2n，但分配mats[2n]会提示access violations
#define n 128
#define m 3
#define nmi 7
#define scale 6 //scale分解级数，
#define wlen 6 //wlen小波长度
#define sp_per_block n
#define block_per_grid 3
__global__ void denoise(float *cd)
{
int tid=threadIdx.x;
int xr=threadIdx.x+blockIdx.x*blockDim.x;
int sca[nmi]={128,64,32,16,8,4,2};
int flg[nmi]={0,128,192,224,240,248,252};
float h[wlen]={0.332671,0.806892,0.459878,-0.135011,-0.085441,0.035226};
float g[wlen]={0.035226,0.085441,-0.135011,-0.459878,0.806892,-0.33267};
__shared__ float mats[n*2];//前n个为原序列，后n个单元存放处理结果
float p,q;
int mid;
while (xr<m*n)
{
mats[tid]=cd[xr]; //此处提示access violations
__syncthreads();
for (int j=1;j<scale;j++)
{
if (tid<sca[j-1] && tid%2==0)
{
p=0;
for (int k=0;k<wlen;k++)
{
mid=tid+k;
if (mid>=sca[j-1])
mid-=sca[j-1];
p+=h[k]*mats[flg[j-1]+mid];
q+=g[k]*mats[flg[j-1]+mid];
}
mats[flg[j]+tid/2]=p;
}
}
cd[xr]=mats[tid+n]; //此处提示access violations，并且只有前64个元素传递成功
__syncthreads();
xr += blockDim.x*gridDim.x;
}
}
LZ您好：
看到您12#给出的新的实现，有很大改进！
您之前的invoke kernel时block也是128线程的，只是我当时被您定义的一个叫blocksize的宏暂时迷惑了一阵，该宏也一直没有用到。不过在新版代码中已经修正了这一点。
下面将就您的代码给出问题分析和建议：
1：[优化建议]18行，mats_sum的读入和第一次规约，这个可以直接读前面已经读入并已经同步过的mats的值，而不必去读global memory。当然您这样写逻辑也是正确的。
2：[概念分析]21行，不使用__syncthreads()的直接原因在于只有一个warp干活，而不是数据的多少。同时，在单warp规约，使用shared memory的时候，根据手册说明，需要指定volatile，通过指定volatile保证shared memory不被寄存器缓冲，从而每一步在不使用__syncthreads()的情况下能得到正确的结果。（如果这里不指定volatile，但是每行规约后都加上__syncthreads()也可以，但是会慢一些）
3：[错误分析]29行，在您的if()的单warp规约之后，应该加上__syncthreads()，这个同步不是为了保证shared memory的完整写入，而是为了保证没有参加规约的那些线程先等在这里，等规约结果出来以后再向下执行，否则读到的mats_sum[0]的结果就是错误的。（而参与规约的那个warp的32个线程是真正等规约结束才执行到这里的，这也是您之前只有前32个线程结果正确的原因）
您的问题正在于此，请修正。
4：[其他建议]建议您将22行的“volatile float* s_ptr=mats_sum;”提前放置在“if(tid<32)”之前，使之作用域和生存周期得到扩展，并在30行用“s-ptr[0]”代替原有的“mats_sum[0]”，这样写更为合理一些，但因为需要在29行规约结束后添加__syncthreads()，所以也可以不改，实测结果依然是正确的。
大致建议如上，祝您调试顺利~
LZ您好：
之前说的几条是根据您1#提供的叙述给出的，仅对1#负责。
以及，请您将13#中代码对应的host端部分补充完整，以便看到您的缓冲区分配，kernel调用等的实际情况。
祝您好运~
听君一席话，胜读十一年书！我以为if（tid<32)之后，其它线程就不用干了呢，没想到还在拼命跑啊！
不过还有一个小小的问题，为了便于调试，我上面的归约程序的行数为M=3，但实际使用时M都是很大的。当我把M改成一个比较大的数，如10000，程序就会无法运行，也无法启动调试。弹窗提示为：stack overflow.
程序中是哪条指令暗中使用了这么多堆栈呢？在我浅薄的知识范围里，我只知道调用函数时要使用堆栈。
我发现问题了，这个程序跟上面的归约程序是一样的错误，少了一条syncthreads.使得那些没有被使用的线程不断地跑。改正之后测试成功，也不提示access violations了。就是这条syncthreads困扰了我这么久。
另外我想问一下，13#第35行的tid%2能否有更高效的表述方式。cuda sample中归约优化的例子中，首先就是对这个取余符号进行优化，可见这个指令的运算有多么的慢啊。
还有，下标中嵌套下标，例如mat[flg[j]+tid/2],GPU对这种下标的处理能力与Cpu有没有明显区别。
楼主您的：
int main()
{
float mat[M][N]; //本句将在stack中进行空间分配，VC默认1MB的stack
…
}
此问题实际上和CUDA无关。
感谢来访。
此外，关于您的%2操作是否有更高效的方式的问题，实际上这个就很高效。（在开启了优化以及release之类的配置下将会被编译成&1的）。
关于为何你指出的其他来源资料表示%2很低效，我表示无法赞同此说法。
感谢来访。
同样的，我对您的：
“可见这个指令的运算有多么的慢啊”。
和
“例如mat[flg[j]+tid/2],GPU对这种下标的处理能力与Cpu有没有明显区别”。
的说法表示不赞同。
但论坛依然保留您的原帖，不做删除。（为了鼓励大家踊跃发言，百家争鸣）