[clang] a46ef7d - Revert "[CUDA][HIP] Always defer diagnostics for wrong-sided reference"
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 17 05:11:29 PDT 2020
Author: Yaxun (Sam) Liu
Date: 2020-07-17T08:10:56-04:00
New Revision: a46ef7d42dc8aa5083319bef678262bddf299f82
URL: https://github.com/llvm/llvm-project/commit/a46ef7d42dc8aa5083319bef678262bddf299f82
DIFF: https://github.com/llvm/llvm-project/commit/a46ef7d42dc8aa5083319bef678262bddf299f82.diff
LOG: Revert "[CUDA][HIP] Always defer diagnostics for wrong-sided reference"
This reverts commit 4fc752b30b9acac73a282cb844a6240e6cb70cca.
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 e2190fc42de4..283a04683a32 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -715,8 +715,9 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
CallerKnownEmitted] {
switch (IdentifyCUDAPreference(Caller, Callee)) {
case CFP_Never:
+ return DeviceDiagBuilder::K_Immediate;
case CFP_WrongSide:
- assert(Caller && "Never/wrongSide calls require a non-null caller");
+ assert(Caller && "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 c01a687e12c0..814fda2ac7d3 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=host %s
+// RUN: -fsyntax-only -verify %s
// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
// RUN: -aux-triple x86_64-unknown-unknown \
-// RUN: -fsyntax-only -verify=dev %s
+// RUN: -fsyntax-only -verify %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(); // 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}}
+ 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}}
x = __builtin_abs(1);
}
__attribute__((device)) void df() {
int x = __nvvm_read_ptx_sreg_tid_x();
- 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}}
+ 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}}
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 900efcef43b8..c89037c52bff 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 nvptx -emit-llvm -o - \
-// RUN: -verify -fcuda-is-device -fsyntax-only -verify-ignore-unexpected=note
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN: -verify -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 191268c9a5f1..b9efd1c09e69 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=host,expected %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify=dev,expected %s
+// 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
#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(); }
-// host-note at -1 1+ {{'d' declared here}}
+// expected-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(); }
-// dev-note at -1 1+ {{'h' declared here}}
+// expected-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() {}
-// dev-note at -1 1+ {{'g' declared here}}
+// expected-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(); }
-// host-note at -1 1+ {{'cd' declared here}}
+// expected-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(); }
-// dev-note at -1 1+ {{'ch' declared here}}
+// expected-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; // host-error {{reference to __device__ function 'd' in __host__ function}}
+ DeviceFnPtr fp_d = d; // expected-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; // host-error {{reference to __device__ function 'cd' in __host__ function}}
+ DeviceFnPtr fp_cd = cd; // expected-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; // dev-error {{reference to __host__ function 'h' in __device__ function}}
+ HostFnPtr fp_h = h; // expected-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; // dev-error {{reference to __host__ function 'ch' in __device__ function}}
+ HostFnPtr fp_ch = ch; // expected-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; // dev-error {{reference to __global__ function 'g' in __device__ function}}
+ GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
g(); // expected-error {{no matching function for call to 'g'}}
- g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __device__ function}}
+ g<<<0,0>>>(); // expected-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; // dev-error {{reference to __host__ function 'h' in __global__ function}}
+ HostFnPtr fp_h = h; // expected-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; // dev-error {{reference to __host__ function 'ch' in __global__ function}}
+ HostFnPtr fp_ch = ch; // expected-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; // dev-error {{reference to __global__ function 'g' in __global__ function}}
+ GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
g(); // expected-error {{no matching function for call to 'g'}}
- g<<<0,0>>>(); // dev-error {{reference to __global__ function 'g' in __global__ function}}
+ g<<<0,0>>>(); // expected-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 48f7229df21f..83dce50b4af8 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=dev,expected %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %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); // dev-note {{'d1g' declared here}}
+__global__ void d1g(void); // expected-note {{'d1g' declared here}}
__device__ void d1(void) {
d1h(); // expected-error {{no matching function}}
d1d();
d1hd();
- d1g<<<1, 1>>>(); // dev-error {{reference to __global__ function 'd1g' in __device__ function}}
+ d1g<<<1, 1>>>(); // expected-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 d2e59b8033c3..8e5b7ddddb8f 100644
--- a/clang/test/SemaCUDA/implicit-device-lambda.cu
+++ b/clang/test/SemaCUDA/implicit-device-lambda.cu
@@ -1,7 +1,5 @@
-// 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
+// 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
#include "Inputs/cuda.h"
@@ -104,5 +102,5 @@ __device__ void foo() {
void foo() {}
};
X x;
- x.foo(); // dev-error {{reference to __host__ function 'foo' in __device__ function}}
+ x.foo(); // expected-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 85c27ce43632..8e17daa0c123 100644
--- a/clang/test/SemaCUDA/method-target.cu
+++ b/clang/test/SemaCUDA/method-target.cu
@@ -1,5 +1,4 @@
-// RUN: %clang_cc1 -fsyntax-only -verify=host,expected %s
-// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,expected %s
+// RUN: %clang_cc1 -fsyntax-only -verify %s
#include "Inputs/cuda.h"
@@ -7,11 +6,11 @@
// Test 1: host method called from device function
struct S1 {
- void method() {} // dev-note {{'method' declared here}}
+ void method() {} // expected-note {{'method' declared here}}
};
__device__ void foo1(S1& s) {
- s.method(); // dev-error {{reference to __host__ function 'method' in __device__ function}}
+ s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
}
//------------------------------------------------------------------------------
@@ -30,22 +29,22 @@ __device__ void foo2(S2& s, int i, float f) {
// Test 3: device method called from host function
struct S3 {
- __device__ void method() {} // host-note {{'method' declared here}}
+ __device__ void method() {} // expected-note {{'method' declared here}}
};
void foo3(S3& s) {
- s.method(); // host-error {{reference to __device__ function 'method' in __host__ function}}
+ s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}}
}
//------------------------------------------------------------------------------
// Test 4: device method called from host&device function
struct S4 {
- __device__ void method() {} // host-note {{'method' declared here}}
+ __device__ void method() {} // expected-note {{'method' declared here}}
};
__host__ __device__ void foo4(S4& s) {
- s.method(); // host-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+ s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
}
//------------------------------------------------------------------------------
@@ -64,9 +63,9 @@ __device__ void foo5(S5& s, S5& t) {
// Test 6: call method through pointer
struct S6 {
- void method() {} // dev-note {{'method' declared here}};
+ void method() {} // expected-note {{'method' declared here}};
};
__device__ void foo6(S6* s) {
- s->method(); // dev-error {{reference to __host__ function 'method' in __device__ function}}
+ s->method(); // expected-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 70a1cda6ab0c..e502d134b086 100644
--- a/clang/test/SemaCUDA/reference-to-kernel-fn.cu
+++ b/clang/test/SemaCUDA/reference-to-kernel-fn.cu
@@ -1,14 +1,12 @@
-// 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: %clang_cc1 -std=c++11 -fsyntax-only -verify \
// 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 {};
@@ -19,11 +17,13 @@ typedef void (*fn_ptr_t)();
__host__ __device__ fn_ptr_t get_ptr_hd() {
return kernel;
- // dev-error at -1 {{reference to __global__ function}}
+#ifdef DEVICE
+ // expected-error at -2 {{reference to __global__ function}}
+#endif
}
__host__ fn_ptr_t get_ptr_h() {
return kernel;
}
__device__ fn_ptr_t get_ptr_d() {
- return kernel; // dev-error {{reference to __global__ function}}
+ return kernel; // expected-error {{reference to __global__ function}}
}
More information about the cfe-commits
mailing list