[CUDA][HIP] shared variable not allowed with trivial ctor

[aywala]

Similar with https://github.com/llvm/llvm-project/issues/72261

struct A {
  A(){}
  ~A() {}
  int x;
};

#define __device__ __attribute__((device))
#define __shared__ __attribute__((shared))

__device__ A x;

__device__ void foo() {
  __shared__ A y;  // error with clang, ok with nvcc
}

https://github.com/llvm/llvm-project/pull/73140 fixed global device var init, but non-global shared var still fails.

nvcc does not check host device attribute for either global device var or shared var.

nvcc even does not report error when constructor for shared var is not empty ( just report warning, but the constructor will never be called).

@yxsamliu

https://godbolt.org/z/5T9fKdavn https://godbolt.org/z/GqW6PTG18

[Artem-B]

nvcc even does not report error when constructor for shared var is not empty ( just report warning, but the constructor will never be called).

That's arguably a bug in NVCC. This results in obviously broken code. Presumably the stuff non-empty constructor is intended to do is important, but it will never be called because we have no dynamic initializers on the GPU. This should've been a hard error, not a warning.

[Artem-B]

@aywala It's not a bug. It's a portability issue in your example.

<source>:14:16: error: no matching constructor for initialization of 'A'
   14 |   __shared__ A y;
      |                ^
<source>:2:3: note: candidate constructor not viable: call to __host__ function from __device__ function
    2 |   A(){}
      |   ^
<source>:1:8: note: candidate constructor (the implicit copy constructor) not viable: requires 1 argument, but 0 were provided
    1 | struct A {
      |        ^

The problem is not that trivial constructor is not allowed, but rather than the only defined constructors are __host__ and can't be accessed on the GPU side at all.

This is the area where NVCC and clang, by design, differ in their behavior. NVCC physically removes host-side code from the sources, so compilation succeeds because struct A ends up calling the default constructor.

Clang, on the other hand, sees the code as written, and, as far as it's concerned, there's only a __host__ constructor, and, it can not be called from the GPU code.

Making the constructors/destructors __host__ __device__ will make the code work with both nvcc and clang. https://godbolt.org/z/4c8j6o6x6

Alternatively, you could define separate __device__ overloads for the constructor/destructor. It will get the code to compile with clang, but NVCC does not support attribute-based overloads, so that change would not be portable back to nvcc.

[Artem-B]

Closing as "not a bug".

[aywala]

@Artem-B It is totally the same thing as https://github.com/llvm/llvm-project/issues/72261. Yeah, it can be considered not a bug but a different behavior from nvcc. I noticed some work has been done (https://github.com/llvm/llvm-project/pull/73140) to match nvcc, so I think this code sample should be allowed too.

[aywala]

@Artem-B

That's arguably a bug in NVCC. This results in obviously broken code. Presumably the stuff non-empty constructor is intended to do is important, but it will never be called because we have no dynamic initializers on the GPU. This should've been a hard error, not a warning.

I totally agree it should be an error.

This is the area where NVCC and clang, by design, differ in their behavior. NVCC physically removes host-side code from the sources, so compilation succeeds because struct A ends up calling the default constructor.

That may be not right. When there is a non-empty constructor without __device__ attr for global device-side variables or shared variables, nvcc still reports error/warning about dynamic initialization, so it may not been removed.

https://godbolt.org/z/cGWbefzs3