1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
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;
}
|
