在 K20 上没有 -G 选项时 CUDA C 返回不确定且奇怪的结果

Question

我遇到的问题是，我的代码在 K20 上没有 -G 选项时返回不正确的结果。原来的代码太复杂，所以我将其简化以方便阅读和测试。为了在减少代码时保留不正确的结果，有些语句看起来很奇怪，例如:

r_rhs2 = (r_p - r_p)-r_c;

如果我将其替换为

r_rhs2 =-r_c;

得出错误结果的可能性会大大减少。代码首先在主机上分配3个数组(2D，in*jn)，然后将cudaMemcpy分配给设备，然后启动内核“test<<>>”进行一些计算。这是最简化的代码:

#include <stdio.h>
#include <math.h>
#include <time.h>
#include <stdlib.h>

#define in 16
#define jn 16
#define N0 (in*jn*3)
#define N1 (in*jn)

           double  h_dt, ***h_w;
__device__ double  d_dt,   *d_w;



__global__ void test()
{
  int id0,nn;
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  int j = blockIdx.y*blockDim.y + threadIdx.y;
  double r_a, r_b, r_c, r_rhs1, r_rhs2;

  nn = 1;
  id0 = j*in + i;       //id0 is the ID of element in an array
  r_a = d_w[     id0];  //d_w[id0] is the first array
  r_b = d_w[  N1+id0];  //d_w[N1+id0] is the second array
  r_c = d_w[2*N1+id0];  //d_w[2*N1+id0] is the third array

  if (r_b <= 0.0){           //some calculation
    r_rhs1 = 0.0;            //
  }else{                     //
    r_rhs1 =-((r_c)/nn)*r_b; //the "()" increase the error rate
  }

  if(abs(r_b+r_a)>1.e20)printf("weird result=%e,%e,%e,%d,%d\n", r_b+r_a, r_b, r_a, i, j);  //one of the two printf statement, I show the result later that sometimes r_b+r_a is a large number while r_b and r_a are normal

  if (r_b+r_a <= 0.0){          //some calculation
    r_rhs2 = 0.0;               //
  }else{                        //
    r_rhs2 = (r_a - r_a) - r_c; //(r_a - r_a) increase the error rate
  }

  d_w[  N1+id0] = d_w[  N1+id0]+r_rhs2;  //update d_w
  d_w[2*N1+id0] = d_w[2*N1+id0]+r_rhs1;
}

                     
                     
double *** C_mymalloc(int d4, int d3, int d2)
{
  double ***a       = (double ***)malloc(sizeof(double)*d4*d3*d2 + sizeof(double *)*d4*d3 + sizeof(double **)*d4);
  double  **start_l = (double **)a + d4;
  double   *start_k = (double *)a  + d4 + d4*d3;
  int  k, l;


  for (l = 0; l < d4; l++){
    a[l] = start_l + l*d3;
    for (k = 0; k < d3; k++){
      a[l][k] = start_k + l*d3*d2 + k*d2;
    }
  }
  return a;
}


int main()
{
  double *w_m;
  int i,j,n;

  h_w = C_mymalloc(3,jn,in);
  cudaMalloc((void**)&w_m, sizeof(double)*N0);
  cudaMemcpyToSymbol(d_w, &w_m, sizeof(double *));

  for (j = 0; j <= jn-1; j++){
    for (i = 0; i <= in-1; i++){
      h_w[0][j][i] = 1.0;
      h_w[1][j][i] = 0.0;
      h_w[2][j][i] = 1.0/(i*i+j*j+1.0) - 1.0/((i-in)*(i-in)+(j-jn)*(j-jn)+1.0);  //no infinite number
    }
  }
  cudaMemcpy(w_m, h_w[0][0],  N0*sizeof(double), cudaMemcpyHostToDevice);

  dim3 dim_G(1 , 1 );
  dim3 dim_B(16, 16);

  n=0;

  while (n<=10){

    test<<<dim_G, dim_B>>>();

    printf("n=%d\n",n);

    n = n + 1;

    h_dt = 1.0;                                    //Though nonsense, 
    cudaMemcpyToSymbol(d_dt, &h_dt, sizeof(double)); //the weird result
  }                                                //disappear without
                                                   //the two statement

  free(h_w);

  return 0;
}

我的设备:K20； CUDA5编译命令:nvcc main.cu -lm -o exe -archcompute_35 -code sm_35； cuda-memcheck ./exe；典型结果:

n=0
n=1
n=2
n=3
n=4
n=5
n=6
n=7
n=8
dsa=3.741112e+117,-1.938073e-01,1.000000e+00,0,6
dsa=3.449943e+222,-1.859864e-01,1.000000e+00,1,6
dsa=-5.398272e+183,-1.681859e-01,1.000000e+00,2,6
dsa=3.317999e+214,-6.978805e-02,1.000000e+00,6,6
dsa=-2.596131e+264,-4.906721e-02,1.000000e+00,7,6
dsa=-3.011521e+154,-3.072307e-02,1.000000e+00,8,6
dsa=3.665153e+35,0.000000e+00,1.000000e+00,10,6
dsa=5.476628e+246,1.271596e-02,1.000000e+00,11,6
dsa=4.741912e+222,4.596547e-02,1.000000e+00,15,6
n=9
n=10
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors

其他一些方面:

1. 并非每次运行都会产生错误结果，大约有 50% 的概率。
2. 如果加上-G选项，结果是正确的(可能是测试次数不够)
3. 如果加上__syncthread()，结果是正确的
4. 如果我在C2050上运行代码，并将编译命令转换为:nvcc main.cu -lm -o exe -archcompute_20 -code sm_20，结果是正确的

Answer 1

将CUDA5.0更新到CUDA5.5即可解决该问题(我没有尝试过更高版本)。我仍然不知道原因。据 Robert Crovella 称，这可能是 CUDA5.0 中的一个错误。不管怎样，CUDA5.5 工作得很好。特别感谢 njuffa 和 Robert Crovella。