Eli Bendersky | 7325e56 | 2014-09-03 15:27:03 +0000 | [diff] [blame] | 1 | //===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===// |
| 2 | // |
| 3 | // The LLVM Compiler Infrastructure |
| 4 | // |
| 5 | // This file is distributed under the University of Illinois Open Source |
| 6 | // License. See LICENSE.TXT for details. |
| 7 | // |
| 8 | //===----------------------------------------------------------------------===// |
| 9 | /// \file |
| 10 | /// \brief This file implements semantic analysis for CUDA constructs. |
| 11 | /// |
| 12 | //===----------------------------------------------------------------------===// |
| 13 | |
| 14 | #include "clang/Sema/Sema.h" |
| 15 | #include "clang/AST/ASTContext.h" |
| 16 | #include "clang/AST/Decl.h" |
| 17 | #include "clang/Sema/SemaDiagnostic.h" |
| 18 | using namespace clang; |
| 19 | |
| 20 | ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, |
| 21 | MultiExprArg ExecConfig, |
| 22 | SourceLocation GGGLoc) { |
| 23 | FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl(); |
| 24 | if (!ConfigDecl) |
| 25 | return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use) |
| 26 | << "cudaConfigureCall"); |
| 27 | QualType ConfigQTy = ConfigDecl->getType(); |
| 28 | |
| 29 | DeclRefExpr *ConfigDR = new (Context) |
| 30 | DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc); |
| 31 | MarkFunctionReferenced(LLLLoc, ConfigDecl); |
| 32 | |
| 33 | return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr, |
| 34 | /*IsExecConfig=*/true); |
| 35 | } |
| 36 | |
| 37 | /// IdentifyCUDATarget - Determine the CUDA compilation target for this function |
| 38 | Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { |
| 39 | // Implicitly declared functions (e.g. copy constructors) are |
| 40 | // __host__ __device__ |
| 41 | if (D->isImplicit()) |
| 42 | return CFT_HostDevice; |
| 43 | |
| 44 | if (D->hasAttr<CUDAGlobalAttr>()) |
| 45 | return CFT_Global; |
| 46 | |
| 47 | if (D->hasAttr<CUDADeviceAttr>()) { |
| 48 | if (D->hasAttr<CUDAHostAttr>()) |
| 49 | return CFT_HostDevice; |
| 50 | return CFT_Device; |
| 51 | } |
| 52 | |
| 53 | return CFT_Host; |
| 54 | } |
| 55 | |
| 56 | bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget, |
| 57 | CUDAFunctionTarget CalleeTarget) { |
| 58 | // CUDA B.1.1 "The __device__ qualifier declares a function that is... |
| 59 | // Callable from the device only." |
| 60 | if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device) |
| 61 | return true; |
| 62 | |
| 63 | // CUDA B.1.2 "The __global__ qualifier declares a function that is... |
| 64 | // Callable from the host only." |
| 65 | // CUDA B.1.3 "The __host__ qualifier declares a function that is... |
| 66 | // Callable from the host only." |
| 67 | if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) && |
| 68 | (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)) |
| 69 | return true; |
| 70 | |
| 71 | if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) |
| 72 | return true; |
| 73 | |
| 74 | return false; |
| 75 | } |
| 76 | |