我之前提问过,就是在emudebug 模式下,程序能够正确的运行,但是在debug模式下,程序不能正确的运行。运行的时候电脑就蓝屏了,不得不重启。调试的时候错误大概是内存的分配错误。
所以,我想大家帮我看看我的内核代码中到底是怎么个问题,并且指出来,大家相互学习,呵呵。
代码如下:
global void result(unsigned char id, unsigned char itempd , int nMaxWidth, int nMaxHeight, float *dMaxR, float *dSigmaST,float *dSigmaS)
{
long m,n;
float R=0.;
unsigned char P=0;
unsigned char TempP=0;
long y=blockDim.yblockIdx.y+threadIdx.y;//高
long x=blockDim.xblockIdx.x+threadIdx.x;//宽
if(0<=y<256-48)
if(0<=x<256-48)
{
*dSigmaST=0.;
*dSigmaS=0.;
for(m=0;m<48;m++)
{
for(n=0;n<48;n++)
{
P=id[(y+m)256+x+n];
TempP=itempd[m48+n];
*dSigmaS+=(float)(P*P);
*dSigmaST+=(float)(Pixel*TempP);
}
}
R=*dSigmaST/(sqrt(*dSigmaS)*sqrt(32621576));
if(R>*dMaxR)
{
*dMaxR=R;
*nMaxWidth=x;
*nMaxHeight=y;
}
}
}
其中 dim3 dimgrad(16,16,1) ,dimblock(16,16,1); * id 为为256*256 数组的一维表示形式,itempd为大小为4848的数组的一维表示形式。
我想问一哈大家,到底是哪里超限了呢?
急切盼望大家的回答。。。。。
虽然有人跟我说了是内存错误吧,但我总是找不到到底错误在哪里?
要不你把所有的代码和相关的材料给我看看!光这一段代码还真的不能说明太多问题!呵呵!而且我也想不出来你这段代码想干吗?
我可以把我的程序发给你吗?
不过程序要能够运行的话,必须安装opencv。
可以告诉我你的邮箱吗?
再说一下吧,就是相当于找这256256矩阵中和 4848矩阵中相关系数最大的那个矩阵的起始位置。呵呵!谢谢
[ 本帖最后由 snapshotwhu 于 2010-3-11 23:11 编辑 ]
这段代码:
if(0<=y<256-48)
if(0<=x<256-48)
{
我不记得C支持这种语法,不过如果编译通过,那么就只能是
if( ( 0<=y )<256-48 ){
if( ( 0<=x )<256-48 ){
其实等价于
if( 256-48>0 or 1 )
if( 256-48>0 or 1 )
这个逻辑~~~
不应该是这个问题啊,因为在emudebug下可以运行成功啊!
另外其他所有的指针变量我都cudamalloc 了正确的内存。
再说x和y本身就是大于等于0的,我可以不需要前面的那句话啊!
到底是什么问题呢?谢谢楼主了,会不会是有什么同步的问题呢?
id[(255+47)*256 + 255+47]; id 定义的空间多少?:)
对不起LZ那几个分支的逻辑:0<=y<256-48 ,0<=x<256-48 C、C++确实不支持这种语法,我不知道你是怎么编译通过的,在我这里直接被杀
system
10
没有这个问题,x,y的值是不超过 256-48 。
system
11
if( (y< (256-48)) && (x<(256-48))) 这个地方不需要判断大于0就可以
x和y铁定大于0,要不然就是硬件挂了,,,,
dSigmaST=0.0f;
dSigmaS=0.0f;
换成register,用local 会很慢的
system
13
问题是我每计算到一个地方,dSigmaST,dSigmaS 都要置0一次。所以我也是用cudamalloc 分配的。
用cudamalloc 分配的就是 local吗?我记得好像是global 啊!
system
14
这两个都是临时变量,最终也没有记录作用的,所以这两个都是可以直接用register来存储,性能会快很多,cudamalloc分配的是global的内存,会很慢,现在的global和local都是同样的global内存,没有缓冲的。
这个代码最后还有一个问题,就是大家找到的并一不一定都是最大的,因为所有的线程都在写最后的输出的变量,这个地方需要修改一下,可以先得到16*16个,然后再得到最大的哪一个;
system
15
global void result(unsigned char id, unsigned char itempd , int nMaxWidth, int nMaxHeight, float* dMaxR)
{
long m,n;
//float R=0.0f;
unsigned char P=0;
unsigned char TempP=0;
float dSigmaST = 0.0f;
float dSigmaS = 0.0f;
__shared__ float R[16][9];
float nMaxWidth_ = x;
float nMaxHeight_ = y;
long y=blockDim.y*blockIdx.y+threadIdx.y;//高
long x=blockDim.x*blockIdx.x+threadIdx.x;//宽
if( y< (256-48)) && x< (256-48))
{
for(m=0;m<48;m++)
{
for(n=0;n<48;n++)
{
P=id[(y+m)*256+x+n];
TempP=itempd[m*48+n];
dSigmaS+=(float)(P*P);
dSigmaST+=(float)(Pixel*TempP);
}
}
R[threadIdx.y][threadIdx.x]=dSigmaST/(sqrt(dSigmaS)*sqrt(32621576));
}
if (threadIdx.x <8)
{
if (R[threadIdx.y][threadIdx.x] < R[threadIdx.y][threadIdx.x+8])
…
部分代码,其他的自己补充了