Kernel behavior with variable initialization, index and printf()

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi.

I didn’t find any explanation in NVIDIA CUDA C Programming Guide B.16.

I run the kernel below <<< 1,1 >>> …

uint64_t idx = threadIdx.x;
uint64_t x = 0x0123456789abcdef;
uint64_t y = 0;
uint64_t z = idx;
if (y == z) {
    printf("x : %016llX \n", x);
    printf("y : %016llX \n", y);
    printf("z : %016llX \n", z);
}

// x : 0123456789ABCDEF
// y : 0000000000000000
// z : 0000000000000000

All as expected! Now …

uint64_t idx = threadIdx.x;
uint64_t x = 0x0123456789abcdef;
uint64_t y;
uint64_t z = idx;
if (y == z) {
    printf("x : %016llX \n", x);
    printf("y : %016llX \n", y);
    printf("z : %016llX \n", z);
}

// x : 0123456789ABCDEF
// y : 0123456789ABCDEF
// z : 0000000000000000

Why is omitting the initialization of y changing the printed result of y. Or the other way around, how can (y == z) be true if it isn’t true as shown by printf()?

Then this …

uint64_t x = 0x0123456789abcdef;
uint64_t y;
uint64_t z = 0;
printf("y : %016llX \n", y);
if (y == z) {
    printf("x : %016llX \n", x);
    printf("y : %016llX \n", y);
    printf("z : %016llX \n", z);
}

// y : 0000000000000000

Why is (y == z) not true? z is initialized to 0 and y shows 0 as well.

There are quite some printf() inside kernel posts out there, but I didn’t find any explaining above.
If I do all of the above in main.c, all works as expected.
What am I missing?

2

The value of expressions incorporating uninitialized variables is undefined. One cannot rely on undefined behavior matching any particular expectation, though it might occasionally, by chance.

3

Is it possible that uninitialized variable (UB) spits out value x in source code line 20, and the same variable value y in line 25?

4

Undefined is undefined. That means one cannot assume anything about the value of such an expression. One cannot reason about it in any way. It is futile to attempt to impose any kind of consistency requirement on undefined-ness.

What is the value of y at line N? Undefined. What is the value of y at line M? Undefined. What is the value of y == z? Undefined.

If it helps with a mental model of undefined-ness, maybe imagine an invisible hand tossing a 232-sided die each and every time there is a reference to y and assigning it the value that comes out on top.

1 Like
5

Understood. I always thought that undefined is related to the value as such, but that this value remains identical throughout the code. Thanks for the explain!