r217043 - Split off CUDA-specific Sema parts to a new file
Eli Bendersky
eliben at google.com
Wed Sep 3 08:27:03 PDT 2014
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 )
More information about the cfe-commits
mailing list