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

Fixed regression in test builtin-amdgcn-atomic-inc-dec-failure.cpp.

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.
@@ -740,9 +739,10 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
 
   DeviceDiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
       << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
-  DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
-                    Caller, *this)
-      << Callee;
+  if (!Callee->getBuiltinID())
+    DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
+                      Caller, *this)
+        << Callee;
   return DiagKind != DeviceDiagBuilder::K_Immediate &&
          DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
 }

clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp

@@ -1,19 +1,26 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: not %clang_cc1 %s -x hip -fcuda-is-device -o - -emit-llvm -triple=amdgcn-amd-amdhsa 2>&1 | FileCheck %s
+// RUN: %clang_cc1 %s -x hip -fcuda-is-device -o - \
+// RUN:   -triple=amdgcn-amd-amdhsa -fsyntax-only \
+// RUN:   -verify=dev
+// RUN: %clang_cc1 %s -x hip -triple x86_64 -o - \
+// RUN:   -aux-triple amdgcn-amd-amdhsa -fsyntax-only \
+// RUN:   -verify=host
+
+// dev-no-diagnostics
 
 void test_host() {
   __UINT32_TYPE__ val32;
   __UINT64_TYPE__ val64;
 
-  // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_inc32' in __host__ function
+  // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_inc32' in __host__ function}}
   val32 = __builtin_amdgcn_atomic_inc32(&val32, val32, __ATOMIC_SEQ_CST, "");
 
-  // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_inc64' in __host__ function
+  // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_inc64' in __host__ function}}
   val64 = __builtin_amdgcn_atomic_inc64(&val64, val64, __ATOMIC_SEQ_CST, "");
 
-  // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_dec32' in __host__ function
+  // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_dec32' in __host__ function}}
   val32 = __builtin_amdgcn_atomic_dec32(&val32, val32, __ATOMIC_SEQ_CST, "");
 
-  // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_dec64' in __host__ function
+  // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_dec64' in __host__ function}}
   val64 = __builtin_amdgcn_atomic_dec64(&val64, val64, __ATOMIC_SEQ_CST, "");
 }

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,13 @@
 
 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-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}}
   x = __builtin_abs(1);
 }

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"
 

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() {

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}}
 }

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}}
 }

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}}
 }

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}}
 }