clang/lib/Sema/SemaAMDGPU.cpp

   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