[clang][cuda] clang++ does not support __managed__ variables

[cbuchner1]

Hi,

I am having difficulty declaring__managed__ variables with clang++. This storage type means the variable should be accessible from both the host and the device. I am trying this with clang-20 and CUDA toolkit 12.1 which is supposed to be supported.

The declaration (in global scope)

__managed__ int val;
leads to clang++-20 (also clang++-17) output the warning

test.cu:1:1: warning: 'managed' attribute ignored [-Wignored-attributes]

The variable val will not be accessible from device functions and global kernels and it leads to compilation errors.

Here is a short reproducer that should be saved in a .cu file and passed to clang++ (possibly also defining a CUDA include path and linkage to the CUDA runtime)

#include <cstdio>

__managed__ int val;

__device__ void test_devicefunc() {
    printf("devicefunc val=%d\n", val);
}

__global__ void test_kernel() {
    printf("kernel val=%d\n", val);
    test_devicefunc();
}

int main(int argc, char *argv[]) {
    val = 100;
    test_kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

The same code compiles fine with nvcc and creates the output
kernel val=100
devicefunc val=100

[Artem-B]

Support for __managed__ variables has not been implemented in clang at all. NVIDIA does not document the interactions between those variables and the CUDA runtime. Combined with the fact that there is virtually no existing code that uses those managed variables, it's a very low priority feature.

If someone is motivated enough to figure it out and send a patch in, I'd be happy to review it.