[clang] [AMDGPU] Add sema check for global_atomic_fadd_v2f16 builtin (PR #158145)

[via cfe-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-amdgpu

Author: Tim Gu (tcgu-amd)

<details>
<summary>Changes</summary>

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. 

---
Full diff: https://github.com/llvm/llvm-project/pull/158145.diff


2 Files Affected:

- (modified) clang/include/clang/Sema/SemaAMDGPU.h (+2) 
- (modified) clang/lib/Sema/SemaAMDGPU.cpp (+32) 


``````````diff
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

``````````

</details>


https://github.com/llvm/llvm-project/pull/158145