用户
 找回密码
 立即注册
freshmen 该用户已被删除
发表于 2013-8-23 09:54:13
64415
版主你好
       我的这段代码在处理不同的视频相邻帧的时候,得到的结果,一帧正确一阵错误,用C得出的两帧结果都是正确的
,kernel函数的输入直方图的统计也是正确的。最后输出整形阈值常数,
    为什么处理有的帧结果正确,有的帧结果错误呢,
  您看看是不是那个地方出现bug了
//size=(float)width*height;
  __global__ void thresholdkernel(float*dev_hist,float size,int*threshold)
{
        __shared__ float hist[256];
        __shared__ float avg_hist[256];
        __shared__ float total[256];
        __shared__ int   Site[256];

        float maxVa,p,a;

        int tid=threadIdx.x;
        //write in shared memory
        hist[tid]=dev_hist[tid]/size;
        Site[tid]=tid;
        __syncthreads();

        avg_hist[tid]=(float)tid*hist[tid];
        __syncthreads();

        //计算均值
        for(int mid=blockDim.x/2;mid>0;mid>>=1)
        {
                if(tid<mid)
                {
                        avg_hist[tid]+=avg_hist[tid+mid];
                }
                __syncthreads();
        }
        __syncthreads();

        //
        p=a=0;
        for(int i=0;i<=tid;i++)
        {
                p+=hist[i];
                a+=i*hist[i];
        }
        __syncthreads();
        total[tid]=(avg_hist[0]*p-a)*(avg_hist[0]*p-a)/p/(1-p);
        __syncthreads();

        //求出最大值及位置
        for(int mid=blockDim.x/2;mid>0;mid/=2)
        {
                if(tid<mid)
                {
                        if(total[tid]<total[tid+mid])
                        {
                                total[tid]=total[tid+mid];
                                Site[tid]=Site[tid+mid];
                        }
                        __syncthreads();
                }
                __syncthreads();
        }
        if(tid==0)
              threshold[0]=Site[tid];
        __syncthreads();

}
使用道具 举报 回复
发表于 2013-8-23 10:42:43
LZ您好:

我无法直接推断出您的问题所在,以及您仅提供了kernel的代码,下面就该代码的目测情况做简要说明。需要指出的是,您需要自行保证您算法逻辑的正确性以及host端调用的正确性。

根据您kenrel的写法,您使用了一维的block,每block有256个线程。
您的dev_hist[]是global memory中的数组。
您的写法“hist[tid]=dev_hist[tid]/size;”表明,所有的block都对同样的dev_hist[]的前256个数据做处理,或者您只使用了1个block,这里不太符合一般的习惯,可能是有问题的,请您注意。

下面是对avg_hist[]的初始化和规约求和,目测没有明显问题。

下面是对p和a的计算以及对total[]的赋值,目测没有明显问题。

最后是求最大值和位置,规约部分目测无明显问题,能求出block内的最大值(虽然没有处理相等的情况)但是您将结果回写到global memory的时候“if(tid==0) threshold[0]=Site[tid];”这种写法一般是不正确的,多个block都会对threshold[0]写入,而且还会发生竟写,除非您只使用一个block。

您的代码大致目测分析如上,目测能力有限,仅供参考。
建议您在调试工具的帮助下,对您的代码进行分析,以获得确切的情况。

祝您调试顺利。
使用道具 举报 回复 支持 反对
发表于 2013-8-23 10:55:30
ice 发表于 2013-8-23 10:42
LZ您好:

我无法直接推断出您的问题所在,以及您仅提供了kernel的代码,下面就该代码的目测情况做简要说明 ...

谢谢版主的回答,我的分块就是一个block,来处理256个直方图数据,最后得出一个阈值,
使用道具 举报 回复 支持 反对
发表于 2013-8-23 10:56:03
ice 发表于 2013-8-23 10:42
LZ您好:

我无法直接推断出您的问题所在,以及您仅提供了kernel的代码,下面就该代码的目测情况做简要说明 ...

谢谢版主的回答,我的分块就是一个block,来处理256个直方图数据,最后得出一个阈值,
使用道具 举报 回复 支持 反对
发表于 2013-8-23 10:57:56
ice 发表于 2013-8-23 10:42
LZ您好:

我无法直接推断出您的问题所在,以及您仅提供了kernel的代码,下面就该代码的目测情况做简要说明 ...

不好意思,是我没说清楚
host段代码
bool Threshold(float*hos_hist,int width,int height,int*threshold)
{
        float *dev_hist=NULL;
        int   *dev_threshold=NULL;
        float size=(float)width*height;

        checkCudaErrors(cudaMalloc((void**)&dev_hist,256*sizeof(float)));
        checkCudaErrors(cudaMalloc((void**)&dev_threshold,sizeof(int)));

        checkCudaErrors(cudaMemcpy(dev_hist,hos_hist,256*sizeof(int),cudaMemcpyHostToDevice));

        thresholdkernel<<<1,256>>>(dev_hist,size,dev_threshold);

        checkCudaErrors(cudaMemcpy(threshold,dev_threshold,sizeof(int),cudaMemcpyDeviceToHost));

        checkCudaErrors(cudaFree(dev_hist));
    checkCudaErrors(cudaFree(dev_threshold));
        return 0;
}
使用道具 举报 回复 支持 反对
发表于 2013-8-23 11:55:27
freshmen 发表于 2013-8-23 10:57
不好意思,是我没说清楚
host段代码
bool Threshold(float*hos_hist,int width,int height,int*threshold ...

LZ您好:

未能目测出更多问题,请您以您实机上的调试为准。
如果您深刻怀疑是硬件故障(虽然这个一般来说是小概率事件),请您更换硬件环境继续尝试。

以及,如果您所谓的“错误”是和CPU结果相比较的话,那么作为浮点数的计算,计算顺序不同会轻微影响结果,如果您的算法较为敏感,可能会产生不同的结果,但如果这样的话,一般不能认为是其中一方的计算错误。

大致如此,祝您好运~
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册