I’ve run into some odd behaviour on the part of the NVIDIA compiler. It might be justified but I can’t quite spot the logic.
int i = ...;
if(threadIdx.x > i) { ... }
This is a stripped down version of a real kernel. The conditional test cannot take place directly because one parameter is unsigned (threadIdx.x) while the other is signed - and they’re possibly different sizes too.
A C compiler would perform a type conversion to make the test valid. The CUDA compiler appears to be reinterpreting the signed integer as an unsigned integer without performing any type conversion, or missing sign extension, or something. The compiler will even admit to this if you work directly with a constant:
if(threadIdx.x > -10) { ... }
warning: integer conversion resulted in a change of sign
The solution is to cast threadIdx.x to a signed integer by hand. Am I wrong to expect this behaviour from the compiler?
Edit: The use case for this is I have an integer parameter that may or may not be negative. The thread index test should always succeed for negative numbers but not always for positive numbers. Without manual casting the test fails for some negative numbers.
– The rank of any unsigned integer type shall equal the rank of the corresponding signed integer type, if any.
[ Thus ‘signed int’ and ‘unsigned int’ have the same rank which is used in section 6.3.1.8 to determine what direction promotions go. ]
6.3.1.8
[…]
Otherwise, if the operand that has unsigned integer type has rank greater or equal to the rank of the type of the other operand, then the operand with signed integer type is converted to the type of the operand with unsigned integer type.
[ Thus, the behavior specified by the standard is to convert the signed int to an unsigned int; hence the message from the compiler that you’re changing the sign of that value. ]
When you compare an unsigned int and a signed int , the signed int is converted to unsigned int. Converting a negative signed int to a unsigned int is done by adding
UINT_MAX + 1, so your negative signed number may become a big unsigned int.
No, it is just a way to explain it to people who do know nothing about twos complement I think, and using “real-world” numbers (i.e. not C ints) UINT_MAX+1 is not 0.
I did not check up, but I expect it does not matter anyway, converting a negative number to unsigned is almost certain to be either undefined or implementation-defined, and the same for converting a number > INT_MAX to signed, so if you make any assumptions about these cases you are not writing C but some compiler-specific code (and yes such code would not work e.g. on older Crays that use a different number encoding).
I think my confusion arose from expecting a warning from the compiler in a comparison of signed and unsigned numbers, as the GCC and Intel C/C++ compilers do.
NVCC does warn in the case of comparing a variable with a negative constant. Just not in the case where signed/unsigned variables are compared. That cost me a couple of hours of debugging.
I was going to start another rant about ivory tower trolls and the Standard.
But ok, fine. How about this simple fix:
Why is threadIdx unsigned in the first place? Is some future arch going to support 3 bn threads-per-block? Let’s make threadIdx signed, and all problems are solved. I think this won’t use extra PTX instructions.