From f3d02c144a250d213f6e2882eea6d1e3ca5a995c Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Wed, 12 Oct 2016 01:30:08 +0000 Subject: [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 --- lib/Sema/SemaCUDA.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) (limited to 'lib/Sema/SemaCUDA.cpp') 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. -- cgit v1.2.1