0

This is a kernel to calculate the normal vector.

__global__ void getNormalMapKernel(const ushort *dev_depth, float3 *normalMap, const float *K ,const float *T)
{
    int x = threadIdx.x;                             
    int y = blockIdx.x;                                   

    float3 normal = make_float3(0.0, 0.0, 0.0);

    if (x < IMG_WIDTH - 1 && y < IMG_HEIGHT - 1)
    {
        ushort depth = dev_depth[y * IMG_WIDTH + x];
        ushort depth_right = dev_depth[y * IMG_WIDTH + x + 1];
        ushort depth_down = dev_depth[(y + 1) * IMG_WIDTH + x];

        if ( depth && absDevUshort(depth, depth_right) < 20 && absDevUshort(depth, depth_down) < 20)
        {
            float3 cameraPos = make_float3((x - K[1]) * depth / K[0] - T[0], (y - K[3]) * depth / K[2] - T[1], depth - T[2]);
            float3 cameraPosRight = make_float3((x + 1 - K[1]) * depth_right / K[0] - T[0], (y - K[3]) * depth_right / K[2] - T[1], depth_right - T[2]);
            float3 cameraPosDown = make_float3((x - K[1]) * depth_down / K[0] - T[0], (y + 1 - K[3]) * depth_down / K[2] - T[1], depth_down - T[2]);

            normal = normalize(cross(cameraPosRight - cameraPos, cameraPosDown - cameraPos));
        }
    }

    normalMap[y * IMG_WIDTH + x] = normal;
}

When executing the kernel function, the computer screen begin to blink and then blue and then crash.

If I comment the two ifs, everything is right.

Then I try to move the last assignment into the if like:

__global__ void getNormalMapKernel(const ushort *dev_depth, float3 *normalMap, const float *K ,const float *T)
{
    int x = threadIdx.x;                             
    int y = blockIdx.x;                                   

    float3 normal = make_float3(0.0, 0.0, 0.0);

    if (x < IMG_WIDTH - 1 && y < IMG_HEIGHT - 1)
    {
        ushort depth = dev_depth[y * IMG_WIDTH + x];
        ushort depth_right = dev_depth[y * IMG_WIDTH+ x + 1];
        ushort depth_down = dev_depth[(y + 1) * IMG_WIDTH + x];

        if ( depth && absDevUshort(depth, depth_right) < 20 && absDevUshort(depth, depth_down) < 20)
        {
            float3 cameraPos = make_float3((x - K[1]) * depth / K[0] - T[0], (y - K[3]) * depth / K[2] - T[1], depth - T[2]);
            float3 cameraPosRight = make_float3((x + 1 - K[1]) * depth_right / K[0] - T[0], (y - K[3]) * depth_right / K[2] - T[1], depth_right - T[2]);
            float3 cameraPosDown = make_float3((x - K[1]) * depth_down / K[0] - T[0], (y + 1 - K[3]) * depth_down / K[2] - T[1], depth_down - T[2]);

            normalMap[y * IMG_WIDTH + x]= normalize(cross(cameraPosRight - cameraPos, cameraPosDown - cameraPos));
        }
    }
}

It crashes again. It seems that if I change the value of normal in the ifs or assign a new value to the normalMap[y * IMG_WIDTH + x] in the ifs, it will cause crash.

Anyone can help? I will be deeply grateful to who can solve the problem.

Kill Console
  • 1,933
  • 18
  • 15
  • Do you mean, the computer crashes? If yes, then you probably have a driver issue, try updating them. If the screen simply freezes, goes to black, but comes back eventually, then your kernel is likely taking too long to execute and Windows is killing the driver. Lookup Timeout Detection and Recovery (TDR). – user703016 Oct 15 '14 at 11:31
  • @Cicada the second condition. In the first code block above, if I do not change the value of normal, like I assign the normalized result to another local variable, the kernel takes the same amount of calculation, everything is OK. I just don't know why. – Kill Console Oct 15 '14 at 11:49
  • What OS? If you want help with a code that is not working, you should provide a MCVE. – Robert Crovella Oct 15 '14 at 12:35
  • @RobertCrovella Windows 7. This kernel is just one function in my project. Like I said above, if I comment the normalize operation or assign the normalized result to another local variable instead of normal, the program won't crash. – Kill Console Oct 15 '14 at 12:45
  • Have you looked at the thread that I've linked? Did it help? – Michal Hosala Oct 20 '14 at 12:45
  • @MichalHosala If it helps, I won't start the bounty. I've seen that question before I asked this one. – Kill Console Oct 21 '14 at 11:27
  • 1
    Then it would be nice to sum up your investigation up until now in the description of this question, so the potential askers don't have to investigate paths that you already have. – Michal Hosala Oct 21 '14 at 11:32
  • 2
    Are you 100% sure it crashes in *getNormalMapKernel* (and not in a latter kernel) ? Maybe the two *if* prevent *normalMap* to be filled for some *x* and *y* so in a latter kernel, you get an UMR which causes the crash. Have you tried a *cuda-memcheck* on your program ? – GaTTaCa Oct 22 '14 at 08:23

0 Answers0