xref: /llvm-project/clang/lib/Sema/SemaAMDGPU.cpp (revision a2c3e0c4cb08e9cf26f4a745777295e1eb562e1b) 
1 //===------ SemaAMDGPU.cpp ------- AMDGPU target-specific routines --------===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 //  This file implements semantic analysis functions specific to AMDGPU.
10 //
11 //===----------------------------------------------------------------------===//
12 
13 #include "clang/Sema/SemaAMDGPU.h"
14 #include "clang/Basic/DiagnosticSema.h"
15 #include "clang/Basic/TargetBuiltins.h"
16 #include "clang/Sema/Ownership.h"
17 #include "clang/Sema/Sema.h"
18 #include "llvm/Support/AtomicOrdering.h"
19 #include <cstdint>
20 
21 namespace clang {
22 
23 SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
24 
25 bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
26                                                 CallExpr *TheCall) {
27   // position of memory order and scope arguments in the builtin
28   unsigned OrderIndex, ScopeIndex;
29 
30   const auto *FD = SemaRef.getCurFunctionDecl();
31   assert(FD && "AMDGPU builtins should not be used outside of a function");
32   llvm::StringMap<bool> CallerFeatureMap;
33   getASTContext().getFunctionFeatureMap(CallerFeatureMap, FD);
34   bool HasGFX950Insts =
35       Builtin::evaluateRequiredTargetFeatures("gfx950-insts", CallerFeatureMap);
36 
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);
42     [[maybe_unused]] ExprResult R =
43         SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);
44     assert(!R.isInvalid());
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:
57       Diag(ArgExpr->getExprLoc(),
58            diag::err_amdgcn_global_load_lds_size_invalid_value)
59           << ArgExpr->getSourceRange();
60       Diag(ArgExpr->getExprLoc(),
61            diag::note_amdgcn_global_load_lds_size_valid_value)
62           << HasGFX950Insts << ArgExpr->getSourceRange();
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:
81     return checkMovDPPFunctionCall(TheCall, 5, 1);
82   case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
83     return checkMovDPPFunctionCall(TheCall, 2, 1);
84   case AMDGPU::BI__builtin_amdgcn_update_dpp: {
85     return checkMovDPPFunctionCall(TheCall, 6, 2);
86   }
87   default:
88     return false;
89   }
90 
91   ExprResult Arg = TheCall->getArg(OrderIndex);
92   auto ArgExpr = Arg.get();
93   Expr::EvalResult ArgResult;
94 
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();
99 
100   // Check validity of memory ordering as per C11 / C++11's memody model.
101   // Only fence needs check. Atomic dec/inc allow all memory orders.
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_cast<llvm::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();
123   Expr::EvalResult ArgResult1;
124   // Check that sync scope is a constant literal
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 
132 bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
133                                          unsigned NumDataArgs) {
134   assert(NumDataArgs <= 2);
135   if (SemaRef.checkArgCountRange(TheCall, NumArgs, NumArgs))
136     return true;
137   Expr *Args[2];
138   QualType ArgTys[2];
139   for (unsigned I = 0; I != NumDataArgs; ++I) {
140     Args[I] = TheCall->getArg(I);
141     ArgTys[I] = Args[I]->getType();
142     // TODO: Vectors can also be supported.
143     if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
144       SemaRef.Diag(Args[I]->getBeginLoc(),
145                    diag::err_typecheck_cond_expect_int_float)
146           << ArgTys[I] << Args[I]->getSourceRange();
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())) &&
160       getASTContext().getTypeSize(ArgTys[0]) ==
161           getASTContext().getTypeSize(ArgTys[1]))
162     return false;
163 
164   SemaRef.Diag(Args[1]->getBeginLoc(),
165                diag::err_typecheck_call_different_arg_types)
166       << ArgTys[0] << ArgTys[1];
167   return true;
168 }
169 
170 static bool
171 checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
172                                       const AMDGPUFlatWorkGroupSizeAttr &Attr) {
173   // Accept template arguments for now as they depend on something else.
174   // We'll get to check them when they eventually get instantiated.
175   if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
176     return false;
177 
178   uint32_t Min = 0;
179   if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
180     return true;
181 
182   uint32_t Max = 0;
183   if (!S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
184     return true;
185 
186   if (Min == 0 && Max != 0) {
187     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
188         << &Attr << 0;
189     return true;
190   }
191   if (Min > Max) {
192     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
193         << &Attr << 1;
194     return true;
195   }
196 
197   return false;
198 }
199 
200 AMDGPUFlatWorkGroupSizeAttr *
201 SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
202                                               Expr *MinExpr, Expr *MaxExpr) {
203   ASTContext &Context = getASTContext();
204   AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
205 
206   if (checkAMDGPUFlatWorkGroupSizeArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
207     return nullptr;
208   return ::new (Context)
209       AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
210 }
211 
212 void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
213                                                 const AttributeCommonInfo &CI,
214                                                 Expr *MinExpr, Expr *MaxExpr) {
215   if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
216     D->addAttr(Attr);
217 }
218 
219 void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
220                                                    const ParsedAttr &AL) {
221   Expr *MinExpr = AL.getArgAsExpr(0);
222   Expr *MaxExpr = AL.getArgAsExpr(1);
223 
224   addAMDGPUFlatWorkGroupSizeAttr(D, AL, MinExpr, MaxExpr);
225 }
226 
227 static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
228                                            Expr *MaxExpr,
229                                            const AMDGPUWavesPerEUAttr &Attr) {
230   if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
231       (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
232     return true;
233 
234   // Accept template arguments for now as they depend on something else.
235   // We'll get to check them when they eventually get instantiated.
236   if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
237     return false;
238 
239   uint32_t Min = 0;
240   if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
241     return true;
242 
243   uint32_t Max = 0;
244   if (MaxExpr && !S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
245     return true;
246 
247   if (Min == 0 && Max != 0) {
248     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
249         << &Attr << 0;
250     return true;
251   }
252   if (Max != 0 && Min > Max) {
253     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
254         << &Attr << 1;
255     return true;
256   }
257 
258   return false;
259 }
260 
261 AMDGPUWavesPerEUAttr *
262 SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
263                                        Expr *MinExpr, Expr *MaxExpr) {
264   ASTContext &Context = getASTContext();
265   AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
266 
267   if (checkAMDGPUWavesPerEUArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
268     return nullptr;
269 
270   return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
271 }
272 
273 void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
274                                          Expr *MinExpr, Expr *MaxExpr) {
275   if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
276     D->addAttr(Attr);
277 }
278 
279 void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
280   if (!AL.checkAtLeastNumArgs(SemaRef, 1) || !AL.checkAtMostNumArgs(SemaRef, 2))
281     return;
282 
283   Expr *MinExpr = AL.getArgAsExpr(0);
284   Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
285 
286   addAMDGPUWavesPerEUAttr(D, AL, MinExpr, MaxExpr);
287 }
288 
289 void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
290   uint32_t NumSGPR = 0;
291   Expr *NumSGPRExpr = AL.getArgAsExpr(0);
292   if (!SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))
293     return;
294 
295   D->addAttr(::new (getASTContext())
296                  AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
297 }
298 
299 void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
300   uint32_t NumVGPR = 0;
301   Expr *NumVGPRExpr = AL.getArgAsExpr(0);
302   if (!SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))
303     return;
304 
305   D->addAttr(::new (getASTContext())
306                  AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
307 }
308 
309 static bool
310 checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
311                                      Expr *ZExpr,
312                                      const AMDGPUMaxNumWorkGroupsAttr &Attr) {
313   if (S.DiagnoseUnexpandedParameterPack(XExpr) ||
314       (YExpr && S.DiagnoseUnexpandedParameterPack(YExpr)) ||
315       (ZExpr && S.DiagnoseUnexpandedParameterPack(ZExpr)))
316     return true;
317 
318   // Accept template arguments for now as they depend on something else.
319   // We'll get to check them when they eventually get instantiated.
320   if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
321       (ZExpr && ZExpr->isValueDependent()))
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]) {
328       if (!S.checkUInt32Argument(Attr, Exprs[i], NumWG, i,
329                                  /*StrictlyUnsigned=*/true))
330         return true;
331       if (NumWG == 0) {
332         S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
333             << &Attr << Exprs[i]->getSourceRange();
334         return true;
335       }
336     }
337   }
338 
339   return false;
340 }
341 
342 AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
343     const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
344   ASTContext &Context = getASTContext();
345   AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
346 
347   if (checkAMDGPUMaxNumWorkGroupsArguments(SemaRef, XExpr, YExpr, ZExpr,
348                                            TmpAttr))
349     return nullptr;
350 
351   return ::new (Context)
352       AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
353 }
354 
355 void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
356                                                const AttributeCommonInfo &CI,
357                                                Expr *XExpr, Expr *YExpr,
358                                                Expr *ZExpr) {
359   if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
360     D->addAttr(Attr);
361 }
362 
363 void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
364                                                   const ParsedAttr &AL) {
365   Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
366   Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(2) : nullptr;
367   addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
368 }
369 
370 } // namespace clang
371