Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[CUDA] Reject calls to __device__ functions from host variable global…
… initializers. Reviewers: tra Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23335 llvm-svn: 278196
- Loading branch information
Justin Lebar
committed
Aug 10, 2016
1 parent
7d078bd
commit c989c3e
Showing
3 changed files
with
82 additions
and
28 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,32 @@ | ||
// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify | ||
|
||
#include "Inputs/cuda.h" | ||
|
||
// Check that we get an error if we try to call a __device__ function from a | ||
// module initializer. | ||
|
||
struct S { | ||
__device__ S() {} | ||
// expected-note@-1 {{'S' declared here}} | ||
}; | ||
|
||
S s; | ||
// expected-error@-1 {{reference to __device__ function 'S' in global initializer}} | ||
|
||
struct T { | ||
__host__ __device__ T() {} | ||
}; | ||
T t; // No error, this is OK. | ||
|
||
struct U { | ||
__host__ U() {} | ||
__device__ U(int) {} | ||
// expected-note@-1 {{'U' declared here}} | ||
}; | ||
U u(42); | ||
// expected-error@-1 {{reference to __device__ function 'U' in global initializer}} | ||
|
||
__device__ int device_fn() { return 42; } | ||
// expected-note@-1 {{'device_fn' declared here}} | ||
int n = device_fn(); | ||
// expected-error@-1 {{reference to __device__ function 'device_fn' in global initializer}} |