Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7
mirror-llvm-project/clang/lib/Sema/SemaAMDGPU.cpp
Matt Arsenault b03b97b828 clang/AMDGPU: Require 16-bit-insts for half typed buffer format builtins 
Typed buffer format load/store operations with 16-bit elements require
d16 support which was introduced in gfx8. These builtins previously had
no required features at all, so they were accepted (and then crashed the
backend) on targets without 16-bit support.

Diagnose these in Sema, parallel to the image builtins. The manual
verification here suprised me. The automatic builtin feature verification
is enforced in codegen, which seems like a layering violation which
should be fixed.

Co-Authored-By: Claude (Opus 4.8) <noreply@anthropic.com>
2026-08-01 21:42:57 +02:00

1096 lines
44 KiB
C++
Raw Blame History

//===------ 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/AST/Decl.h"
#include "clang/AST/DynamicRecursiveASTVisitor.h"
#include "clang/AST/Expr.h"
#include "clang/Basic/DiagnosticFrontend.h"
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/TargetBuiltins.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Sema/Ownership.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/Sema.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/StringMap.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/TargetParser/AMDGPUTargetParser.h"
#include <cstdint>
#include <utility>
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_raw_ptr_buffer_load_async_lds:
case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_lds:
case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_async_lds:
case AMDGPU::BI__builtin_amdgcn_load_to_lds:
case AMDGPU::BI__builtin_amdgcn_load_async_to_lds:
case AMDGPU::BI__builtin_amdgcn_global_load_lds:
case AMDGPU::BI__builtin_amdgcn_global_load_async_lds: {
constexpr const int SizeIdx = 2;
llvm::APSInt Size;
Expr *ArgExpr = TheCall->getArg(SizeIdx);
// Check for instantiation-dependent expressions (e.g., involving template
// parameters). These will be checked again during template instantiation.
if (ArgExpr->isInstantiationDependent())
return false;
[[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:
SemaRef.targetDiag(ArgExpr->getExprLoc(),
diag::err_amdgcn_load_lds_size_invalid_value)
<< ArgExpr->getSourceRange();
SemaRef.targetDiag(ArgExpr->getExprLoc(),
diag::note_amdgcn_load_lds_size_valid_value)
<< HasGFX950Insts << ArgExpr->getSourceRange();
return true;
}
}
case AMDGPU::BI__builtin_amdgcn_uicmp:
case AMDGPU::BI__builtin_amdgcn_uicmpl:
case AMDGPU::BI__builtin_amdgcn_sicmp:
case AMDGPU::BI__builtin_amdgcn_sicmpl:
case AMDGPU::BI__builtin_amdgcn_fcmp:
case AMDGPU::BI__builtin_amdgcn_fcmpf: {
// These builtins are deprecated in favor of
// __builtin_amdgcn_ballot_{w32|w64}. Suggest the replacement matching the
// wavefront size of the calling function.
bool IsWave32 = Builtin::evaluateRequiredTargetFeatures("wavefrontsize32",
CallerFeatureMap);
Diag(TheCall->getBeginLoc(), diag::warn_deprecated_builtin)
<< getASTContext().BuiltinInfo.getQuotedName(BuiltinID)
<< (IsWave32 ? "__builtin_amdgcn_ballot_w32"
: "__builtin_amdgcn_ballot_w64");
return false;
}
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_s_setreg:
return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0,
/*High=*/UINT16_MAX);
case AMDGPU::BI__builtin_amdgcn_s_wait_event: {
llvm::APSInt Result;
if (SemaRef.BuiltinConstantArg(TheCall, 0, Result))
return true;
bool IsGFX12Plus = Builtin::evaluateRequiredTargetFeatures(
"gfx12-insts", CallerFeatureMap);
// gfx11 -> gfx12 changed the interpretation of the bitmask. gfx12 inverted
// the intepretation for export_ready, but shifted the used bit by 1. Thus
// waiting for the export_ready event can use a value of 2 universally.
if (((IsGFX12Plus && !Result[1]) || (!IsGFX12Plus && Result[0])) ||
Result.getZExtValue() > 2) {
Expr *ArgExpr = TheCall->getArg(0);
SemaRef.targetDiag(ArgExpr->getExprLoc(),
diag::warn_amdgpu_s_wait_event_mask_no_effect_target)
<< ArgExpr->getSourceRange();
SemaRef.targetDiag(ArgExpr->getExprLoc(),
diag::note_amdgpu_s_wait_event_suggested_value)
<< ArgExpr->getSourceRange();
}
return false;
}
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, 15);
case AMDGPU::BI__builtin_amdgcn_av_load_b128:
return checkAVLoadStore(TheCall, /*IsStore=*/false);
case AMDGPU::BI__builtin_amdgcn_av_store_b128:
return checkAVLoadStore(TheCall, /*IsStore=*/true);
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);
case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128:
case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
return checkAtomicMonitorLoad(TheCall);
case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_format_v4f16:
case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_format_v4f16:
case AMDGPU::BI__builtin_amdgcn_struct_buffer_load_format_v4f16:
case AMDGPU::BI__builtin_amdgcn_struct_buffer_store_format_v4f16: {
StringRef FeatureList(
getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
CallerFeatureMap)) {
Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
<< FD->getDeclName() << FeatureList;
return false;
}
return false;
}
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: {
StringRef FeatureList(
getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
CallerFeatureMap)) {
Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
<< FD->getDeclName() << FeatureList;
return false;
}
unsigned ArgCount = TheCall->getNumArgs() - 1;
llvm::APSInt Result;
// Compilain about dmask values which are too huge to fully fit into 4 bits
// (which is the actual size of the dmask in corresponding HW instructions).
constexpr unsigned DMaskArgNo = 0;
constexpr int Low = 0;
constexpr int High = 15;
if (SemaRef.BuiltinConstantArg(TheCall, DMaskArgNo, Result) ||
SemaRef.BuiltinConstantArgRange(TheCall, DMaskArgNo, Low, High,
/* RangeIsError = */ true))
return true;
// Dmask indicates which elements should be returned and it is not possible
// to return more values than there are elements in return type.
int NumElementsInRetTy = 1;
const Type *RetTy = TheCall->getType().getTypePtr();
if (auto *VTy = dyn_cast<VectorType>(RetTy))
NumElementsInRetTy = VTy->getNumElements();
int NumActiveBitsInDMask =
llvm::popcount(static_cast<uint8_t>(Result.getExtValue()));
if (NumActiveBitsInDMask > NumElementsInRetTy) {
Diag(TheCall->getBeginLoc(),
diag::err_amdgcn_dmask_has_too_many_bits_set);
return true;
}
// For gather, only one bit can be set indicating which exact component to
// return.
bool ExtraGatherChecks =
BuiltinID == AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32 &&
SemaRef.BuiltinConstantArgPower2(TheCall, 0);
return ExtraGatherChecks ||
(SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
(SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
}
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
StringRef FeatureList(
getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
CallerFeatureMap)) {
Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
<< FD->getDeclName() << FeatureList;
return false;
}
unsigned ArgCount = TheCall->getNumArgs() - 1;
llvm::APSInt Result;
// Complain about dmask values which are too huge to fully fit into 4 bits
// (which is the actual size of the dmask in corresponding HW instructions).
constexpr unsigned DMaskArgNo = 1;
return SemaRef.BuiltinConstantArgRange(TheCall, DMaskArgNo, /*Low=*/0,
/*High=*/15,
/*RangeIsError=*/true) ||
SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result) ||
SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result);
}
case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8) {
if (SemaRef.checkArgCountRange(TheCall, 7, 8))
return true;
if (TheCall->getNumArgs() == 7)
return false;
} else if (BuiltinID ==
AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8) {
if (SemaRef.checkArgCountRange(TheCall, 8, 9))
return true;
if (TheCall->getNumArgs() == 8)
return false;
}
// Check if the last argument (clamp operand) is a constant and is
// convertible to bool.
Expr *ClampArg = TheCall->getArg(TheCall->getNumArgs() - 1);
// 1) Ensure clamp argument is a constant expression
llvm::APSInt ClampValue;
if (!SemaRef.VerifyIntegerConstantExpression(ClampArg, &ClampValue)
.isUsable())
return true;
// 2) Check if the argument can be converted to bool type
if (!SemaRef.Context.hasSameType(ClampArg->getType(),
SemaRef.Context.BoolTy)) {
// Try to convert to bool
QualType BoolTy = SemaRef.Context.BoolTy;
ExprResult ClampExpr(ClampArg);
SemaRef.CheckSingleAssignmentConstraints(BoolTy, ClampExpr);
if (ClampExpr.isInvalid())
return true;
}
return false;
}
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0,
/*High=*/0) ||
SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/2, /*Low=*/0,
/*High=*/0);
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 memory 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::checkAtomicOrderingCABIArg(Expr *E, bool MayLoad,
bool MayStore) {
Expr::EvalResult AtomicOrdArgRes;
if (!E->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) ||
(!(MayLoad && MayStore) && (Ord == llvm::AtomicOrderingCABI::acq_rel)) ||
(!MayLoad && Ord == llvm::AtomicOrderingCABI::acquire) ||
(!MayStore && Ord == llvm::AtomicOrderingCABI::release)) {
return Diag(E->getBeginLoc(), diag::warn_atomic_op_has_invalid_memory_order)
<< 0 << E->getSourceRange();
}
return false;
}
// Check that the first argument to TheCall is a global or generic pointer.
static bool checkGlobalOrFlatPointerArg(SemaAMDGPU &S, CallExpr *TheCall) {
Expr *PtrArg = TheCall->getArg(0);
QualType PtrTy = PtrArg->getType()->getPointeeType();
unsigned AS =
S.getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace());
if (AS != llvm::AMDGPUAS::FLAT_ADDRESS &&
AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) {
return S.Diag(TheCall->getBeginLoc(),
diag::err_amdgcn_global_or_flat_pointer_required)
<< PtrArg->getSourceRange();
}
return false;
}
static bool checkScopeAsInt(SemaAMDGPU &S, Expr *Scope) {
if (Scope->isValueDependent())
return false;
auto ScopeModel = AtomicScopeModel::create(AtomicScopeModelKind::Generic);
if (std::optional<llvm::APSInt> Result =
Scope->getIntegerConstantExpr(S.SemaRef.Context)) {
if (!ScopeModel->isValid(Result->getZExtValue())) {
return S.Diag(Scope->getBeginLoc(),
diag::err_atomic_op_has_invalid_sync_scope)
<< Scope->getSourceRange();
}
}
return false;
}
bool SemaAMDGPU::checkAVLoadStore(CallExpr *TheCall, bool IsStore) {
if (checkGlobalOrFlatPointerArg(*this, TheCall))
return true;
Expr *Scope = TheCall->getArg(TheCall->getNumArgs() - 1);
return checkScopeAsInt(*this, Scope);
}
bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) {
bool Fail = checkGlobalOrFlatPointerArg(*this, TheCall);
Expr *AO = TheCall->getArg(IsStore ? 2 : 1);
Expr *Scope = TheCall->getArg(TheCall->getNumArgs() - 1);
if (AO->isValueDependent() || Scope->isValueDependent())
return false;
// Check atomic ordering
Fail |=
checkAtomicOrderingCABIArg(TheCall->getArg(IsStore ? 2 : 1),
/*MayLoad=*/!IsStore, /*MayStore=*/IsStore);
// Last argument is the syncscope as a string literal.
if (!isa<StringLiteral>(Scope->IgnoreParenImpCasts())) {
Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal)
<< Scope->getSourceRange();
Fail = true;
}
return Fail;
}
bool SemaAMDGPU::checkAtomicMonitorLoad(CallExpr *TheCall) {
Expr *AO = TheCall->getArg(1);
Expr *Scope = TheCall->getArg(TheCall->getNumArgs() - 1);
if (AO->isValueDependent() || Scope->isValueDependent())
return false;
bool Fail = checkAtomicOrderingCABIArg(AO, /*MayLoad=*/true,
/*MayStore=*/false);
Fail |= checkScopeAsInt(*this, Scope);
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) {
Diag(AL.getLoc(), diag::warn_amdgpu_num_reg_attr_deprecated) << 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) {
Diag(AL.getLoc(), diag::warn_amdgpu_num_reg_attr_deprecated) << 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);
assert(!SemaRef.isSFINAEContext() &&
"Can't produce SFINAE diagnostic pointing to temporary attribute");
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);
}
Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) {
CallExpr *CE = cast<CallExpr>(E->IgnoreParens());
ASTContext &Ctx = getASTContext();
QualType BoolTy = Ctx.getLogicalOperationType();
SourceLocation Loc = CE->getExprLoc();
if (!CE->getBuiltinCallee())
return *ExpandedPredicates
.insert(SemaRef.BuildBoolLiteral(Loc, false).get())
.first;
bool P = false;
unsigned BI = CE->getBuiltinCallee();
if (Ctx.BuiltinInfo.isAuxBuiltinID(BI))
BI = Ctx.BuiltinInfo.getAuxBuiltinID(BI);
if (BI == AMDGPU::BI__builtin_amdgcn_processor_is) {
auto *GFX = dyn_cast<StringLiteral>(CE->getArg(0)->IgnoreParenCasts());
if (!GFX) {
Diag(Loc, diag::err_amdgcn_processor_is_arg_not_literal);
return nullptr;
}
StringRef N = GFX->getString();
const TargetInfo &TI = Ctx.getTargetInfo();
if (llvm::AMDGPU::parseArchAMDGCN(N) == llvm::AMDGPU::GK_NONE) {
Diag(Loc, diag::err_amdgcn_processor_is_arg_invalid_value) << N;
SmallVector<StringRef, 64> ValidList;
llvm::AMDGPU::fillValidArchListAMDGCN(ValidList);
if (!ValidList.empty())
Diag(Loc, diag::note_amdgcn_processor_is_valid_options)
<< llvm::join(ValidList, ", ");
return nullptr;
}
if (TI.getTriple().isSPIRV()) {
CE->setType(BoolTy);
return *ExpandedPredicates.insert(CE).first;
}
P = TI.isProcessorName(N);
} else {
Expr *Arg = CE->getArg(0);
if (!Arg || Arg->getType() != Ctx.BuiltinFnTy) {
Diag(Loc, diag::err_amdgcn_is_invocable_arg_invalid_value) << Arg;
return nullptr;
}
if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
CE->setType(BoolTy);
return *ExpandedPredicates.insert(CE).first;
}
auto *FD = cast<FunctionDecl>(Arg->getReferencedDeclOfCallee());
StringRef RF = Ctx.BuiltinInfo.getRequiredFeatures(FD->getBuiltinID());
llvm::StringMap<bool> CF;
Ctx.getFunctionFeatureMap(CF, FD);
P = Builtin::evaluateRequiredTargetFeatures(RF, CF);
}
return *ExpandedPredicates.insert(SemaRef.BuildBoolLiteral(Loc, P).get())
.first;
}
bool SemaAMDGPU::IsPredicate(Expr *E) const {
return ExpandedPredicates.contains(E);
}
void SemaAMDGPU::AddPotentiallyUnguardedBuiltinUser(FunctionDecl *FD) {
PotentiallyUnguardedBuiltinUsers.insert(FD);
}
bool SemaAMDGPU::HasPotentiallyUnguardedBuiltinUsage(FunctionDecl *FD) const {
return PotentiallyUnguardedBuiltinUsers.contains(FD);
}
namespace {
/// This class implements -Wamdgpu-unguarded-builtin-usage.
///
/// This is done with a traversal of the AST of a function that includes a
/// call to a target specific builtin. Whenever we encounter an \c if of the
/// form: \c if(__builtin_amdgcn_is_invocable), we consider the then statement
/// guarded.
class DiagnoseUnguardedBuiltins : public DynamicRecursiveASTVisitor {
// TODO: this could eventually be extended to consider what happens when there
// are multiple target architectures specified via target("arch=gfxXXX")
// target("arch=gfxyyy") etc., as well as feature disabling via "-XXX".
Sema &SemaRef;
SmallVector<StringRef> TargetFeatures;
SmallVector<std::pair<SourceLocation, StringRef>> CurrentGFXIP;
SmallVector<unsigned> GuardedBuiltins;
static Expr *FindPredicate(Expr *Cond) {
if (auto *CE = dyn_cast<CallExpr>(Cond)) {
if (CE->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_is_invocable ||
CE->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_processor_is)
return Cond;
} else if (auto *UO = dyn_cast<UnaryOperator>(Cond)) {
return FindPredicate(UO->getSubExpr());
} else if (auto *BO = dyn_cast<BinaryOperator>(Cond)) {
if ((Cond = FindPredicate(BO->getLHS())))
return Cond;
return FindPredicate(BO->getRHS());
}
return nullptr;
}
bool EnterPredicateGuardedContext(CallExpr *P);
void ExitPredicateGuardedContext(bool WasProcessorCheck);
bool TraverseGuardedStmt(Stmt *S, CallExpr *P);
public:
DiagnoseUnguardedBuiltins(Sema &SemaRef) : SemaRef(SemaRef) {
if (auto *TAT = SemaRef.getCurFunctionDecl(true)->getAttr<TargetAttr>()) {
// We use the somewhat misnamed x86 accessors because they provide exactly
// what we require.
TAT->getX86AddedFeatures(TargetFeatures);
if (auto GFXIP = TAT->getX86Architecture())
CurrentGFXIP.emplace_back(TAT->getLocation(), *GFXIP);
}
}
bool TraverseLambdaExpr(LambdaExpr *LE) override {
if (SemaRef.AMDGPU().HasPotentiallyUnguardedBuiltinUsage(
LE->getCallOperator()))
return true; // We have already handled this.
return DynamicRecursiveASTVisitor::TraverseLambdaExpr(LE);
}
bool TraverseStmt(Stmt *S) override {
if (!S)
return true;
return DynamicRecursiveASTVisitor::TraverseStmt(S);
}
void IssueDiagnostics(Stmt *S) { TraverseStmt(S); }
bool TraverseIfStmt(IfStmt *If) override {
if (auto *CE = dyn_cast_or_null<CallExpr>(FindPredicate(If->getCond())))
return TraverseGuardedStmt(If, CE);
return DynamicRecursiveASTVisitor::TraverseIfStmt(If);
}
bool TraverseCaseStmt(CaseStmt *CS) override {
return TraverseStmt(CS->getSubStmt());
}
bool TraverseConditionalOperator(ConditionalOperator *CO) override {
if (auto *CE = dyn_cast_or_null<CallExpr>(FindPredicate(CO->getCond())))
return TraverseGuardedStmt(CO, CE);
return DynamicRecursiveASTVisitor::TraverseConditionalOperator(CO);
}
bool VisitAsmStmt(AsmStmt *ASM) override;
bool VisitCallExpr(CallExpr *CE) override;
};
bool DiagnoseUnguardedBuiltins::EnterPredicateGuardedContext(CallExpr *P) {
bool IsProcessorCheck =
P->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_processor_is;
if (IsProcessorCheck) {
StringRef G = cast<clang::StringLiteral>(P->getArg(0))->getString();
// TODO: handle generic ISAs.
if (!CurrentGFXIP.empty() && G != CurrentGFXIP.back().second) {
SemaRef.Diag(P->getExprLoc(),
diag::err_amdgcn_conflicting_is_processor_options)
<< P;
SemaRef.Diag(CurrentGFXIP.back().first,
diag::note_amdgcn_previous_is_processor_guard);
}
CurrentGFXIP.emplace_back(P->getExprLoc(), G);
} else {
auto *FD = cast<FunctionDecl>(
cast<DeclRefExpr>(P->getArg(0))->getReferencedDeclOfCallee());
GuardedBuiltins.push_back(FD->getBuiltinID());
}
return IsProcessorCheck;
}
void DiagnoseUnguardedBuiltins::ExitPredicateGuardedContext(bool WasProcCheck) {
if (WasProcCheck)
CurrentGFXIP.pop_back();
else
GuardedBuiltins.pop_back();
}
inline std::pair<Stmt *, Stmt *> GetTraversalOrder(Stmt *S) {
std::pair<Stmt *, Stmt *> Ordered;
Expr *Condition = nullptr;
if (auto *CO = dyn_cast<ConditionalOperator>(S)) {
Condition = CO->getCond();
Ordered = {CO->getTrueExpr(), CO->getFalseExpr()};
} else if (auto *If = dyn_cast<IfStmt>(S)) {
Condition = If->getCond();
Ordered = {If->getThen(), If->getElse()};
}
if (auto *UO = dyn_cast<UnaryOperator>(Condition))
if (UO->getOpcode() == UnaryOperatorKind::UO_LNot)
std::swap(Ordered.first, Ordered.second);
return Ordered;
}
bool DiagnoseUnguardedBuiltins::TraverseGuardedStmt(Stmt *S, CallExpr *P) {
assert(S && "Unexpected missing Statement!");
assert(P && "Unexpected missing Predicate!");
auto [Guarded, Unguarded] = GetTraversalOrder(S);
bool WasProcessorCheck = EnterPredicateGuardedContext(P);
bool Continue = TraverseStmt(Guarded);
ExitPredicateGuardedContext(WasProcessorCheck);
return Continue && TraverseStmt(Unguarded);
}
bool DiagnoseUnguardedBuiltins::VisitAsmStmt(AsmStmt *ASM) {
// TODO: should we check if the ASM is valid for the target? Can we?
if (!CurrentGFXIP.empty())
return true;
std::string S = ASM->generateAsmString(SemaRef.getASTContext());
SemaRef.Diag(ASM->getAsmLoc(), diag::warn_amdgcn_unguarded_asm_stmt) << S;
SemaRef.Diag(ASM->getAsmLoc(), diag::note_amdgcn_unguarded_asm_silence) << S;
return true;
}
bool DiagnoseUnguardedBuiltins::VisitCallExpr(CallExpr *CE) {
unsigned ID = CE->getBuiltinCallee();
Builtin::Context &BInfo = SemaRef.getASTContext().BuiltinInfo;
if (!ID)
return true;
if (!BInfo.isTSBuiltin(ID))
return true;
if (ID == AMDGPU::BI__builtin_amdgcn_processor_is ||
ID == AMDGPU::BI__builtin_amdgcn_is_invocable)
return true;
if (llvm::find(GuardedBuiltins, ID) != GuardedBuiltins.end())
return true;
StringRef FL(BInfo.getRequiredFeatures(ID));
llvm::StringMap<bool> FeatureMap;
if (CurrentGFXIP.empty()) {
for (auto &&F : TargetFeatures)
FeatureMap[F] = true;
for (auto &&GID : GuardedBuiltins)
for (auto &&F : llvm::split(BInfo.getRequiredFeatures(GID), ','))
FeatureMap[F] = true;
} else {
static const llvm::Triple AMDGCN(llvm::Triple::amdgpu,
llvm::Triple::NoSubArch, llvm::Triple::AMD,
llvm::Triple::AMDHSA);
llvm::AMDGPU::fillAMDGPUFeatureMap(CurrentGFXIP.back().second, AMDGCN,
FeatureMap);
}
FunctionDecl *BI = CE->getDirectCallee();
SourceLocation BICallLoc = CE->getExprLoc();
if (Builtin::evaluateRequiredTargetFeatures(FL, FeatureMap)) {
SemaRef.Diag(BICallLoc, diag::warn_amdgcn_unguarded_builtin) << BI;
SemaRef.Diag(BICallLoc, diag::note_amdgcn_unguarded_builtin_silence) << BI;
} else {
StringRef GFXIP = CurrentGFXIP.empty() ? "" : CurrentGFXIP.back().second;
SemaRef.Diag(BICallLoc, diag::err_amdgcn_incompatible_builtin)
<< BI << FL << !CurrentGFXIP.empty() << GFXIP;
if (!CurrentGFXIP.empty())
SemaRef.Diag(CurrentGFXIP.back().first,
diag::note_amdgcn_previous_is_processor_guard);
}
return true;
}
} // Unnamed namespace
void SemaAMDGPU::DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD) {
DiagnoseUnguardedBuiltins(SemaRef).IssueDiagnostics(FD->getBody());
}
} // namespace clang
Reference in a new issue View git blame Copy permalink