https://github.com/tcgu-amd created https://github.com/llvm/llvm-project/pull/158145
Addresses https://github.com/ROCm/ROCm/issues/5253. The builtin expects a vector _Float16 of length 2, but clang does not emit errors when the wrong argument types is supplied (e.g. half2). This causes the compilation to pass but error during LTO. This fix the issue by adding sema checks to the builtins. >From 2192052dff2df318a35b8e2b28c7cf4f75070734 Mon Sep 17 00:00:00 2001 From: Tim Gu <[email protected]> Date: Thu, 11 Sep 2025 15:26:31 -0400 Subject: [PATCH] Added clang sema check to ensure atomic_fadd_v2f16 builtin uses v2f16 arg. --- clang/include/clang/Sema/SemaAMDGPU.h | 2 ++ clang/lib/Sema/SemaAMDGPU.cpp | 32 +++++++++++++++++++++++++++ 2 files changed, 34 insertions(+) diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index bac812a9d4fcf..9ca4418349fff 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -31,6 +31,8 @@ class SemaAMDGPU : public SemaBase { bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs, unsigned NumDataArgs); + bool checkAMDGCNAtomicFaddV2F16Type(CallExpr *TheCall); + /// Create an AMDGPUWavesPerEUAttr attribute. AMDGPUFlatWorkGroupSizeAttr * CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min, diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index baba503239e9f..c0f17f185982e 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -109,6 +109,10 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, 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_global_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: + return checkAMDGCNAtomicFaddV2F16Type(TheCall); default: return false; } @@ -436,4 +440,32 @@ void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D, addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr); } + +bool SemaAMDGPU::checkAMDGCNAtomicFaddV2F16Type(CallExpr *TheCall) { + // Check that the pointer argument is a pointer to v2f16 + + Expr *Arg = TheCall->getArg(1); + QualType ArgType = Arg->getType(); + + // Check if it's a vector type + if (!ArgType->isVectorType()) { + Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types) + << "expected _Float16 vector of length 2" << ArgType + << Arg->getSourceRange(); + return true; + } + + const VectorType *VT = ArgType->getAs<VectorType>(); + + // Check element type (should be _Float16) and vector length (should be 2) + QualType ElementType = VT->getElementType(); + if (!ElementType->isFloat16Type() || VT->getNumElements() != 2) { + Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types) + << "expected _Float16 vector of length 2" << ArgType + << Arg->getSourceRange(); + return true; + } + + return false; +} } // namespace clang _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
