diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 283a04683a32..6203edea7112 100644 --- 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. @@ -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; } diff --git a/clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp b/clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp index 9351b4ecb032..88fcbd716ef4 100644 --- a/clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp +++ b/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, ""); } diff --git a/clang/test/SemaCUDA/builtins.cu b/clang/test/SemaCUDA/builtins.cu index 814fda2ac7d3..78a333e511a5 100644 --- 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,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); } diff --git a/clang/test/SemaCUDA/call-kernel-from-kernel.cu b/clang/test/SemaCUDA/call-kernel-from-kernel.cu index c89037c52bff..900efcef43b8 100644 --- 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 b9efd1c09e69..191268c9a5f1 100644 --- 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 83dce50b4af8..48f7229df21f 100644 --- 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 8e5b7ddddb8f..d2e59b8033c3 100644 --- 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 8e17daa0c123..85c27ce43632 100644 --- 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 e502d134b086..70a1cda6ab0c 100644 --- 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}} }