Почему Cuda Nsight показывает неправильное содержание структуры после возвращения из функции?

Question

Для отладки мне нужно увидеть значения структуры.Я новичок в Cuda, поэтому возможны ошибки при использовании отладчика Nsight.Моя установка - Developper Studio 2015, Cuda 9.

Я установил точку останова в Nsight в функции.Внутри этой функции все выглядит хорошо.Функция называется TraceTheRay_GPU ().Он заполняет структуру под названием «AbstractIntersection».Пока я остаюсь внутри значений функции для normal, normal_geom и т. Д. Хороши.

Когда я выхожу, Nsight показывает случайные значения для структуры.В чем может быть причина?

Визуализация самой структуры работает, поэтому, когда я показываю нормаль как falsecolor в glview, она выглядит хорошо.Это просто Nsight, который отображает неправильные значения.

Внутри функции значения (например):

isect->point        { x = -57.28 , y = - 53.74, z = 0}
isect->normal       { x = 0,  y = 0,  z = 1}
isect->normal_geom  { x = 0,  y = 0,  z = 1}
isect->dist         126.87

Выход из них:

is.point        { x = -149.02 , y = -24.74, z = -56.82}
is.normal       { x = 95.28,  y = 24.74,  z = 2.8e-45}
is.normal_geom  { x = 0,  y = 0,  z = 0}
is.dist         100

Это выглядиткак Nsight видит другой кусок памяти.Вот еще код, включая определение структуры и способ вызова функции.

// pseudo code:

kernel function() {


  AbstractIntersection is;

  res = TraceTheRay_GPU(r, num_tri, &is);

  // looking at "is"  inside Nsight shows wrong values
  // looking at "is"  inside TraceTheRay_GPU shows good values 
  // displaying is.normal as false color shows a good normal.

}


// real code



struct AbstractIntersection {
    float3 point;
    float3 normal;
    float3 normal_geom;
    float dist;
};


__device__ int TraceTheRay_GPU(const Ray &r, const int number_of_triangles , AbstractIntersection *isect) {

    float min_t = UINT_MAX;
    int hit_index = -1;

    for (int i = 0; i < number_of_triangles; i++)
    {
        float4 v0 = tex1Dfetch(triangle_texture, i * 3);
        float4 e1 = tex1Dfetch(triangle_texture, i * 3 + 1);
        float4 e2 = tex1Dfetch(triangle_texture, i * 3 + 2);

        float t = RayTriangleIntersection(r, make_float3(v0.x, v0.y, v0.z), make_float3(e1.x, e1.y, e1.z), make_float3(e2.x, e2.y, e2.z));

        //float t = 0;

        if (t < min_t && t > 0.001)
        {
            min_t = t;
            hit_index = i;
        }
    }

    if (hit_index > -1) {
        float4 e1 = tex1Dfetch(triangle_texture, hit_index * 3 + 1);
        float4 e2 = tex1Dfetch(triangle_texture, hit_index * 3 + 2);

        isect->normal       = cross(make_float3(e1.x, e1.y, e1.z), make_float3(e2.x, e2.y, e2.z));
        isect->normal       = normalize(isect->normal);
        isect->normal_geom  = isect->normal;  // to start with a valid structure.
        isect->point        = r.ori + r.dir *min_t;
        isect->dist         = min_t; 
        return 1;
    }
    else
        return 0; 

}



__global__ void raytrace( unsigned int *out_data,
                           const int w,
                           const int h,
                           const int number_of_triangles,
                           const float3 a, const float3 b, const float3 c, 
                           const float3 campos,
                           const float3 light_pos,
                           const float3 light_color,
                           const float3 scene_aabb_min, 
                           const float3 scene_aabb_max,
                           curandState  *  devState)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    float xf = (x-0.5)/((float)w);
    float yf = (y-0.5)/((float)h);

    int ray_depth = 0;
    bool inbox = true;

    float3 t1 = c+(a*xf);
    float3 t2 = b*yf;
    float3 image_pos = t1 + t2;
    Ray r(image_pos,image_pos-campos);

    float3 result = make_float3(0, 0, 0);


    AbstractIntersection is; // there is no init at the moment
                           // full code catches this with a test for "res". 
    int res = 0;


    int num_tri; //  = number_of_triangles;
    num_tri = 12;


  res = TraceTheRay_GPU(r, num_tri, &is);

  // content of "is" is not shown correctly in Nsight.

  // display the normal as falsecolor (like normalmap)
   result = make_float3((is.normal.x+1.0)*0.5, (is.normal.y+1.0)*0.5, (is.normal.z+1.0)*0.5);

  int val = rgbToInt(result.x*255,result.y*255,result.z*255);
    out_data[y * w + x] = val;

    }

Редактировать: после ответа Робертса я добавил дополнительную функцию, чтобы отладчик действительно показывал значения.Чтобы избежать того, что какой-то оптимизатор может удалить его, это делает несколько бессмысленных операций, которые также добавляются к конечному результату.

__device__ int checkMyData(AbstractIntersection * is) {

    int ret_val = 0;

    if (is->normal.x > 0.2)
        ret_val = 17;
    else
        ret_val = 18;

    return ret_val; 

}

В Nsight я могу увидеть правильные значения «is» для первой строки checkMyData ().Я также могу сделать один setp (f10) и все еще видеть значения.Они исчезают в первом операторе if ().Я надеюсь, что это может помочь кому-то, кто плохо знаком с Cuda и Nsight.Таким образом, область видимости переменной в отладчике - это точка останова + следующие итерации, пока не произойдет первое попадание if.