同一算法的直方图均衡化cuda实现结果不正确，而C实现结果是正确的，为什么？

[wenqingkaituozhe]

直方图均衡化的核函数如下，代码可以实现，算法也是对的，但结果却不正确，不知道是什么原因，很费解，希望有高手可以解答一下，谢谢！
__global__ void HistoGramEqualkernel(unsigned char* input_data,unsigned char* output_data, int*histOrg, int*histEqual,int width,int height)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if(x<width && y<height)
{

const unsigned int value=input_data[y*width+x];
atomicAdd(&(histOrg[value]),1);

}
__syncthreads();

float histAdd[256] = {0};//累计直方图
int grayVal[256] = {0};//变换后的灰度值
int notal = width*height;

//累计直方图信息

for (int i=0; i<256; i++)
{
if (i == 0)
{
histAdd[i] = histOrg[i];
}
else
{
histAdd[i] = histAdd[i-1] + histOrg[i];
}
float Val=(histAdd[i]/notal)*255;
grayVal[i] = (int)(Val+0.5);
}
//将变换后的灰度值反映到图像并统计均衡化后的直方图
if(x<width && y<height)
{
output_data[y*width+x] = grayVal[input_data[y*width+x]];

atomicAdd(&(histEqual[output_data[tid]]),1);
}
__syncthreads();
}

[adagio_chen]

错误的结果是什么？

[wenqingkaituozhe]

错误的结果就是帖子里程序后面附上的图像，正确结果如下图：

[adagio_chen]

结果大体上是对的，但是分成了16层。联想到256=16*16，你看看是不是设置grid和block错了

[wenqingkaituozhe]

你讲的我也想过，线程设置如下： #define TILE_WIDTH 16 dim3 dimBlock( TILE_WIDTH, TILE_WIDTH); //线程块的维度 dim3 dimGrid((width+TILE_WIDTH-1)/ TILE_WIDTH, (height+TILE_WIDTH-1)/ TILE_WIDTH); //线程网格的维度 HistoGramEqualkernel<<< dimGrid, dimBlock >>>(Md,Nd,Or,Eq,width, height); //调用内核函数

[随机过程]

直方图累加后，变换前还要进行同步

[wenqingkaituozhe]

在计算累计直方图时是一个串行计算过程，并没有用到线程，加上同步__syncthreads()没有作用的，我也试了，确实无用，现在我有个疑问： 我将串行和并行的变换后的灰度值 grayVal[i] 输出，结果是一样的，说明这之前是对的，我输出output_data[i]时数据不同，说明在最后一步查找映射表错了，但是我不知道为什么？ 即if(x<width && y<height) { output_data[y*width+x] = grayVal[input_data[y*width+x]]; atomicAdd(&(histEqual[output_data[tid]]),1); } 过程出错

[wenqingkaituozhe]

关于这个问题已解决，我用两个核函数来做，一个负责统计并计算灰度映射表，另一个进行新灰度映射并赋值，应该是消除了i（0—256）的影响，这样做出来的结果是对的.

[JeneyWang]

请问楼主您的问题解决了吗？我现在正在学习cuda直方图均衡化，请问您的代码在哪里可以下载，或者可不可以把您的代码发我一份，我想学习一下~感谢感谢~（[e-mail]）