CUDA kernel seemingly ignoring “if” statement

Question

What follows is the part of my kernel that does not behave properly, then an explanation of what I've found while debugging.

__global__ void Mangler(float *matrix, int *map)
{
    __shared__ signed int localMap[N];

    if(0 == threadIdx.x) 
    {
        for(int i=0; i<N; i++) 
            localMap[i] = -1;
    }

    __syncthreads();

    int fn = ...; // a lot of code goes into this number, skipped for clarity
    int rnumber = threadIdx.x;

    int X = atomicCAS(&localMap[fn], -1, rnumber); // Spot of bother 1

    if(X == -1) // Spot of bother 2
    {
        // some code
    }
    else 
    {
        // other code
    }
}

I've found in the documentation that atomicCAS(*address, compare, value) basically returns (and saves to the given address) the result of (old == compare ? value : old), where old is the value at the address before executing the function.

Going with this, I believe that executing int X = atomicCAS(&localMap[fn], -1, rnumber); should have two possible outcomes (according to NVidia Cuda C Programming Guide):

if localMap[fn] == -1 then X should have a value of rnumber and localMap[fn] should have a value of rnumber. This does not happen.
if localMap[fn] != -1 then X should be set to the value of localMap[fn] and said value should be left intact.

What happens instead, as debugging with NSight has shown me, is that X is being assigned -1, while localMap[fn] is being assigned the value of rnumber. I do not understand that, but as you can see in my code, I've changed the if to catch this situation.

Which brings me to spot of bother number 2: though NSight shows the value of X as -1, the if {} is being completely skipped (no breakpoints within hit whatsoever) and execution jumps straight to else.

My questions:

~~Do I misunderstand atomicCAS completely?~~ yes, I did
What could cause and if which should evaluate as true to jump straight into else in device code?

I'm using NVidia CUDA 5.5, Visual Studio 2012 x64 on Windows 8, NVidia Nsight Monitor Visual Studio Edition 3.1. The GPU for the machine is NVidia GeForce GTX 550 Ti.

I've tried changing the syntax to if(X!=-1); the true branch of the if is still not being executed.

While I appreciate the philosophical aspect of this question, could you suggest a way of confirming whether it does? I can't think of any, seeing as debugging device code is a pain to begin with. It generally didn't seem to tell any thinly veiled lies thus far. — Pawel J. Wal
Also, the value of X is later used as an array index (within else), and it results in a memory access violation, which is consistent with a -1 slipping in there. — Pawel J. Wal
One way to tell if the debugger is "telling the truth" is to insert some printf statements in your if and else paths. Are you compiling this code with the -G switch? Also, if X evaluates differently for different threads within a warp, things may be confusing depending on which thread you have in focus, or which version of X you are inspecting. — Robert Crovella
I seem to remember that if there is divergence in an if conditional, the else part is executed first. — Roger Dahl

flolo flolo · Accepted Answer · 2013-09-11T09:41:39

From the doc, atomicCAS returns the old value, that means, that in your list, your two outcomes are wrong. Your X will always be set to the old value of localMap[fn], regardless which value it had. What is set according to the comparison with the -1, is the new value of localMap[fn]. When it is -1, it is set to rnumber, else it is left intact.

So the behaviour you see with the values of X, rnumber and localMap are as expected.

I cannot help with your second problem, as I dont use NSight, and dont know how it works - according to your code, your true branch should be evaluated (but be careful: your false branch also - as it is multi threaded some threads can have the condition evaluated to true, and some to false - my guess/assumption would be that you must tell somehow your debugger which thread/warp/block you want to debug and you looked at the false).

CUDA kernel seemingly ignoring “if” statement

1 Answers