integer division with Constant

Question

I'm writing a CUDA kernel and I have been running into some peculiar behavior with integer division and constants that's leaving me scratching my head.
I'll illustrate what happens below:

#define X 8
#define Y 4
#define K X/Y
...code....
int var = 8;
...code....
printf("K = %d, var = %d, var/K = %d\n", K, var, var/K);

I am using var/K as a terminating condition in a loop somewhere, and I would suspect that var/K would give a value of 4, but It gives a value of 1 instead, thus breaking my kernel. Here is the output of the print statement:

K = 2, var = 8, var/K = 1

When I change K from a #define to a global constant of type int:

__device__ const int K=X/Y;

the problem goes away, the division result gives 4, and my kernel functions correctly. I've tried casting K to int (which shouldn't be necessary) and that doesn't change anything.
I'm wondering if there is some restriction on dividing constants or if I have a basic misunderstanding somewhere.

preprocessor is just text replacement, so your definitions give `8/8/4` for `var/K` but why does that result in `1`? — deamentiaemundi, Feb 18 '17 at 05:13
This really has nothing to do with CUDA, just preprocessor quirk... — CygnusX1, Feb 18 '17 at 07:08
Another example of [why macros are dangerous, mmmkay](http://stackoverflow.com/questions/14041453/why-are-preprocessor-macros-evil-and-what-are-the-alternatives)? — einpoklum, Feb 18 '17 at 14:41

Robert Crovella · Accepted Answer · 2017-02-18T14:28:56.513

The basic problem here has nothing to do with CUDA; it would manifest itself with ordinary C/C++ compilers.

As indicated in the comments, creating "constants" in these two different ways:

#define K X/Y

and

const int K=X/Y;

does not yield the same thing in all cases, as you've discovered. The first methodology uses the C or C++ preprocessor to go through your code and make a literal (text) substitution of X/Y any place that it finds K. So when you do this:

int var = 8;
printf("K = %d, var = %d, var/K = %d\n", K, var, var/K);

the expression var/K quite literally gets replaced with var/X/Y, before the language processing begins. And since C and C++ would evaluate such an expression from left to right, it does not yield the result you expected.

The second methodology, of course, creates a real variable in the language, and works as you expect. Again, none of this is unique to CUDA.

Although it may not remove every kind of quirky behavior, a common recommendation when using preprocessor macros (i.e. #define) that involve expressions, is to always use parenthesis around such macros, like so:

#define K (X/Y)

This generally forces the text expression in your macro to be evaluated via the language processor in a way that is consistent with your intent. In particular, it prevents the unwanted side effect you have here, where some parts of your macro are participating in adjacent operations before the expression in the macro itself gets evaluated. And of course, simply defining a global const variable works as you've already discovered, and may be an even better choice.

As an aside, and again referred to in the comments, I would expect var/K as you've shown here to be evaluated as 0 not 1 for the case you have given. However you've not shown a complete code, so there may be something else at work here, and regardless of that difference of opinion, your usage of the macro as you have shown would not give your expected result of 8/2 = 4. So whether it should be 0, as I claim, or 1, as you claim, it would not give your expected result of 4.

integer division with Constant

1 Answers1