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

直方图均衡化的核函数如下，代码可以实现，算法也是对的，但结果却不正确，不知道是什么原因，很费解，希望有高手可以解答一下，谢谢！
__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();
}