NVCC编译Debug与Release结果不同的解决方案

早些时候我使用NVCC编译代码的时候有些遇见一个很头疼的问题，Debug和Release模式下编译执行的结果竟然不同。下面详细叙述下我的问题和解决方案，以及给诸位同仁的建议。

先贴出我出错的代码：

//Prewitt
__global__ void kernelGrayPrewitt(uchar *src, uchar *dst,unsigned int width, unsigned int heigth, const uchar thresh)
{
  int col = threadIdx.x + blockIdx.x*blockDim.x;  
  int row = threadIdx.y + blockIdx.y*blockDim.y;  

  if (col < 1 || row < 1 || row > heigth - 1 || col > width - 2)  
  {
     return;   
  }
  int offset = row*width + col;    
  int offsetUp = (row-1) * width + col;   
  int offsetDown = (row+1) * width + col;  
  uchar data[3][3] =   
  { {src[offsetUp-1], src[offsetUp], src[offsetUp+1]},     
    {src[offset-1],   src[offset],  src[offset+1]   },    
    {src[offsetDown-1], src[offsetDown], src[offsetDown+1]}};//prepare the data.    
  unsigned int po1 = data[0][2]+data[1][2]+data[2][2];      
  unsigned int po2 = data[0][0]+data[1][0]+data[2][0];      
  uchar dx = po1 > po2? po1-po2:po2-po1;       

  po1 = data[0][0]+data[0][1]+data[0][2];       
  po2 = data[2][0]+data[2][1]+data[2][2];        
  uchar dy = po1 > po2? po1-po2:po2-po1;        
  dst[offset] = dy > dx? dy:dx;     
  if (dst[offset] > thresh)   
  {  
     dst[offset] = 255u;  
  }    
  else  
  {   
     dst[offset] = 0;   
  }   
}  

给出我处理的图像：

1、原始图像

2、Debug模式下编译后的处理结果

3、Release模式下处理结果

上面展示了让人无语的处理结果。当初我遇到这个问题的时候，崩溃了好久。。。

我解决的方法是：将

  unsigned int po1 = data[0][2]+data[1][2]+data[2][2];      
  unsigned int po2 = data[0][0]+data[1][0]+data[2][0]; 

换成了

  int po1 = data[0][2]+data[1][2]+data[2][2];      
  int po2 = data[0][0]+data[1][0]+data[2][0]; 

其实现在我也没有真正弄清楚，为什么三个uchar 类型的数据相加溢出后的结果可以正确的放入int类型的数据中，而放入unsigned int类型的数据中是错误的那。所以我坚定的认为这个是NVCC的一个bug！

继续分析：如果保持源程序不动，在修改以个地方就是最后赋值的地方，将dst[offset] = dy > dx? dy:dx;修改为dst[offset] = dy; 或者dst[offset] = dx; 那么Release结果是正确的，也就是说dx和dy实际上的值是正确的，为什么执行dst[offset] = dy > dx? dy:dx;这一句会有问题那？而且我修改正确的方式也并没有直接修改dy，dx的值。真的让人难以置信！这是不是可以知道它们的执行原理才行那！

今天我只能给出一点建议，那就是尽量少的在CUDA程序中使用类型转化。这个很可能让你发现非常无语的错误的。