Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-04-04 14:23:00 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

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

Browse Source 
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
This commit is contained in:
Yaxun (Sam) Liu 2020-07-15 13:25:32 -04:00
parent 9a0689e072
commit 4fc752b30b
8 changed files with 51 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/lib/Sema/SemaCUDA.cpp
View File

@ -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.

12
clang/test/SemaCUDA/builtins.cu
View File

@ -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);
}

4
clang/test/SemaCUDA/call-kernel-from-kernel.cu
View File

@ -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"

34
clang/test/SemaCUDA/function-overload.cu
View File

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

6
clang/test/SemaCUDA/function-target.cu
View File

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

8
clang/test/SemaCUDA/implicit-device-lambda.cu
View File

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

19
clang/test/SemaCUDA/method-target.cu
View File

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

14
clang/test/SemaCUDA/reference-to-kernel-fn.cu
View File

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