diff options
| author | Justin Lebar <jlebar@google.com> | 2016-10-12 01:30:08 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-10-12 01:30:08 +0000 |
| commit | f3d02c144a250d213f6e2882eea6d1e3ca5a995c (patch) | |
| tree | d963bab1b39962f2ecc34be93f7523f699038146 /lib/Sema/SemaCUDA.cpp | |
| parent | 683bf4db05fb5149c1f2bc3e8c508938ad6dccfa (diff) | |
| download | clang-f3d02c144a250d213f6e2882eea6d1e3ca5a995c.tar.gz | |
[CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.
Previously, this was an immediate, don't pass go, don't collect $200
error. But this precludes us from writing code like
__host__ __device__ void launch_kernel() {
kernel<<<...>>>();
}
Such code isn't wrong, following our notions of right and wrong in CUDA,
unless it's codegen'ed.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283963 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'lib/Sema/SemaCUDA.cpp')
| -rw-r--r-- | lib/Sema/SemaCUDA.cpp | 3 |
1 files changed, 1 insertions, 2 deletions
diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp index 9e101d16da..2a66124080 100644 --- a/lib/Sema/SemaCUDA.cpp +++ b/lib/Sema/SemaCUDA.cpp @@ -120,8 +120,7 @@ Sema::IdentifyCUDAPreference(const FunctionDecl *Caller, // (a) Can't call global from some contexts until we support CUDA's // dynamic parallelism. if (CalleeTarget == CFT_Global && - (CallerTarget == CFT_Global || CallerTarget == CFT_Device || - (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice))) + (CallerTarget == CFT_Global || CallerTarget == CFT_Device)) return CFP_Never; // (b) Calling HostDevice is OK for everyone. |