[Clang][OpenCL][AMDGPU] Allow _Float16 and half vector type compatibility

[ranapratap55]

Summary

Allowing implicit compatibility between _Float16 vector types and half vector types in OpenCL mode. This enables AMDGPU builtins to work correctly across OpenCL, HIP, and C++ without requiring separate builtin definitions.

Problem Statement

When using AMDGPU image builtins that return half-precision vectors in OpenCL, users encounter type incompatibility errors:
Builtin Definition:
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4xiiQtii", "nc", "image-insts")

Test Case:

typedef half half4 __attribute__((ext_vector_type(4)));
half4 test_builtin_image_load_1d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {
  return __builtin_amdgcn_image_load_1d_v4f16_i32(100, i32, tex, 120, i32);
}

Error:

error: returning '__attribute__((__vector_size__(4 * sizeof(_Float16)))) _Float16' 
(vector of 4 '_Float16' values) from a function with incompatible result type 
'half4' (vector of 4 'half' values)

Solution

In OpenCL, allow implicit compatibility between _Float16 vector types and half vector types. This is needed for AMDGPU builtins that may return _Float16 vectors to work correctly with OpenCL half vector types.

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang

Author: Rana Pratap Reddy (ranapratap55)

Changes

In OpenCL, allow implicit compatibility between _Float16 vector types and half vector types. This is needed for AMDGPU builtins that may return _Float16 vectors to work correctly with OpenCL half vector types.

---

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

2 Files Affected:

• (modified) clang/lib/AST/ASTContext.cpp (+15)
• (modified) clang/lib/Sema/SemaExpr.cpp (+2-1)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index b359fc8350375..7d000f8a8764f 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -10527,6 +10527,21 @@ bool ASTContext::areCompatibleVectorTypes(QualType FirstVec,
       Second->getVectorKind() != VectorKind::RVVFixedLengthMask_4)
     return true;
 
+  // In OpenCL, treat half and _Float16 vector types as compatible.
+  if (getLangOpts().OpenCL &&
+      First->getNumElements() == Second->getNumElements()) {
+    QualType FirstElt = First->getElementType();
+    QualType SecondElt = Second->getElementType();
+
+    if ((FirstElt->isFloat16Type() && SecondElt->isHalfType()) ||
+        (FirstElt->isHalfType() && SecondElt->isFloat16Type())) {
+      if (First->getVectorKind() != VectorKind::AltiVecPixel &&
+          First->getVectorKind() != VectorKind::AltiVecBool &&
+          Second->getVectorKind() != VectorKind::AltiVecPixel &&
+          Second->getVectorKind() != VectorKind::AltiVecBool)
+        return true;
+    }
+  }
   return false;
 }
 
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index cfabd1b76c103..741bcb7e41db2 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -7819,7 +7819,8 @@ ExprResult Sema::CheckExtVectorCast(SourceRange R, QualType DestTy,
   if (SrcTy->isVectorType()) {
     if (!areLaxCompatibleVectorTypes(SrcTy, DestTy) ||
         (getLangOpts().OpenCL &&
-         !Context.hasSameUnqualifiedType(DestTy, SrcTy))) {
+         !Context.hasSameUnqualifiedType(DestTy, SrcTy) &&
+         !Context.areCompatibleVectorTypes(DestTy, SrcTy))) {
       Diag(R.getBegin(),diag::err_invalid_conversion_between_ext_vectors)
         << DestTy << SrcTy << R;
       return ExprError();

[shiltian]

We will need more folks (preferably in OpenCL language committee and/or outside AMD) to review.

Can you also add more context in the description?

[arsenm]

Missing tests?

[github-actions]

✅ With the latest revision this PR passed the C/C++ code formatter.

[ranapratap55]

Missing tests?

Added tests.

[wenju-he]

LGTM