[clang] 4fc752b - [CUDA][HIP] Always defer diagnostics for wrong-sided reference
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 17 04:52:29 PDT 2020
Author: Yaxun (Sam) Liu
Date: 2020-07-17T07:51:55-04:00
New Revision: 4fc752b30b9acac73a282cb844a6240e6cb70cca
URL: https://github.com/llvm/llvm-project/commit/4fc752b30b9acac73a282cb844a6240e6cb70cca
DIFF: https://github.com/llvm/llvm-project/commit/4fc752b30b9acac73a282cb844a6240e6cb70cca.diff
LOG: [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
Added:
Modified:
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
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 283a04683a32..e2190fc42de4 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.
diff --git a/clang/test/SemaCUDA/builtins.cu b/clang/test/SemaCUDA/builtins.cu
index 814fda2ac7d3..c01a687e12c0 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,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 at -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 at -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 at 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 at 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 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 at -1 1+ {{'d' declared here}}
+// host-note at -1 1+ {{'d' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
__host__ HostReturnTy h() { return HostReturnTy(); }
-// expected-note at -1 1+ {{'h' declared here}}
+// dev-note at -1 1+ {{'h' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
// expected-note at -4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
__global__ void g() {}
-// expected-note at -1 1+ {{'g' declared here}}
+// dev-note at -1 1+ {{'g' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
// expected-note at -4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
-// expected-note at -1 1+ {{'cd' declared here}}
+// host-note at -1 1+ {{'cd' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
-// expected-note at -1 1+ {{'ch' declared here}}
+// dev-note at -1 1+ {{'ch' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
// expected-note at -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 at -2 {{reference to __global__ function}}
-#endif
+ // dev-error at -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}}
}
More information about the cfe-commits
mailing list