请教版主:
我用的是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]
system
10
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#中“如前所述”含义何在?
system
11
快速回复主题的窗口和高级模式的窗口下,工具栏上都有一个“<>”样的工具图标,点开将代码粘贴进去即可。
system
12
//我的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]); //输出第一行结果以检查是否正确
}
system
13
//还是没有理解版主说的越界问题,不知道我的思维落入了哪个陷阱
//特此贴出我写的小波分解核函数<<<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;
}
}
system
14
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(),所以也可以不改,实测结果依然是正确的。
大致建议如上,祝您调试顺利~
system
15
LZ您好:
之前说的几条是根据您1#提供的叙述给出的,仅对1#负责。
以及,请您将13#中代码对应的host端部分补充完整,以便看到您的缓冲区分配,kernel调用等的实际情况。
祝您好运~
system
16
听君一席话,胜读十一年书!我以为if(tid<32)之后,其它线程就不用干了呢,没想到还在拼命跑啊!
不过还有一个小小的问题,为了便于调试,我上面的归约程序的行数为M=3,但实际使用时M都是很大的。当我把M改成一个比较大的数,如10000,程序就会无法运行,也无法启动调试。弹窗提示为:stack overflow.
程序中是哪条指令暗中使用了这么多堆栈呢?在我浅薄的知识范围里,我只知道调用函数时要使用堆栈。
system
17
我发现问题了,这个程序跟上面的归约程序是一样的错误,少了一条syncthreads.使得那些没有被使用的线程不断地跑。改正之后测试成功,也不提示access violations了。就是这条syncthreads困扰了我这么久。
另外我想问一下,13#第35行的tid%2能否有更高效的表述方式。cuda sample中归约优化的例子中,首先就是对这个取余符号进行优化,可见这个指令的运算有多么的慢啊。
还有,下标中嵌套下标,例如mat[flg[j]+tid/2],GPU对这种下标的处理能力与Cpu有没有明显区别。
system
18
楼主您的:
int main()
{
float mat[M][N]; //本句将在stack中进行空间分配,VC默认1MB的stack
…
}
此问题实际上和CUDA无关。
感谢来访。
system
19
此外,关于您的%2操作是否有更高效的方式的问题,实际上这个就很高效。(在开启了优化以及release之类的配置下将会被编译成&1的)。
关于为何你指出的其他来源资料表示%2很低效,我表示无法赞同此说法。
感谢来访。
system
20
同样的,我对您的:
“可见这个指令的运算有多么的慢啊”。
和
“例如mat[flg[j]+tid/2],GPU对这种下标的处理能力与Cpu有没有明显区别”。
的说法表示不赞同。
但论坛依然保留您的原帖,不做删除。(为了鼓励大家踊跃发言,百家争鸣)