[CUDA] Only allow __global__ on free functions and static member functions.
Summary:
Warn for NVCC compatibility if you declare a static member function or
inline function as __global__.
Reviewers: tra
Subscribers: jhen, echristo, cfe-commits
Differential Revision: http://reviews.llvm.org/D16261
llvm-svn: 258263
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 61dfdd3..568c765 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -273,12 +273,9 @@
resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
Sema::CUDAFunctionTarget Target2,
Sema::CUDAFunctionTarget *ResolvedTarget) {
- if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
- // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
- // Clang should detect this earlier and produce an error. Then this
- // condition can be changed to an assertion.
- return true;
- }
+ // Only free functions and static member functions may be global.
+ assert(Target1 != Sema::CFT_Global);
+ assert(Target2 != Sema::CFT_Global);
if (Target1 == Sema::CFT_HostDevice) {
*ResolvedTarget = Target2;