I am getting a very strange error that I hope someone can explain to me.
I have code similar to this:
__device__ foo()
{
bool isTrue = false;
int i = 0;
do
{
if(i%2)
isTrue = true;
i++;
}while(i < 10);
if(isTrue)
doSomething();
}
I realize that in this example is silly because ‘isTrue’ will always take ‘true’, but it illustrates the problem nicely…
The problem is that if I try to set ‘isTrue’ inside an if in the do-while loop and then access it outside of the loop my kernel execution time gets reduced by a full order of magnitude, even to the point of causing the kernel to not run. If I comment out the ‘if(isTrue)’ part and change nothing else then the kernel runs at full speed. If I set ‘isTrue’ to ‘true’ after the loop is finished executing then it also runs at full speed. I don’t understand why accessing the variable should cause such a slowdown compared with all of the other calculation I am doing. Also, in my code ‘doSomething()’ simply changes the value of another variable so it is not like I am getting a major divergence or anything.
Also, my code has nested ‘if’ statements in the do-while loop and if I try to set ‘isTrue’ in a deeper ‘if’ statement the problem gets even worse.
I am at a total loss right now because my code runs really fast… until I try to read anything that I may have set in the loop, and it only happens if I read the variable, not if I write it first.
Thanks.