SemaAMDGPU.cpp

Go to the documentation of this file.

//===------ SemaAMDGPU.cpp ------- AMDGPU target-specific routines --------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
//  This file implements semantic analysis functions specific to AMDGPU.
//
//===----------------------------------------------------------------------===//
 
#include "clang/Sema/SemaAMDGPU.h"
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/TargetBuiltins.h"
#include "clang/Sema/Ownership.h"
#include "clang/Sema/Sema.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/AtomicOrdering.h"
#include <cstdint>
 
namespace clang {
 
SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
 
bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
                                                CallExpr *TheCall) {
  // position of memory order and scope arguments in the builtin
  unsigned OrderIndex, ScopeIndex;
 
  const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
  assert(FD && "AMDGPU builtins should not be used outside of a function");
  llvm::StringMap<bool> CallerFeatureMap;
  getASTContext().getFunctionFeatureMap(CallerFeatureMap, FD);
  bool HasGFX950Insts =
      Builtin::evaluateRequiredTargetFeatures("gfx950-insts", CallerFeatureMap);
 
  switch (BuiltinID) {
  case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_load_lds:
  case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_lds:
  case AMDGPU::BI__builtin_amdgcn_load_to_lds:
  case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
    constexpr const int SizeIdx = 2;
    llvm::APSInt Size;
    Expr *ArgExpr = TheCall->getArg(SizeIdx);
    [[maybe_unused]] ExprResult R =
        SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);
    assert(!R.isInvalid());
    switch (Size.getSExtValue()) {
    case 1:
    case 2:
    case 4:
      return false;
    case 12:
    case 16: {
      if (HasGFX950Insts)
        return false;
      [[fallthrough]];
    }
    default:
      Diag(ArgExpr->getExprLoc(), diag::err_amdgcn_load_lds_size_invalid_value)
          << ArgExpr->getSourceRange();
      Diag(ArgExpr->getExprLoc(), diag::note_amdgcn_load_lds_size_valid_value)
          << HasGFX950Insts << ArgExpr->getSourceRange();
      return true;
    }
  }
  case AMDGPU::BI__builtin_amdgcn_get_fpenv:
  case AMDGPU::BI__builtin_amdgcn_set_fpenv:
    return false;
  case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
  case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
  case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
  case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
    OrderIndex = 2;
    ScopeIndex = 3;
    break;
  case AMDGPU::BI__builtin_amdgcn_fence:
    OrderIndex = 0;
    ScopeIndex = 1;
    break;
  case AMDGPU::BI__builtin_amdgcn_mov_dpp:
    return checkMovDPPFunctionCall(TheCall, 5, 1);
  case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
    return checkMovDPPFunctionCall(TheCall, 2, 1);
  case AMDGPU::BI__builtin_amdgcn_update_dpp:
    return checkMovDPPFunctionCall(TheCall, 6, 2);
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp8:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp8:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_bf8:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_bf8:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp4:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp4:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp8:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_bf8:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp4:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_fp6:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_fp6:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_bf6:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_bf6:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
  case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
    return SemaRef.BuiltinConstantArgRange(TheCall, 2, 0, 7);
  case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
  case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
  case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
    return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/false);
  case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
  case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
  case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
    return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
  default:
    return false;
  }
 
  ExprResult Arg = TheCall->getArg(OrderIndex);
  auto ArgExpr = Arg.get();
  Expr::EvalResult ArgResult;
 
  if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))
    return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
           << ArgExpr->getType();
  auto Ord = ArgResult.Val.getInt().getZExtValue();
 
  // Check validity of memory ordering as per C11 / C++11's memody model.
  // Only fence needs check. Atomic dec/inc allow all memory orders.
  if (!llvm::isValidAtomicOrderingCABI(Ord))
    return Diag(ArgExpr->getBeginLoc(),
                diag::warn_atomic_op_has_invalid_memory_order)
           << 0 << ArgExpr->getSourceRange();
  switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
  case llvm::AtomicOrderingCABI::relaxed:
  case llvm::AtomicOrderingCABI::consume:
    if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
      return Diag(ArgExpr->getBeginLoc(),
                  diag::warn_atomic_op_has_invalid_memory_order)
             << 0 << ArgExpr->getSourceRange();
    break;
  case llvm::AtomicOrderingCABI::acquire:
  case llvm::AtomicOrderingCABI::release:
  case llvm::AtomicOrderingCABI::acq_rel:
  case llvm::AtomicOrderingCABI::seq_cst:
    break;
  }
 
  Arg = TheCall->getArg(ScopeIndex);
  ArgExpr = Arg.get();
  Expr::EvalResult ArgResult1;
  // Check that sync scope is a constant literal
  if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))
    return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
           << ArgExpr->getType();
 
  return false;
}
 
bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) {
  bool Fail = false;
 
  // First argument is a global or generic pointer.
  Expr *PtrArg = TheCall->getArg(0);
  QualType PtrTy = PtrArg->getType()->getPointeeType();
  unsigned AS = getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace());
  if (AS != llvm::AMDGPUAS::FLAT_ADDRESS &&
      AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) {
    Fail = true;
    Diag(TheCall->getBeginLoc(), diag::err_amdgcn_coop_atomic_invalid_as)
        << PtrArg->getSourceRange();
  }
 
  // Check atomic ordering
  Expr *AtomicOrdArg = TheCall->getArg(IsStore ? 2 : 1);
  Expr::EvalResult AtomicOrdArgRes;
  if (!AtomicOrdArg->EvaluateAsInt(AtomicOrdArgRes, getASTContext()))
    llvm_unreachable("Intrinsic requires imm for atomic ordering argument!");
  auto Ord =
      llvm::AtomicOrderingCABI(AtomicOrdArgRes.Val.getInt().getZExtValue());
 
  // Atomic ordering cannot be acq_rel in any case, acquire for stores or
  // release for loads.
  if (!llvm::isValidAtomicOrderingCABI((unsigned)Ord) ||
      (Ord == llvm::AtomicOrderingCABI::acq_rel) ||
      Ord == (IsStore ? llvm::AtomicOrderingCABI::acquire
                      : llvm::AtomicOrderingCABI::release)) {
    return Diag(AtomicOrdArg->getBeginLoc(),
                diag::warn_atomic_op_has_invalid_memory_order)
           << 0 << AtomicOrdArg->getSourceRange();
  }
 
  // Last argument is a string literal
  Expr *Arg = TheCall->getArg(TheCall->getNumArgs() - 1);
  if (!isa<StringLiteral>(Arg->IgnoreParenImpCasts())) {
    Fail = true;
    Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal)
        << Arg->getSourceRange();
  }
 
  return Fail;
}
 
bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
                                         unsigned NumDataArgs) {
  assert(NumDataArgs <= 2);
  if (SemaRef.checkArgCountRange(TheCall, NumArgs, NumArgs))
    return true;
  Expr *Args[2];
  QualType ArgTys[2];
  for (unsigned I = 0; I != NumDataArgs; ++I) {
    Args[I] = TheCall->getArg(I);
    ArgTys[I] = Args[I]->getType();
    // TODO: Vectors can also be supported.
    if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
      SemaRef.Diag(Args[I]->getBeginLoc(),
                   diag::err_typecheck_cond_expect_int_float)
          << ArgTys[I] << Args[I]->getSourceRange();
      return true;
    }
  }
  if (NumDataArgs < 2)
    return false;
 
  if (getASTContext().hasSameUnqualifiedType(ArgTys[0], ArgTys[1]))
    return false;
 
  if (((ArgTys[0]->isUnsignedIntegerType() &&
        ArgTys[1]->isSignedIntegerType()) ||
       (ArgTys[0]->isSignedIntegerType() &&
        ArgTys[1]->isUnsignedIntegerType())) &&
      getASTContext().getTypeSize(ArgTys[0]) ==
          getASTContext().getTypeSize(ArgTys[1]))
    return false;
 
  SemaRef.Diag(Args[1]->getBeginLoc(),
               diag::err_typecheck_call_different_arg_types)
      << ArgTys[0] << ArgTys[1];
  return true;
}
 
static bool
checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
                                      const AMDGPUFlatWorkGroupSizeAttr &Attr) {
  // Accept template arguments for now as they depend on something else.
  // We'll get to check them when they eventually get instantiated.
  if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
    return false;
 
  uint32_t Min = 0;
  if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
    return true;
 
  uint32_t Max = 0;
  if (!S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
    return true;
 
  if (Min == 0 && Max != 0) {
    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
        << &Attr << 0;
    return true;
  }
  if (Min > Max) {
    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
        << &Attr << 1;
    return true;
  }
 
  return false;
}
 
AMDGPUFlatWorkGroupSizeAttr *
SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
                                              Expr *MinExpr, Expr *MaxExpr) {
  ASTContext &Context = getASTContext();
  AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
 
  if (checkAMDGPUFlatWorkGroupSizeArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
    return nullptr;
  return ::new (Context)
      AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
}
 
void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
                                                const AttributeCommonInfo &CI,
                                                Expr *MinExpr, Expr *MaxExpr) {
  if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
    D->addAttr(Attr);
}
 
void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
                                                   const ParsedAttr &AL) {
  Expr *MinExpr = AL.getArgAsExpr(0);
  Expr *MaxExpr = AL.getArgAsExpr(1);
 
  addAMDGPUFlatWorkGroupSizeAttr(D, AL, MinExpr, MaxExpr);
}
 
static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
                                           Expr *MaxExpr,
                                           const AMDGPUWavesPerEUAttr &Attr) {
  if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
      (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
    return true;
 
  // Accept template arguments for now as they depend on something else.
  // We'll get to check them when they eventually get instantiated.
  if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
    return false;
 
  uint32_t Min = 0;
  if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
    return true;
 
  uint32_t Max = 0;
  if (MaxExpr && !S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
    return true;
 
  if (Min == 0 && Max != 0) {
    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
        << &Attr << 0;
    return true;
  }
  if (Max != 0 && Min > Max) {
    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
        << &Attr << 1;
    return true;
  }
 
  return false;
}
 
AMDGPUWavesPerEUAttr *
SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
                                       Expr *MinExpr, Expr *MaxExpr) {
  ASTContext &Context = getASTContext();
  AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
 
  if (checkAMDGPUWavesPerEUArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
    return nullptr;
 
  return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
}
 
void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
                                         Expr *MinExpr, Expr *MaxExpr) {
  if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
    D->addAttr(Attr);
}
 
void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
  if (!AL.checkAtLeastNumArgs(SemaRef, 1) || !AL.checkAtMostNumArgs(SemaRef, 2))
    return;
 
  Expr *MinExpr = AL.getArgAsExpr(0);
  Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
 
  addAMDGPUWavesPerEUAttr(D, AL, MinExpr, MaxExpr);
}
 
void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
  uint32_t NumSGPR = 0;
  Expr *NumSGPRExpr = AL.getArgAsExpr(0);
  if (!SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))
    return;
 
  D->addAttr(::new (getASTContext())
                 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
}
 
void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
  uint32_t NumVGPR = 0;
  Expr *NumVGPRExpr = AL.getArgAsExpr(0);
  if (!SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))
    return;
 
  D->addAttr(::new (getASTContext())
                 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
}
 
static bool
checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
                                     Expr *ZExpr,
                                     const AMDGPUMaxNumWorkGroupsAttr &Attr) {
  if (S.DiagnoseUnexpandedParameterPack(XExpr) ||
      (YExpr && S.DiagnoseUnexpandedParameterPack(YExpr)) ||
      (ZExpr && S.DiagnoseUnexpandedParameterPack(ZExpr)))
    return true;
 
  // Accept template arguments for now as they depend on something else.
  // We'll get to check them when they eventually get instantiated.
  if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
      (ZExpr && ZExpr->isValueDependent()))
    return false;
 
  uint32_t NumWG = 0;
  Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
  for (int i = 0; i < 3; i++) {
    if (Exprs[i]) {
      if (!S.checkUInt32Argument(Attr, Exprs[i], NumWG, i,
                                 /*StrictlyUnsigned=*/true))
        return true;
      if (NumWG == 0) {
        S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
            << &Attr << Exprs[i]->getSourceRange();
        return true;
      }
    }
  }
 
  return false;
}
 
AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
    const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
  ASTContext &Context = getASTContext();
  AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
 
  if (checkAMDGPUMaxNumWorkGroupsArguments(SemaRef, XExpr, YExpr, ZExpr,
                                           TmpAttr))
    return nullptr;
 
  return ::new (Context)
      AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
}
 
void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
                                               const AttributeCommonInfo &CI,
                                               Expr *XExpr, Expr *YExpr,
                                               Expr *ZExpr) {
  if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
    D->addAttr(Attr);
}
 
void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
                                                  const ParsedAttr &AL) {
  Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
  Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(2) : nullptr;
  addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
}
 
} // namespace clang