maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/SemaCUDA/call-device-fn-from-host.cu

86 lines
2.6 KiB
Plaintext
Raw Normal View History

[CUDA] When we emit an error that might have been deferred, also print a callstack. Summary: Previously, when you did something not allowed in a host+device function and then caused it to be codegen'ed, we would print out an error telling you that you did something bad, but we wouldn't tell you how we decided that the function needed to be codegen'ed. This change causes us to print out a callstack when emitting deferred errors. This is immensely helpful when debugging highly-templated code, where it's often unclear how a function became known-emitted. We only print the callstack once per function, after we print the all deferred errors. This patch also switches all of our hashtables to using canonical FunctionDecls instead of regular FunctionDecls. This prevents a number of bugs, some of which are caught by tests added here, in which we assume that two FDs for the same function have the same pointer value. Reviewers: rnk Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D25704 llvm-svn: 284647
2016-10-20 05:15:01 +08:00
// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
// RUN: -verify -verify-ignore-unexpected=note
[CUDA] Raise an error if a wrong-side call is codegen'ed. Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
2016-08-16 07:00:49 +08:00
// Note: This test won't work with -fsyntax-only, because some of these errors
// are emitted during codegen.
#include "Inputs/cuda.h"
__device__ void device_fn() {}
[CUDA] Simplify some repeated diagnostic expectations in CUDA tests. Instead of repeating the diagnostic, use "expected-note N". Test-only change. llvm-svn: 284882
2016-10-22 04:50:47 +08:00
// expected-note@-1 5 {{'device_fn' declared here}}
[CUDA] Raise an error if a wrong-side call is codegen'ed. Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
2016-08-16 07:00:49 +08:00
struct S {
__device__ S() {}
[CUDA] Simplify some repeated diagnostic expectations in CUDA tests. Instead of repeating the diagnostic, use "expected-note N". Test-only change. llvm-svn: 284882
2016-10-22 04:50:47 +08:00
// expected-note@-1 2 {{'S' declared here}}
[CUDA] Raise an error if a wrong-side call is codegen'ed. Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
2016-08-16 07:00:49 +08:00
__device__ ~S() { device_fn(); }
[CUDA] Fix "declared here" note on deferred wrong-side errors. Previously we weren't deferring these "declared here" notes, which is obviously wrong. llvm-svn: 278767
2016-08-16 08:48:21 +08:00
// expected-note@-1 {{'~S' declared here}}
[CUDA] Raise an error if a wrong-side call is codegen'ed. Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
2016-08-16 07:00:49 +08:00
int x;
};
struct T {
__host__ __device__ void hd() { device_fn(); }
// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
// No error; this is (implicitly) inline and is never called, so isn't
// codegen'ed.
__host__ __device__ void hd2() { device_fn(); }
__host__ __device__ void hd3();
__device__ void d() {}
[CUDA] Fix "declared here" note on deferred wrong-side errors. Previously we weren't deferring these "declared here" notes, which is obviously wrong. llvm-svn: 278767
2016-08-16 08:48:21 +08:00
// expected-note@-1 {{'d' declared here}}
[CUDA] Raise an error if a wrong-side call is codegen'ed. Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
2016-08-16 07:00:49 +08:00
};
__host__ __device__ void T::hd3() {
device_fn();
// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
}
template <typename T> __host__ __device__ void hd2() { device_fn(); }
// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
void host_fn() { hd2<int>(); }
__host__ __device__ void hd() { device_fn(); }
// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
// No error because this is never instantiated.
template <typename T> __host__ __device__ void hd3() { device_fn(); }
__host__ __device__ void local_var() {
S s;
// expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}}
}
__host__ __device__ void placement_new(char *ptr) {
::new(ptr) S();
// expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}}
}
__host__ __device__ void explicit_destructor(S *s) {
s->~S();
// expected-error@-1 {{reference to __device__ function '~S' in __host__ __device__ function}}
}
__host__ __device__ void hd_member_fn() {
T t;
// Necessary to trigger an error on T::hd. It's (implicitly) inline, so
// isn't codegen'ed until we call it.
t.hd();
}
__host__ __device__ void h_member_fn() {
T t;
t.d();
// expected-error@-1 {{reference to __device__ function 'd' in __host__ __device__ function}}
}
__host__ __device__ void fn_ptr() {
auto* ptr = &device_fn;
// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
}
template <typename T>
__host__ __device__ void fn_ptr_template() {
auto* ptr = &device_fn; // Not an error because the template isn't instantiated.
}