[CUDA] should report error for inline __device__ variable

[aywala]

see CUDA - inline variable
A namespace scope inline variable declared with device or constant or managed memory space specifier must have internal linkage, if the code is compiled with nvcc in whole program compilation mode.

#include <stdio.h>

__device__ inline int a = 10;

__global__ void kernel() {
    printf("%d\n",a);
}

int main() {
    kernel<<<1,1>>>();
    cudaDeviceSynchronize();
}

nvcc report_error.cu -o report_error
report_error.cu(3): error: An inline __device__/__constant__/__managed__ variable must have internal linkage when the program is compiled in whole program mode (-rdc=false)
  __attribute__((device)) inline int a = 10;
                                     ^

1 error detected in the compilation of "report_error.cu".

but no error report with
clang++ report_error.cu -o report_error --cuda-gpu-arch=sm_52 -L /usr/local/cuda/lib64 -lcudart -pthread

[Artem-B]

I'm not quite sure what's the reason nvcc insists on making a special case for inline variables for the whole compilation mode.

AFAICT, in that mode it just does not matter whether it has internal linking or not, as there's only one TU that gets linked into the GPU executable.

I'm sure NVCC has their reasons, but I'm not convinced that those reasons apply to clang. I currently see no specific reason why compiler must treat inline vars differently in RDC and non-RDC compilations.

@yxsamliu -- what do you think?

[t-lutz]

This has to do with the way the CUDA runtime maps host addresses and device addresses. Each device variable has a corresponding host definition which is registered by the runtime and used for lookup. In whole program mode, each translation unit is compiled independently so an inline device variable behaves like a static one and will have multiple definitions (across multiple programs); the corresponding shadow definition would have linkonce linkage so a unique address across the whole host program; that would break the lookup (multiple device variables are registered against the same host variable). Demoting linkage of the shadow definition on the host is doable but breaks C++ semantics in several ways and would introduce odr-violations and other subtle bugs; so emitting an error is the best way to mitigate that.

Clang would have the same limitation as nvcc since it is using the same driver-compiler interface.

[Artem-B]

@t-lutz Thank you for the details.

I would argue that this situation fits better under "undefined behavior" category, not a hard error.

• as a user, I would expect "cc -c a.cc -o a.o" and cc a,cc -o a.exe to give me exactly the same compilation errors. Seeing a hard error is certainly an unexpected and inconsistent outcome.
• You have a valid point that the view of that __device__ inline variable will differ on the host and on the device, and, when we play by the normal C++ rules on the host, that will result in a potentially invalid behavior. Hard error here is sensible, but comes at the price of making compiler behavior inconsistent.

Demoting linkage of the shadow definition on the host is doable but breaks C++ semantics in several ways and would introduce odr-violations and other subtle bugs;
...
Clang would have the same limitation as nvcc

Agreed.

so emitting an error is the best way to mitigate that.

Some sort of diagnostics is clearly needed here. But...
I think it blocks benign uses, too. E.g. in a program that consists of a single TU or multiple TUs, but with only a single TU with a particular __device__ inline variable.

Considering that the compiler has no awareness about the larger build structure and can't tell whether __device__ inline will cause a problem, I think that a warning would be a better way to handle this situation.

[yxsamliu]

@t-lutz Thank you for the details.

I would argue that this situation fits better under "undefined behavior" category, not a hard error.

• as a user, I would expect "cc -c a.cc -o a.o" and cc a,cc -o a.exe to give me exactly the same compilation errors. Seeing a hard error is certainly an unexpected and inconsistent outcome.
• You have a valid point that the view of that __device__ inline variable will differ on the host and on the device, and, when we play by the normal C++ rules on the host, that will result in a potentially invalid behavior. Hard error here is sensible, but comes at the price of making compiler behavior inconsistent.

Demoting linkage of the shadow definition on the host is doable but breaks C++ semantics in several ways and would introduce odr-violations and other subtle bugs;
...
Clang would have the same limitation as nvcc

Agreed.

so emitting an error is the best way to mitigate that.

Some sort of diagnostics is clearly needed here. But... I think it blocks benign uses, too. E.g. in a program that consists of a single TU or multiple TUs, but with only a single TU with a particular __device__ inline variable.

Considering that the compiler has no awareness about the larger build structure and can't tell whether __device__ inline will cause a problem, I think that a warning would be a better way to handle this situation.

Agree. Since this may break users' working code, better make it a warning that can be disabled.

It is hard to guess users' intention when they add inline to a device variable. Is it to make it one instance across TU's, or just follow a habit to use inline variables in header.