From 7325e56c109bd9741650a388e9f6de8d32b241e1 Mon Sep 17 00:00:00 2001 From: Eli Bendersky Date: Wed, 3 Sep 2014 15:27:03 +0000 Subject: [PATCH] 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 llvm-svn: 217043 --- clang/lib/Sema/CMakeLists.txt | 1 + clang/lib/Sema/SemaCUDA.cpp | 76 ++++++++++++++++++++++++++++++++++ clang/lib/Sema/SemaDeclCXX.cpp | 40 ------------------ clang/lib/Sema/SemaExpr.cpp | 17 -------- 4 files changed, 77 insertions(+), 57 deletions(-) create mode 100644 clang/lib/Sema/SemaCUDA.cpp diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt index 7847d2c36e5b..4a772d8972a9 100644 --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -21,6 +21,7 @@ add_clang_library(clangSema SemaChecking.cpp SemaCodeComplete.cpp SemaConsumer.cpp + SemaCUDA.cpp SemaDecl.cpp SemaDeclAttr.cpp SemaDeclCXX.cpp diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp new file mode 100644 index 000000000000..f8ff42924693 --- /dev/null +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -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()) + return CFT_Global; + + if (D->hasAttr()) { + if (D->hasAttr()) + 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; +} + diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 4ca12cf8bf1a..5da654610593 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -13101,46 +13101,6 @@ Sema::checkExceptionSpecification(ExceptionSpecificationType EST, } } -/// 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()) - return CFT_Global; - - if (D->hasAttr()) { - if (D->hasAttr()) - 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, diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index d9278863534d..c23bbf0b12f3 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -4588,23 +4588,6 @@ Sema::ActOnCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc, 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 )