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 
77