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
Index: clang/test/SemaCUDA/reference-to-kernel-fn.cu
===================================================================
--- clang/test/SemaCUDA/reference-to-kernel-fn.cu
+++ 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 @@
__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}}
}
Index: clang/test/SemaCUDA/method-target.cu
===================================================================
--- clang/test/SemaCUDA/method-target.cu
+++ 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 @@
// 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 @@
// 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}}
}
Index: clang/test/SemaCUDA/implicit-device-lambda.cu
===================================================================
--- clang/test/SemaCUDA/implicit-device-lambda.cu
+++ 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 @@
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}}
}
Index: clang/test/SemaCUDA/function-target.cu
===================================================================
--- clang/test/SemaCUDA/function-target.cu
+++ 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 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}}
}
Index: clang/test/SemaCUDA/function-overload.cu
===================================================================
--- clang/test/SemaCUDA/function-overload.cu
+++ 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 @@
// 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 @@
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 @@
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 @@
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 @@
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() {
Index: clang/test/SemaCUDA/call-kernel-from-kernel.cu
===================================================================
--- clang/test/SemaCUDA/call-kernel-from-kernel.cu
+++ 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"
Index: clang/test/SemaCUDA/builtins.cu
===================================================================
--- clang/test/SemaCUDA/builtins.cu
+++ 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);
}
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -715,9 +715,8 @@
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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
