Author: eliben Date: Wed Sep 3 10:27:03 2014 New Revision: 217043 URL: http://llvm.org/viewvc/llvm-project?rev=217043&view=rev Log: Split off CUDA-specific Sema parts to a new file
In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 Added: cfe/trunk/lib/Sema/SemaCUDA.cpp Modified: cfe/trunk/lib/Sema/CMakeLists.txt cfe/trunk/lib/Sema/SemaDeclCXX.cpp cfe/trunk/lib/Sema/SemaExpr.cpp Modified: cfe/trunk/lib/Sema/CMakeLists.txt URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/CMakeLists.txt?rev=217043&r1=217042&r2=217043&view=diff ============================================================================== --- cfe/trunk/lib/Sema/CMakeLists.txt (original) +++ cfe/trunk/lib/Sema/CMakeLists.txt Wed Sep 3 10:27:03 2014 @@ -21,6 +21,7 @@ add_clang_library(clangSema SemaChecking.cpp SemaCodeComplete.cpp SemaConsumer.cpp + SemaCUDA.cpp SemaDecl.cpp SemaDeclAttr.cpp SemaDeclCXX.cpp Added: cfe/trunk/lib/Sema/SemaCUDA.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=217043&view=auto ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (added) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Sep 3 10:27:03 2014 @@ -0,0 +1,76 @@ +//===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +/// \file +/// \brief This file implements semantic analysis for CUDA constructs. +/// +//===----------------------------------------------------------------------===// + +#include "clang/Sema/Sema.h" +#include "clang/AST/ASTContext.h" +#include "clang/AST/Decl.h" +#include "clang/Sema/SemaDiagnostic.h" +using namespace clang; + +ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, + MultiExprArg ExecConfig, + SourceLocation GGGLoc) { + FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl(); + if (!ConfigDecl) + return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use) + << "cudaConfigureCall"); + QualType ConfigQTy = ConfigDecl->getType(); + + DeclRefExpr *ConfigDR = new (Context) + DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc); + MarkFunctionReferenced(LLLLoc, ConfigDecl); + + return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr, + /*IsExecConfig=*/true); +} + +/// IdentifyCUDATarget - Determine the CUDA compilation target for this function +Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { + // Implicitly declared functions (e.g. copy constructors) are + // __host__ __device__ + if (D->isImplicit()) + return CFT_HostDevice; + + if (D->hasAttr<CUDAGlobalAttr>()) + return CFT_Global; + + if (D->hasAttr<CUDADeviceAttr>()) { + if (D->hasAttr<CUDAHostAttr>()) + return CFT_HostDevice; + return CFT_Device; + } + + return CFT_Host; +} + +bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget, + CUDAFunctionTarget CalleeTarget) { + // CUDA B.1.1 "The __device__ qualifier declares a function that is... + // Callable from the device only." + if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device) + return true; + + // CUDA B.1.2 "The __global__ qualifier declares a function that is... + // Callable from the host only." + // CUDA B.1.3 "The __host__ qualifier declares a function that is... + // Callable from the host only." + if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) && + (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)) + return true; + + if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) + return true; + + return false; +} + Modified: cfe/trunk/lib/Sema/SemaDeclCXX.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclCXX.cpp?rev=217043&r1=217042&r2=217043&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDeclCXX.cpp (original) +++ cfe/trunk/lib/Sema/SemaDeclCXX.cpp Wed Sep 3 10:27:03 2014 @@ -13101,46 +13101,6 @@ Sema::checkExceptionSpecification(Except } } -/// IdentifyCUDATarget - Determine the CUDA compilation target for this function -Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { - // Implicitly declared functions (e.g. copy constructors) are - // __host__ __device__ - if (D->isImplicit()) - return CFT_HostDevice; - - if (D->hasAttr<CUDAGlobalAttr>()) - return CFT_Global; - - if (D->hasAttr<CUDADeviceAttr>()) { - if (D->hasAttr<CUDAHostAttr>()) - return CFT_HostDevice; - return CFT_Device; - } - - return CFT_Host; -} - -bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget, - CUDAFunctionTarget CalleeTarget) { - // CUDA B.1.1 "The __device__ qualifier declares a function that is... - // Callable from the device only." - if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device) - return true; - - // CUDA B.1.2 "The __global__ qualifier declares a function that is... - // Callable from the host only." - // CUDA B.1.3 "The __host__ qualifier declares a function that is... - // Callable from the host only." - if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) && - (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)) - return true; - - if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) - return true; - - return false; -} - /// HandleMSProperty - Analyze a __delcspec(property) field of a C++ class. /// MSPropertyDecl *Sema::HandleMSProperty(Scope *S, RecordDecl *Record, Modified: cfe/trunk/lib/Sema/SemaExpr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=217043&r1=217042&r2=217043&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaExpr.cpp (original) +++ cfe/trunk/lib/Sema/SemaExpr.cpp Wed Sep 3 10:27:03 2014 @@ -4588,23 +4588,6 @@ Sema::ActOnCallExpr(Scope *S, Expr *Fn, ExecConfig, IsExecConfig); } -ExprResult -Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, - MultiExprArg ExecConfig, SourceLocation GGGLoc) { - FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl(); - if (!ConfigDecl) - return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use) - << "cudaConfigureCall"); - QualType ConfigQTy = ConfigDecl->getType(); - - DeclRefExpr *ConfigDR = new (Context) DeclRefExpr( - ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc); - MarkFunctionReferenced(LLLLoc, ConfigDecl); - - return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr, - /*IsExecConfig=*/true); -} - /// ActOnAsTypeExpr - create a new asType (bitcast) from the arguments. /// /// __builtin_astype( value, dst type ) _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
