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

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.

Current clang behavior causes false alarms for valid use cases.

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

Differential Revision: https://reviews.llvm.org/D83893

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 283a046..e2190fc 100644 (file)
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -715,9 +715,8 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
                                       CallerKnownEmitted] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {
     case CFP_Never:
-      return DeviceDiagBuilder::K_Immediate;
     case CFP_WrongSide:
-      assert(Caller && "WrongSide calls require a non-null caller");
+      assert(Caller && "Never/wrongSide calls require a non-null caller");
       // If we know the caller will be emitted, we know this wrong-side call
       // will be emitted, so it's an immediate error.  Otherwise, defer the
       // error until we know the caller is emitted.

diff --git a/clang/test/SemaCUDA/builtins.cu b/clang/test/SemaCUDA/builtins.cu
index 814fda2..c01a687 100644 (file)
--- a/clang/test/SemaCUDA/builtins.cu
+++ b/clang/test/SemaCUDA/builtins.cu
@@ -7,10 +7,10 @@
 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -triple x86_64-unknown-unknown \
 // RUN:     -aux-triple nvptx64-unknown-cuda \
-// RUN:     -fsyntax-only -verify %s
+// RUN:     -fsyntax-only -verify=host %s
 // RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
 // RUN:     -aux-triple x86_64-unknown-unknown \
-// RUN:     -fsyntax-only -verify %s
+// RUN:     -fsyntax-only -verify=dev %s
 
 #if !(defined(__amd64__) && defined(__PTX__))
 #error "Expected to see preprocessor macros from both sides of compilation."
@@ -18,14 +18,14 @@
 
 void hf() {
   int x = __builtin_ia32_rdtsc();
-  int y = __nvvm_read_ptx_sreg_tid_x(); // expected-note  {{'__nvvm_read_ptx_sreg_tid_x' declared here}}
-  // expected-error@-1 {{reference to __device__ function '__nvvm_read_ptx_sreg_tid_x' in __host__ function}}
+  int y = __nvvm_read_ptx_sreg_tid_x(); // host-note  {{'__nvvm_read_ptx_sreg_tid_x' declared here}}
+  // host-error@-1 {{reference to __device__ function '__nvvm_read_ptx_sreg_tid_x' in __host__ function}}
   x = __builtin_abs(1);
 }
 
 __attribute__((device)) void df() {
   int x = __nvvm_read_ptx_sreg_tid_x();
-  int y = __builtin_ia32_rdtsc(); // expected-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
-                                  // expected-note@20 {{'__builtin_ia32_rdtsc' declared here}}
+  int y = __builtin_ia32_rdtsc(); // dev-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
+                                  // dev-note@20 {{'__builtin_ia32_rdtsc' declared here}}
   x = __builtin_abs(1);
 }

diff --git a/clang/test/SemaCUDA/call-kernel-from-kernel.cu b/clang/test/SemaCUDA/call-kernel-from-kernel.cu
index c89037c..900efce 100644 (file)
--- a/clang/test/SemaCUDA/call-kernel-from-kernel.cu
+++ b/clang/test/SemaCUDA/call-kernel-from-kernel.cu
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
-// RUN:   -verify -fsyntax-only -verify-ignore-unexpected=note
+// RUN: %clang_cc1 %s --std=c++11 -triple nvptx -emit-llvm -o - \
+// RUN:   -verify -fcuda-is-device -fsyntax-only -verify-ignore-unexpected=note
 
 #include "Inputs/cuda.h"
 

diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu
index b9efd1c..191268c 100644 (file)
--- a/clang/test/SemaCUDA/function-overload.cu
+++ b/clang/test/SemaCUDA/function-overload.cu
@@ -1,8 +1,8 @@
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify=host,expected %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify=dev,expected %s
 
 #include "Inputs/cuda.h"
 
@@ -75,37 +75,37 @@ extern "C" __host__ __device__ int chhd2() { return 0; }
 
 // Helper functions to verify calling restrictions.
 __device__ DeviceReturnTy d() { return DeviceReturnTy(); }
-// expected-note@-1 1+ {{'d' declared here}}
+// host-note@-1 1+ {{'d' declared here}}
 // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
 
 __host__ HostReturnTy h() { return HostReturnTy(); }
-// expected-note@-1 1+ {{'h' declared here}}
+// dev-note@-1 1+ {{'h' declared here}}
 // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
 // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
 
 __global__ void g() {}
-// expected-note@-1 1+ {{'g' declared here}}
+// dev-note@-1 1+ {{'g' declared here}}
 // expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
 // expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
 // expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
 
 extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
-// expected-note@-1 1+ {{'cd' declared here}}
+// host-note@-1 1+ {{'cd' declared here}}
 // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
 
 extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
-// expected-note@-1 1+ {{'ch' declared here}}
+// dev-note@-1 1+ {{'ch' declared here}}
 // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
 // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
 // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
 
 __host__ void hostf() {
-  DeviceFnPtr fp_d = d;         // expected-error {{reference to __device__ function 'd' in __host__ function}}
+  DeviceFnPtr fp_d = d;         // host-error {{reference to __device__ function 'd' in __host__ function}}
   DeviceReturnTy ret_d = d();   // expected-error {{no matching function for call to 'd'}}
-  DeviceFnPtr fp_cd = cd;       // expected-error {{reference to __device__ function 'cd' in __host__ function}}
+  DeviceFnPtr fp_cd = cd;       // host-error {{reference to __device__ function 'cd' in __host__ function}}
   DeviceReturnTy ret_cd = cd(); // expected-error {{no matching function for call to 'cd'}}
 
   HostFnPtr fp_h = h;
@@ -129,9 +129,9 @@ __device__ void devicef() {
   DeviceFnPtr fp_cd = cd;
   DeviceReturnTy ret_cd = cd();
 
-  HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __device__ function}}
+  HostFnPtr fp_h = h;         // dev-error {{reference to __host__ function 'h' in __device__ function}}
   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
-  HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __device__ function}}
+  HostFnPtr fp_ch = ch;       // dev-error {{reference to __host__ function 'ch' in __device__ function}}
   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
 
   DeviceFnPtr fp_dh = dh;
@@ -139,9 +139,9 @@ __device__ void devicef() {
   DeviceFnPtr fp_cdh = cdh;
   DeviceReturnTy ret_cdh = cdh();
 
-  GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
+  GlobalFnPtr fp_g = g; // dev-error {{reference to __global__ function 'g' in __device__ function}}
   g(); // expected-error {{no matching function for call to 'g'}}
-  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
+  g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __device__ function}}
 }
 
 __global__ void globalf() {
@@ -150,9 +150,9 @@ __global__ void globalf() {
   DeviceFnPtr fp_cd = cd;
   DeviceReturnTy ret_cd = cd();
 
-  HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __global__ function}}
+  HostFnPtr fp_h = h;         // dev-error {{reference to __host__ function 'h' in __global__ function}}
   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
-  HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __global__ function}}
+  HostFnPtr fp_ch = ch;       // dev-error {{reference to __host__ function 'ch' in __global__ function}}
   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
 
   DeviceFnPtr fp_dh = dh;
@@ -160,9 +160,9 @@ __global__ void globalf() {
   DeviceFnPtr fp_cdh = cdh;
   DeviceReturnTy ret_cdh = cdh();
 
-  GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
+  GlobalFnPtr fp_g = g; // dev-error {{reference to __global__ function 'g' in __global__ function}}
   g(); // expected-error {{no matching function for call to 'g'}}
-  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
+  g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __global__ function}}
 }
 
 __host__ __device__ void hostdevicef() {

diff --git a/clang/test/SemaCUDA/function-target.cu b/clang/test/SemaCUDA/function-target.cu
index 83dce50..48f7229 100644 (file)
--- a/clang/test/SemaCUDA/function-target.cu
+++ b/clang/test/SemaCUDA/function-target.cu
@@ -1,5 +1,5 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify=dev,expected %s
 
 #include "Inputs/cuda.h"
 
@@ -23,11 +23,11 @@ __host__ void h1(void) {
 __host__ void d1h(void); // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
 __device__ void d1d(void);
 __host__ __device__ void d1hd(void);
-__global__ void d1g(void); // expected-note {{'d1g' declared here}}
+__global__ void d1g(void); // dev-note {{'d1g' declared here}}
 
 __device__ void d1(void) {
   d1h(); // expected-error {{no matching function}}
   d1d();
   d1hd();
-  d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
+  d1g<<<1, 1>>>(); // dev-error {{reference to __global__ function 'd1g' in __device__ function}}
 }

diff --git a/clang/test/SemaCUDA/implicit-device-lambda.cu b/clang/test/SemaCUDA/implicit-device-lambda.cu
index 8e5b7dd..d2e59b8 100644 (file)
--- a/clang/test/SemaCUDA/implicit-device-lambda.cu
+++ b/clang/test/SemaCUDA/implicit-device-lambda.cu
@@ -1,5 +1,7 @@
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -fsyntax-only -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
-// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify=dev,expected -fsyntax-only \
+// RUN:   -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only \
+// RUN:   -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
 
 #include "Inputs/cuda.h"
 
@@ -102,5 +104,5 @@ __device__ void foo() {
     void foo() {}
   };
   X x;
-  x.foo(); // expected-error {{reference to __host__ function 'foo' in __device__ function}}
+  x.foo(); // dev-error {{reference to __host__ function 'foo' in __device__ function}}
 }

diff --git a/clang/test/SemaCUDA/method-target.cu b/clang/test/SemaCUDA/method-target.cu
index 8e17daa..85c27ce 100644 (file)
--- a/clang/test/SemaCUDA/method-target.cu
+++ b/clang/test/SemaCUDA/method-target.cu
@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -verify=host,expected %s
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,expected %s
 
 #include "Inputs/cuda.h"
 
@@ -6,11 +7,11 @@
 // Test 1: host method called from device function
 
 struct S1 {
-  void method() {} // expected-note {{'method' declared here}}
+  void method() {} // dev-note {{'method' declared here}}
 };
 
 __device__ void foo1(S1& s) {
-  s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+  s.method(); // dev-error {{reference to __host__ function 'method' in __device__ function}}
 }
 
 //------------------------------------------------------------------------------
@@ -29,22 +30,22 @@ __device__ void foo2(S2& s, int i, float f) {
 // Test 3: device method called from host function
 
 struct S3 {
-  __device__ void method() {} // expected-note {{'method' declared here}}
+  __device__ void method() {} // host-note {{'method' declared here}}
 };
 
 void foo3(S3& s) {
-  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}}
+  s.method(); // host-error {{reference to __device__ function 'method' in __host__ function}}
 }
 
 //------------------------------------------------------------------------------
 // Test 4: device method called from host&device function
 
 struct S4 {
-  __device__ void method() {}  // expected-note {{'method' declared here}}
+  __device__ void method() {}  // host-note {{'method' declared here}}
 };
 
 __host__ __device__ void foo4(S4& s) {
-  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+  s.method(); // host-error {{reference to __device__ function 'method' in __host__ __device__ function}}
 }
 
 //------------------------------------------------------------------------------
@@ -63,9 +64,9 @@ __device__ void foo5(S5& s, S5& t) {
 // Test 6: call method through pointer
 
 struct S6 {
-  void method() {} // expected-note {{'method' declared here}};
+  void method() {} // dev-note {{'method' declared here}};
 };
 
 __device__ void foo6(S6* s) {
-  s->method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+  s->method(); // dev-error {{reference to __host__ function 'method' in __device__ function}}
 }

diff --git a/clang/test/SemaCUDA/reference-to-kernel-fn.cu b/clang/test/SemaCUDA/reference-to-kernel-fn.cu
index e502d13..70a1cda 100644 (file)
--- a/clang/test/SemaCUDA/reference-to-kernel-fn.cu
+++ b/clang/test/SemaCUDA/reference-to-kernel-fn.cu
@@ -1,12 +1,14 @@
-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify \
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify=host \
+// RUN:   -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify=dev \
 // RUN:   -verify-ignore-unexpected=note %s
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify \
-// RUN:   -verify-ignore-unexpected=note -DDEVICE %s
 
 // Check that we can reference (get a function pointer to) a __global__
 // function from the host side, but not the device side.  (We don't yet support
 // device-side kernel launches.)
 
+// host-no-diagnostics
+
 #include "Inputs/cuda.h"
 
 struct Dummy {};
@@ -17,13 +19,11 @@ typedef void (*fn_ptr_t)();
 
 __host__ __device__ fn_ptr_t get_ptr_hd() {
   return kernel;
-#ifdef DEVICE
-  // expected-error@-2 {{reference to __global__ function}}
-#endif
+  // dev-error@-1 {{reference to __global__ function}}
 }
 __host__ fn_ptr_t get_ptr_h() {
   return kernel;
 }
 __device__ fn_ptr_t get_ptr_d() {
-  return kernel;  // expected-error {{reference to __global__ function}}
+  return kernel;  // dev-error {{reference to __global__ function}}
 }