cuda 计算结果不正确

system · 2013 年9 月 13 日 14:20

本人根据SDK的矩阵乘的例子，进行修改得到简单的对两个矩阵（145145）做加减运算的程序,wW,wH是矩阵的长和宽，但是发现输出结果不正确，不知道哪里出的问题，希望大家能帮看看。其中的kernel和host程序如下。
global void nvdiCUDA(int Arr, float Op, int wW, int wH)
{ int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockIdx.yblockDim.y;
float value=0;
if(x<=wW&&y<=wH)
{ value=Arr[59wWwH+xwH+Y]+Arr[27wWwH+xwH+Y];
}
Op[xwH+y]=value；
}
nvdi（int *arr, float *op, dim3 &dimsA, dim3 &dimsO）
{ int *d_arr;
float *d_op ; (接着分配空间cudamalloc，数值传递cudaMemcpy…)
threads(blocksize,blocksize);
blocks((dimsA.x+blocksize-1)/blocksize,(dimsA.y+blocksize-1)/blocksize);
nvdiCUDA<<<blocks,threads>>>(d_arr, d_op,dimsA.x,dimsA.y);
cudaFree…
}
其中dimsA(W,H,B); dimsO(W,H,1); W为宽，H为高，B是第几个矩阵数据。blocksize等于16；就是有B个相同大小的矩阵列优先依次存在arr指向的内存空间中，矩阵元素都是非零正整数。现要对第59个矩阵和第27个矩阵进行一下简单运算。但是，我发现输出结果不正确，查看调用运算后的op[0],op[1]始终都是0.000，同时，查看了参照修改的matrixMul例子的结果是正确的，请问斑竹是哪的问题呢？

system · 2013 年9 月 13 日 14:22

楼主您好，

您的代码至少有如下问题：
您在kernel的运行期间进行了cudaFree。

建议您至少应该在cudaFree()前等待（cudaDeviceSynchronize())kernel完成后再尝试释放缓冲区。

否则您提前的缓冲区释放将导致kernel访存失败。

请立刻修正。

感谢深夜来访。

system · 2013 年9 月 13 日 14:26

以及，楼主您还有一个问题：

如果您的wW和wH分别代表宽和高的话，
您这里的：
if(x<=wW&&y<=wH)

可能需要改写成：
if(x<wW&&y<wH)
方能保证安全的不下标越界。

请您考虑这点。

感谢深夜来访。

system · 2013 年9 月 13 日 14:35

LZ您好：

我来稍微补充一下横扫斑竹。

LZ代码中：
if(x<=wW&&y<=wH)
{ value=Arr[59wWwH+xwH+Y]+Arr[27wWwH+xwH+Y];
}
Op[x*wH+y]=value ；

应当将“Op[x*wH+y]=value ”放在if的{}中，这样才合乎算法逻辑。

以及，按照LZ当前的写法，这句赋值对于超出所需范围的线程而言也是执行的，这样很容易出现错误结果乃至直接写越界，导致kernel挂掉。

建议LZ修改此处BUG。

祝您编码顺利~

修正一处笔误，深表歉意。

system · 2013 年9 月 13 日 14:40

非常感谢您，那么及时的给我回复，之前条件是没有等于wW和wH的，但是结果也是不正确，不过您说的对，我会改正过来的。对2#提的问题，是有可能造成问题的原因，因为现在没在那边实验室，程序不在手边，我明早去改了试一下，在与您讨论。非常感谢您的帮助。。

system · 2013 年9 月 13 日 14:46

以及，LZ，根据您的代码实现目测您的矩阵存储是列优先存储顺序的，不知是否这样。

如是，则请无视本楼。
如不是，请检查一下您的实现。

祝您好运~

system · 2013 年9 月 13 日 14:54

啊恩，是的哦，谢谢ice斑竹，我想补充问个问题，怎么样能调试kernel函数呢，在VS2010里无法进入kernel内部函数，也无法查看设备端变量值，就是知道内核函数有问题，可是无法查看该怎么办呢？

system · 2013 年9 月 13 日 14:56

LZ您好：

您可以立即安装nsight 来调试您的device端代码。

祝您调试顺利~

system · 2013 年9 月 14 日 01:56

你好横扫千军斑竹，在调用kernel后补上cudaDeviceSynchronize()的计算结果输出是有值但还是不正确的。。是哪里还有问题吗？

system · 2013 年9 月 14 日 01:58

你好ice斑竹，经过相应修改kernel代码，输出是有值可是结果还是不正确的，不知道哪里出了问题呢

system · 2013 年9 月 14 日 02:03

LZ您好：

添加cudaDeviceSynchronize()仅仅是上述建议之一，就您提供的代码段而言，还有很多其他的BUG，以及在您没有提供的代码中，也可能有BUG。

因此，单独添加cudaDeviceSynchronize()并不能保证您得到正确的结果。

请您根据上述各楼的建议，修正各已知BUG。

祝您好运~

system · 2013 年9 月 14 日 02:04

LZ您好：

如果上述各楼所发现的BUG均已修正，那么请您告知您目前问题的详细情况，以及您修改后最新版代码的详细情况，以便分析。

祝您好运~

system · 2013 年9 月 14 日 02:24

hi,你好，我在主页上无法下载nsight呢，出错显示找不到页面，是怎么回事啊，哪里还可以下载到吗

system · 2013 年9 月 14 日 02:27

LZ您好：

是在这里下载的：

https://developer.nvidia.com/nvidia-nsight-visual-studio-edition

需要注册nvidia 开发者账号（不是本论坛的账号）。

祝您好运~

system · 2013 年9 月 14 日 03:25

因为程序无法拷出来，因此只能在此输一遍

#include<>
…
#define blocksize 16
#define W 145
#define H 145
#define B 220
global void nvdiCUDA(int Arr, float Op, int wW, int wH)
{ int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockIdx.yblockDim.y;
float value=0;
if(x<wW&&y<wH)
{ value=Arr[59wWwH+xwH+Y]+Arr[27wWwH+xwH+Y];
Op[x*wH+y]=value；
}
}
int nvdi（int argc,char argv,int blocksize,int arr, float op, dim3 &dimsA, dim3 &dimsO）
{ int d_arr;
float d_op ;
unsigned int mem_size_A=dimsA.xdimsA.ydimsA.zsizeof(int);
unsigned int mem_size_O=dimsA.xdimsA.ysizeof(float);
cudaEror_t error;
error=cudaMalloc((void**)&d_Arr,mem_size_A);
if(error!=cudaSuccess)
{printf(“cudaMalloc d_Arr return error code %d, line(%d)\n”,error,LINE);
exit(EXIT_FAILURE);
}
error=cudaMalloc((void**)&d_Op,mem_size_O);
if(error!=cudaSuccess)
{printf(“cudaMalloc d_Op return error code %d, line(%d)\n”,error,LINE);
exit(EXIT_FAILURE);
}
error=cudaMemcpy(d_Arr,Arr,mem_size_A,cudaMemcpyHostToDevice);
if(error!=cudaSuccess)
{printf(“cudaMemcpy(d_Arr,Arr) return error code %d, line(%d)\n”,error,LINE);
exit(EXIT_FAILURE);
}

dim3 threads(blocksize,blocksize);
dim3 blocks((dimsA.x+blocksize-1)/blocksize,(dimsA.y+blocksize-1)/blocksize);

nvdiCUDA<<<blocks,threads>>>(d_Arr, d_Op,dimsA.x,dimsA.y);
cudaDeviceSynchronize();
cudaEvent_t start;
error=cudaEventCreat(&start);
if(error!=cudaSuccess)
{printf(stderr,“failed to creat start event(errorcode %s)!\n”,cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
cudaEvent_t stop;
error=cudaEventCreat(&stop);
if(error!=cudaSuccess)
{printf(stderr,“failed to creat stop event(errorcode %s)!\n”,cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
//recore start event
error=cudaEventRecord(start,NULL);
if(error!=cudaSuccess)
{printf(stderr,“failed to record start event(errorcode %s)!\n”,cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
//excute the kernel
int nter=300;
for(int j=0;j<nter;j++){
nvdiCUDA<<<blocks,threads>>>(d_Arr,d_Op,dimsA.x,dimsA.y);
cudaDeviceSynchornize();}
//record stop event
error=cudaEventRecord(stop,NULL);
if(error!=cudaSuccess)
{printf(stderr,“failed to record stop event(errorcode %s)!\n”,cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
//wait for the stop event
error=cudaEventSynchornize(stop);
if(error!=cudaSuccess)
{printf(stderr,“failed to synchronize on the stop event(errorcode %s)!\n”,cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
float msecTotal = 0.0f;
error=cudaEventElapsedTime(&msecTotal,start,stop);
if(error!=cudaSuccess)
{printf(stderr,“failed to get time elapsed between events(errorcode %s)!\n”,cudaGetErrorString(error));
exit(EXIT_FAILURE);
}
float mescPerNvdi=msecTotal/nter;
printf(“time=%.3f”,mescperNvdi);
//copy result
error=cudaMemcpy(Op,d_Op,mem_size_O,cudaMemcpyDeviceToHost);
if(error!=cudaSuccess)
{printf(“cudamencpy(Op,d_Op) returned error code %d,line(%d)\n”,error,LINE);
exit(EXIT_FAILURE);
}
//clean up mem
cudaFree(d_Arr);
cudaFree(d_Op);
cudaDeviceReset();
return 1;
}
int main(int argc, char *argv)
{ intdevID = 0;
// 读设备参数
。。。。。
dim3 dimsA（W,H,B）;
dim3 dimsO(W,H,1);
unsigned int sizea_s,sizeo_s,sizea,sizeo;
int i,array;
float oput;
sizea_s=WHB;
sizeo_s=WH;
sizea=sizea_ssizeof(int);
sizeo=sizeo_ssizeof(float);
array=(int)malloc(sizea);
oput=(float)malloc(sizeo);
if(array==NULL||oput==NULL)
{printf(“allocate mem failed\n”);
return 0;
}
FILE *p;
if((fp=fopen(“rcb.raw”,“r”))==NULL)
{printf(“open file failed\n”);
exit(0);
}
if((fread(array,sizeof(int),sizea_s,fp))==NULL)
{printf(“read data failed\n”);
exit(0);
}
fclose(fp);
int nvdi_result = nvdi(argc,argv,blocksize,array,oput,dimsA,dimsO);
FILE *P;
P=FOPEN(“ndvi,txt”,“w”);
if((fwrite(oput,sizeof(float),sizeo_s,p))==NULL)
{printf(“write file failed\n”); exit(0);
}
fclose(p);
free(array);
free(oput);
exit(nvdi_result);
}

rcb.row是读入的外部文件，查看了一下读入结果，前面数据输入正确，但是在第220个矩阵的最后一列从53-144个数据是错误的并且都为-842150451，查看了一个数据array[14514559]=2480,array[14514527]=4305,但是最后得到op[0]=8.7619483e+033,这个会是怎么回事，还有数据读入array为何最后92个数据输入不对

system · 2013 年9 月 14 日 03:29

你好，我是在这儿下的，并且注册了nvidia的账号，但是选择下载都会跳出页面 Let us help you find the page or asset you were looking for ，The system has reported an “Access Denied " error but don’t worry.。。。。

system · 2013 年9 月 14 日 03:36

LZ您好：
我这里下载似乎无问题的。

[attach]3399[/attach]

system · 2013 年9 月 14 日 03:42

啊恩，我这边却不可以呢，那么您能否可以给我发一份邮箱呢？pinx_bu@163.com ，非常感谢啊。。

system · 2013 年9 月 14 日 03:44

LZ您好：

将随后在本论坛资源版提供，请稍候。

system · 2013 年9 月 14 日 03:46

ok 好的:)