[CUDA] Clang does not diagnose invalid call to a host function from device.

[Artem-B]

Clang fails to check whether particular call is allowed and lets the calls to host functions sneak through into the device code.

https://godbolt.org/z/fqnhefGq3

__host__ int host_func();

struct A {
    int x; 
  __host__ A(int) { x = host_func(); }; // Must never be called on the device.
  __host__ __device__ A(int, int) { x = host_func(); }; // host_func() should trigger deferred diags
  __host__ __device__ A(int a, int b, int c) { x = a+b+c; }; // OK to call
};
struct B : A {
  __host__ __device__ B()
      : A(1) {} // A(1) should trigger deferred diag.
};

struct B1 : A {
  __host__ __device__ B1()
      : A(1, 2) {}
};

struct B2 : A {
  __host__ __device__ B2()
      : A(1, 2, host_func()) {} // host_func() call should trigger deferred diag.
};

struct C {
  B b;
  B1 b1;
  B2 b2;
};

__attribute__((global)) void kernel() {
  C c;
}

This results in ptxas failing due to an unresolved reference to the host_func

ptxas fatal   : Unresolved extern function '_Z9host_funcv'
clang++: error: ptxas command failed with exit code 255 (use -v to see invocation)

[Artem-B]

@yxsamliu ^^^ Found an interesting corner case where we allow calling a host function from a GPU w/o diagnosing it.

[yxsamliu]

The deferred diag relies on iterating the AST tree to detect the functions that are to be emitted.

It is supposed to detect emission of host_func() through the chain kernel()-> C::C() -> B::B() -> A::A(int) -> host_func() but did not.

[Artem-B]

The deferred diag relies on iterating the AST tree to detect the functions that are to be emitted.

That's likely the problem. I vaguely recall that constructor calls are special and materialize only during codegen, so we may be skipping the check at that point and need to apply it at the point where we actually call the constructors.