[PATCH] D83893: [CUDA][HIP] Always defer diagnostics for wrong-sided reference

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added a subscriber: kristof.beyls.

When a device function calls a host function or vice versa, this is wrong-sided
reference. Currently clang immediately diagnose it. This is different from nvcc
behavior, where it is diagnosed only if the function is really emitted:

  void foo();
  inline __device__ void bar() { foo(); }

https://godbolt.org/z/4G8P31

To me nvcc behavior makes more sense, since we do not need a diagnostic
unless we really need it. Current clang behavior causes diagnostics for wrong-sided
reference emitted twice, once in host compilation, once in device compilation,
which is unnecessary and cluttering the screen.

More importantly, current clang behavior causes false alarms for valid use cases:

  __device__ void bar(int x);
  
  template<class T>
  __device__ void foo(T x) {
    bar(x);
  }
  
  struct A {
    int a;
    operator int() { return a; }
  };
  
  void foo(A a) {}
  
  template<typename T>
  __device__ __host__ void test(T t) {
    foo(t);
  }
  
  int main() {
    A a;
    test(a);
    return 0;
  } 

In this example, we only emit `test<A>` on host, it should compile.
However since `test<A>` is a host device function, it triggers
instantiation of `foo<A>` in device compilation. `foo<A>` calls
a host `operator int()` which causes immediate diagnostics.

In this case, user may never intend to use `foo<A>` on device side.
The instantiation of `foo<A>` is a side effect of instantiation of
host device function `test<A>`, therefore any diagnostics incurred
by `foo<A>` should be deferred.

This patch let clang always defer diagnostics for wrong-sided
reference.


https://reviews.llvm.org/D83893

Files:
  clang/lib/Sema/SemaCUDA.cpp
  clang/test/SemaCUDA/builtins.cu
  clang/test/SemaCUDA/call-kernel-from-kernel.cu
  clang/test/SemaCUDA/function-overload.cu
  clang/test/SemaCUDA/function-target.cu
  clang/test/SemaCUDA/implicit-device-lambda.cu
  clang/test/SemaCUDA/method-target.cu
  clang/test/SemaCUDA/reference-to-kernel-fn.cu

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D83893.278241.patch
Type: text/x-patch
Size: 15461 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200715/b4a79a37/attachment-0001.bin>