xref: /freebsd/contrib/llvm-project/clang/lib/Sema/SemaAMDGPU.cpp (revision 5036d9652a5701d00e9e40ea942c278e9f77d33d) 
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   switch (BuiltinID) {
30   case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
31     constexpr const int SizeIdx = 2;
32     llvm::APSInt Size;
33     Expr *ArgExpr = TheCall->getArg(SizeIdx);
34     [[maybe_unused]] ExprResult R =
35         SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);
36     assert(!R.isInvalid());
37     switch (Size.getSExtValue()) {
38     case 1:
39     case 2:
40     case 4:
41       return false;
42     default:
43       Diag(ArgExpr->getExprLoc(),
44            diag::err_amdgcn_global_load_lds_size_invalid_value)
45           << ArgExpr->getSourceRange();
46       Diag(ArgExpr->getExprLoc(),
47            diag::note_amdgcn_global_load_lds_size_valid_value)
48           << ArgExpr->getSourceRange();
49       return true;
50     }
51   }
52   case AMDGPU::BI__builtin_amdgcn_get_fpenv:
53   case AMDGPU::BI__builtin_amdgcn_set_fpenv:
54     return false;
55   case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
56   case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
57   case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
58   case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
59     OrderIndex = 2;
60     ScopeIndex = 3;
61     break;
62   case AMDGPU::BI__builtin_amdgcn_fence:
63     OrderIndex = 0;
64     ScopeIndex = 1;
65     break;
66   default:
67     return false;
68   }
69 
70   ExprResult Arg = TheCall->getArg(OrderIndex);
71   auto ArgExpr = Arg.get();
72   Expr::EvalResult ArgResult;
73 
74   if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))
75     return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
76            << ArgExpr->getType();
77   auto Ord = ArgResult.Val.getInt().getZExtValue();
78 
79   // Check validity of memory ordering as per C11 / C++11's memody model.
80   // Only fence needs check. Atomic dec/inc allow all memory orders.
81   if (!llvm::isValidAtomicOrderingCABI(Ord))
82     return Diag(ArgExpr->getBeginLoc(),
83                 diag::warn_atomic_op_has_invalid_memory_order)
84            << 0 << ArgExpr->getSourceRange();
85   switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
86   case llvm::AtomicOrderingCABI::relaxed:
87   case llvm::AtomicOrderingCABI::consume:
88     if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
89       return Diag(ArgExpr->getBeginLoc(),
90                   diag::warn_atomic_op_has_invalid_memory_order)
91              << 0 << ArgExpr->getSourceRange();
92     break;
93   case llvm::AtomicOrderingCABI::acquire:
94   case llvm::AtomicOrderingCABI::release:
95   case llvm::AtomicOrderingCABI::acq_rel:
96   case llvm::AtomicOrderingCABI::seq_cst:
97     break;
98   }
99 
100   Arg = TheCall->getArg(ScopeIndex);
101   ArgExpr = Arg.get();
102   Expr::EvalResult ArgResult1;
103   // Check that sync scope is a constant literal
104   if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))
105     return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
106            << ArgExpr->getType();
107 
108   return false;
109 }
110 
111 static bool
112 checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
113                                       const AMDGPUFlatWorkGroupSizeAttr &Attr) {
114   // Accept template arguments for now as they depend on something else.
115   // We'll get to check them when they eventually get instantiated.
116   if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
117     return false;
118 
119   uint32_t Min = 0;
120   if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
121     return true;
122 
123   uint32_t Max = 0;
124   if (!S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
125     return true;
126 
127   if (Min == 0 && Max != 0) {
128     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
129         << &Attr << 0;
130     return true;
131   }
132   if (Min > Max) {
133     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
134         << &Attr << 1;
135     return true;
136   }
137 
138   return false;
139 }
140 
141 AMDGPUFlatWorkGroupSizeAttr *
142 SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
143                                               Expr *MinExpr, Expr *MaxExpr) {
144   ASTContext &Context = getASTContext();
145   AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
146 
147   if (checkAMDGPUFlatWorkGroupSizeArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
148     return nullptr;
149   return ::new (Context)
150       AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
151 }
152 
153 void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
154                                                 const AttributeCommonInfo &CI,
155                                                 Expr *MinExpr, Expr *MaxExpr) {
156   if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
157     D->addAttr(Attr);
158 }
159 
160 void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
161                                                    const ParsedAttr &AL) {
162   Expr *MinExpr = AL.getArgAsExpr(0);
163   Expr *MaxExpr = AL.getArgAsExpr(1);
164 
165   addAMDGPUFlatWorkGroupSizeAttr(D, AL, MinExpr, MaxExpr);
166 }
167 
168 static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
169                                            Expr *MaxExpr,
170                                            const AMDGPUWavesPerEUAttr &Attr) {
171   if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
172       (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
173     return true;
174 
175   // Accept template arguments for now as they depend on something else.
176   // We'll get to check them when they eventually get instantiated.
177   if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
178     return false;
179 
180   uint32_t Min = 0;
181   if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
182     return true;
183 
184   uint32_t Max = 0;
185   if (MaxExpr && !S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
186     return true;
187 
188   if (Min == 0 && Max != 0) {
189     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
190         << &Attr << 0;
191     return true;
192   }
193   if (Max != 0 && Min > Max) {
194     S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
195         << &Attr << 1;
196     return true;
197   }
198 
199   return false;
200 }
201 
202 AMDGPUWavesPerEUAttr *
203 SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
204                                        Expr *MinExpr, Expr *MaxExpr) {
205   ASTContext &Context = getASTContext();
206   AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
207 
208   if (checkAMDGPUWavesPerEUArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
209     return nullptr;
210 
211   return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
212 }
213 
214 void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
215                                          Expr *MinExpr, Expr *MaxExpr) {
216   if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
217     D->addAttr(Attr);
218 }
219 
220 void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
221   if (!AL.checkAtLeastNumArgs(SemaRef, 1) || !AL.checkAtMostNumArgs(SemaRef, 2))
222     return;
223 
224   Expr *MinExpr = AL.getArgAsExpr(0);
225   Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
226 
227   addAMDGPUWavesPerEUAttr(D, AL, MinExpr, MaxExpr);
228 }
229 
230 void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
231   uint32_t NumSGPR = 0;
232   Expr *NumSGPRExpr = AL.getArgAsExpr(0);
233   if (!SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))
234     return;
235 
236   D->addAttr(::new (getASTContext())
237                  AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
238 }
239 
240 void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
241   uint32_t NumVGPR = 0;
242   Expr *NumVGPRExpr = AL.getArgAsExpr(0);
243   if (!SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))
244     return;
245 
246   D->addAttr(::new (getASTContext())
247                  AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
248 }
249 
250 static bool
251 checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
252                                      Expr *ZExpr,
253                                      const AMDGPUMaxNumWorkGroupsAttr &Attr) {
254   if (S.DiagnoseUnexpandedParameterPack(XExpr) ||
255       (YExpr && S.DiagnoseUnexpandedParameterPack(YExpr)) ||
256       (ZExpr && S.DiagnoseUnexpandedParameterPack(ZExpr)))
257     return true;
258 
259   // Accept template arguments for now as they depend on something else.
260   // We'll get to check them when they eventually get instantiated.
261   if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
262       (ZExpr && ZExpr->isValueDependent()))
263     return false;
264 
265   uint32_t NumWG = 0;
266   Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
267   for (int i = 0; i < 3; i++) {
268     if (Exprs[i]) {
269       if (!S.checkUInt32Argument(Attr, Exprs[i], NumWG, i,
270                                  /*StrictlyUnsigned=*/true))
271         return true;
272       if (NumWG == 0) {
273         S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
274             << &Attr << Exprs[i]->getSourceRange();
275         return true;
276       }
277     }
278   }
279 
280   return false;
281 }
282 
283 AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
284     const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
285   ASTContext &Context = getASTContext();
286   AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
287 
288   if (checkAMDGPUMaxNumWorkGroupsArguments(SemaRef, XExpr, YExpr, ZExpr,
289                                            TmpAttr))
290     return nullptr;
291 
292   return ::new (Context)
293       AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
294 }
295 
296 void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
297                                                const AttributeCommonInfo &CI,
298                                                Expr *XExpr, Expr *YExpr,
299                                                Expr *ZExpr) {
300   if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
301     D->addAttr(Attr);
302 }
303 
304 void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
305                                                   const ParsedAttr &AL) {
306   Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
307   Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(2) : nullptr;
308   addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
309 }
310 
311 } // namespace clang
312