SemaAMDGPU.cpp Source File (original) (raw)

18#include "llvm/Support/AtomicOrdering.h"

19#include

28 unsigned OrderIndex, ScopeIndex;

31 assert(FD && "AMDGPU builtins should not be used outside of a function");

32 llvm::StringMap CallerFeatureMap;

34 bool HasGFX950Insts =

37 switch (BuiltinID) {

38 case AMDGPU::BI__builtin_amdgcn_global_load_lds: {

39 constexpr const int SizeIdx = 2;

40 llvm::APSInt Size;

41 Expr *ArgExpr = TheCall->getArg(SizeIdx);

45 switch (Size.getSExtValue()) {

46 case 1:

47 case 2:

48 case 4:

49 return false;

50 case 12:

51 case 16: {

52 if (HasGFX950Insts)

53 return false;

54 [[fallthrough]];

55 }

56 default:

58 diag::err_amdgcn_global_load_lds_size_invalid_value)

61 diag::note_amdgcn_global_load_lds_size_valid_value)

63 return true;

64 }

65 }

66 case AMDGPU::BI__builtin_amdgcn_get_fpenv:

67 case AMDGPU::BI__builtin_amdgcn_set_fpenv:

68 return false;

69 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:

70 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:

71 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:

72 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:

73 OrderIndex = 2;

74 ScopeIndex = 3;

75 break;

76 case AMDGPU::BI__builtin_amdgcn_fence:

77 OrderIndex = 0;

78 ScopeIndex = 1;

79 break;

80 case AMDGPU::BI__builtin_amdgcn_mov_dpp:

82 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:

84 case AMDGPU::BI__builtin_amdgcn_update_dpp: {

86 }

87 default:

88 return false;

89 }

92 auto ArgExpr = Arg.get();

95 if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))

96 return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)

97 << ArgExpr->getType();

98 auto Ord = ArgResult.Val.getInt().getZExtValue();

100

101

102 if (!llvm::isValidAtomicOrderingCABI(Ord))

103 return Diag(ArgExpr->getBeginLoc(),

104 diag::warn_atomic_op_has_invalid_memory_order)

105 << 0 << ArgExpr->getSourceRange();

106 switch (static_castllvm::AtomicOrderingCABI\(Ord)) {

107 case llvm::AtomicOrderingCABI::relaxed:

108 case llvm::AtomicOrderingCABI::consume:

109 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)

110 return Diag(ArgExpr->getBeginLoc(),

111 diag::warn_atomic_op_has_invalid_memory_order)

112 << 0 << ArgExpr->getSourceRange();

113 break;

114 case llvm::AtomicOrderingCABI::acquire:

115 case llvm::AtomicOrderingCABI::release:

116 case llvm::AtomicOrderingCABI::acq_rel:

117 case llvm::AtomicOrderingCABI::seq_cst:

118 break;

119 }

120

121 Arg = TheCall->getArg(ScopeIndex);

122 ArgExpr = Arg.get();

124

125 if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))

126 return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)

127 << ArgExpr->getType();

128

129 return false;

130}

131

133 unsigned NumDataArgs) {

134 assert(NumDataArgs <= 2);

136 return true;

137 Expr *Args[2];

139 for (unsigned I = 0; I != NumDataArgs; ++I) {

140 Args[I] = TheCall->getArg(I);

141 ArgTys[I] = Args[I]->getType();

142

143 if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {

145 diag::err_typecheck_cond_expect_int_float)

147 return true;

148 }

149 }

150 if (NumDataArgs < 2)

151 return false;

152

153 if (getASTContext().hasSameUnqualifiedType(ArgTys[0], ArgTys[1]))

154 return false;

155

156 if (((ArgTys[0]->isUnsignedIntegerType() &&

157 ArgTys[1]->isSignedIntegerType()) ||

158 (ArgTys[0]->isSignedIntegerType() &&

159 ArgTys[1]->isUnsignedIntegerType())) &&

162 return false;

163

165 diag::err_typecheck_call_different_arg_types)

166 << ArgTys[0] << ArgTys[1];

167 return true;

168}

169

170static bool

172 const AMDGPUFlatWorkGroupSizeAttr &Attr) {

173

174

176 return false;

177

178 uint32_t Min = 0;

180 return true;

181

182 uint32_t Max = 0;

184 return true;

185

186 if (Min == 0 && Max != 0) {

188 << &Attr << 0;

189 return true;

190 }

193 << &Attr << 1;

194 return true;

195 }

196

197 return false;

198}

199

200AMDGPUFlatWorkGroupSizeAttr *

202 Expr *MinExpr, Expr *MaxExpr) {

204 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);

205

207 return nullptr;

208 return ::new (Context)

209 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);

210}

211

214 Expr *MinExpr, Expr *MaxExpr) {

216 D->addAttr(Attr);

217}

218

223

225}

226

228 Expr *MaxExpr,

229 const AMDGPUWavesPerEUAttr &Attr) {

232 return true;

233

234

235

237 return false;

238

239 uint32_t Min = 0;

241 return true;

242

243 uint32_t Max = 0;

245 return true;

246

247 if (Min == 0 && Max != 0) {

249 << &Attr << 0;

250 return true;

251 }

254 << &Attr << 1;

255 return true;

256 }

257

258 return false;

259}

260

261AMDGPUWavesPerEUAttr *

263 Expr *MinExpr, Expr *MaxExpr) {

265 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);

266

268 return nullptr;

269

270 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);

271}

272

274 Expr *MinExpr, Expr *MaxExpr) {

276 D->addAttr(Attr);

277}

278

281 return;

282

285

287}

288

290 uint32_t NumSGPR = 0;

293 return;

294

296 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));

297}

298

300 uint32_t NumVGPR = 0;

303 return;

304

306 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));

307}

308

309static bool

312 const AMDGPUMaxNumWorkGroupsAttr &Attr) {

316 return true;

317

318

319

322 return false;

323

324 uint32_t NumWG = 0;

325 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};

326 for (int i = 0; i < 3; i++) {

327 if (Exprs[i]) {

329 true))

330 return true;

331 if (NumWG == 0) {

332 S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)

334 return true;

335 }

336 }

337 }

338

339 return false;

340}

341

345 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);

346

348 TmpAttr))

349 return nullptr;

350

351 return ::new (Context)

352 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);

353}

354

358 Expr *ZExpr) {

360 D->addAttr(Attr);

361}

362

368}

369

370}

This file declares semantic analysis functions specific to AMDGPU.

Enumerates target-specific builtins in their own namespaces within namespace clang.

Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...

void getFunctionFeatureMap(llvm::StringMap< bool > &FeatureMap, const FunctionDecl *) const

Attr - This represents one attribute.

SourceLocation getLocation() const

SourceLocation getLoc() const

CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).

Expr * getArg(unsigned Arg)

getArg - Return the specified argument.

Decl - This represents one declaration (or definition), e.g.

This represents one expression.

bool isValueDependent() const

Determines whether the value of this expression depends on.

SourceLocation getExprLoc() const LLVM_READONLY

getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...

ParsedAttr - Represents a syntactic attribute.

unsigned getNumArgs() const

getNumArgs - Return the number of actual arguments to this attribute.

Expr * getArgAsExpr(unsigned Arg) const

bool checkAtLeastNumArgs(class Sema &S, unsigned Num) const

Check if the attribute has at least as many args as Num.

bool checkAtMostNumArgs(class Sema &S, unsigned Num) const

Check if the attribute has at most as many args as Num.

A (possibly-)qualified type.

void handleAMDGPUMaxNumWorkGroupsAttr(Decl *D, const ParsedAttr &AL)

void addAMDGPUFlatWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max)

addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size attribute to a particular declar...

void handleAMDGPUFlatWorkGroupSizeAttr(Decl *D, const ParsedAttr &AL)

void handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL)

AMDGPUMaxNumWorkGroupsAttr * CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr)

Create an AMDGPUMaxNumWorkGroupsAttr attribute.

AMDGPUWavesPerEUAttr * CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI, Expr *Min, Expr *Max)

Create an AMDGPUWavesPerEUAttr attribute.

void handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL)

AMDGPUFlatWorkGroupSizeAttr * CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min, Expr *Max)

Create an AMDGPUWavesPerEUAttr attribute.

bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs, unsigned NumDataArgs)

void handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL)

bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall)

void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max)

addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a particular declaration.

void addAMDGPUMaxNumWorkGroupsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr)

addAMDGPUMaxNumWorkGroupsAttr - Adds an amdgpu_max_num_work_groups attribute to a particular declarat...

SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID, bool DeferHint=false)

Emit a diagnostic.

ASTContext & getASTContext() const

Sema - This implements semantic analysis and AST building for C.

ExprResult VerifyIntegerConstantExpression(Expr *E, llvm::APSInt *Result, VerifyICEDiagnoser &Diagnoser, AllowFoldKind CanFold=NoFold)

VerifyIntegerConstantExpression - Verifies that an expression is an ICE, and reports the appropriate ...

FunctionDecl * getCurFunctionDecl(bool AllowLambda=false) const

Returns a pointer to the innermost enclosing function, or nullptr if the current context is not insid...

bool DiagnoseUnexpandedParameterPack(SourceLocation Loc, TypeSourceInfo *T, UnexpandedParameterPackContext UPPC)

If the given type contains an unexpanded parameter pack, diagnose the error.

bool checkUInt32Argument(const AttrInfo &AI, const Expr *Expr, uint32_t &Val, unsigned Idx=UINT_MAX, bool StrictlyUnsigned=false)

If Expr is a valid integer constant, get the value of the integer expression and return success or fa...

bool checkArgCountRange(CallExpr *Call, unsigned MinArgCount, unsigned MaxArgCount)

Checks that a call expression's argument count is in the desired range.

SourceRange getSourceRange() const LLVM_READONLY

SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...

bool evaluateRequiredTargetFeatures(llvm::StringRef RequiredFatures, const llvm::StringMap< bool > &TargetFetureMap)

Returns true if the required target features of a builtin function are enabled.

The JSON file list parser is used to communicate input to InstallAPI.

static bool checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr, Expr *ZExpr, const AMDGPUMaxNumWorkGroupsAttr &Attr)

static bool checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, const AMDGPUFlatWorkGroupSizeAttr &Attr)

static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, const AMDGPUWavesPerEUAttr &Attr)

EvalResult is a struct with detailed info about an evaluated expression.

APValue Val

Val - This is the value the expression can be folded to.