================
@@ -249,7 +249,37 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned
BuiltinID,
unsigned ArgCount = TheCall->getNumArgs() - 1;
llvm::APSInt Result;
- return (SemaRef.BuiltinConstantArg(TheCall, 0, 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;
+ }
+
+ bool ExtraGatherChecks = false;
+ // For gather, only one bit can be set indicating which exact component to
+ // return.
+ if (BuiltinID == AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32)
+ ExtraGatherChecks = SemaRef.BuiltinConstantArgPower2(TheCall, 0);
----------------
arsenm wrote:
```suggestion
// 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);
```
https://github.com/llvm/llvm-project/pull/180949
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
