Author: Rana Pratap Reddy Date: 2025-12-08T21:56:35+05:30 New Revision: b32a2f418ac21cba03d73a87b89cd82bf7b318a6
URL: https://github.com/llvm/llvm-project/commit/b32a2f418ac21cba03d73a87b89cd82bf7b318a6 DIFF: https://github.com/llvm/llvm-project/commit/b32a2f418ac21cba03d73a87b89cd82bf7b318a6.diff LOG: [Clang][OpenCL][AMDGPU] Allow _Float16 and half vector type compatibility (#170605) ## 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. Added: clang/test/CodeGen/half-float16-vector-compatibility.cl clang/test/SemaOpenCL/half-float16-vector-compatibility.cl Modified: clang/include/clang/Basic/BuiltinsAMDGPU.def clang/lib/AST/ASTContext.cpp clang/lib/Sema/SemaExpr.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 8af6ce1528a45..0b78b460c0d6a 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -902,75 +902,75 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*", // Image builtins //===----------------------------------------------------------------------===// TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4xiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4xiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4xiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4xifQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts") -TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts") +TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "extended-image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f16_f32, "V4eifQtV4ibii", "nc", "extended-image-insts") TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "extended-image-insts") diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 404ce3ffd77c7..13018ba8de7e7 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..a8c2e39b49923 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(); @@ -9414,6 +9415,12 @@ AssignConvertType Sema::CheckAssignmentConstraints(QualType LHSType, Kind = CK_IntegralToBoolean; return AssignConvertType::Compatible; } + // In OpenCL, allow compatible vector types (e.g. half to _Float16) + if (Context.getLangOpts().OpenCL && + Context.areCompatibleVectorTypes(LHSType, RHSType)) { + Kind = CK_BitCast; + return AssignConvertType::Compatible; + } return AssignConvertType::Incompatible; } if (RHSType->isArithmeticType()) { diff --git a/clang/test/CodeGen/half-float16-vector-compatibility.cl b/clang/test/CodeGen/half-float16-vector-compatibility.cl new file mode 100644 index 0000000000000..a79ee58475fba --- /dev/null +++ b/clang/test/CodeGen/half-float16-vector-compatibility.cl @@ -0,0 +1,36 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 %s -emit-llvm -o - | FileCheck %s + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +typedef int int4 __attribute__((ext_vector_type(4))); +typedef float float4 __attribute__((ext_vector_type(4))); +typedef _Float16 float16_4 __attribute__((ext_vector_type(4))); +typedef half half4 __attribute__((ext_vector_type(4))); + +// CHECK-LABEL: define dso_local noundef <4 x half> @test_assign_half4_to_float16_4( +// CHECK-SAME: <4 x half> noundef returned [[F16_4:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: ret <4 x half> [[F16_4]] +// +half4 test_assign_half4_to_float16_4(float16_4 f16_4) { + return f16_4; +} + +// CHECK-LABEL: define dso_local noundef <4 x half> @test_assign_float16_4_to_half4( +// CHECK-SAME: <4 x half> noundef returned [[H4:%.*]]) local_unnamed_addr #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: ret <4 x half> [[H4]] +// +float16_4 test_assign_float16_4_to_half4(half4 h4) { + return h4; +} + +// CHECK-LABEL: define dso_local noundef <4 x half> @test_float16_4_to_half4( +// CHECK-SAME: <4 x half> noundef returned [[F16_4:%.*]]) local_unnamed_addr #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: ret <4 x half> [[F16_4]] +// +half4 test_float16_4_to_half4(float16_4 f16_4) { + return (half4)f16_4; +} diff --git a/clang/test/SemaOpenCL/half-float16-vector-compatibility.cl b/clang/test/SemaOpenCL/half-float16-vector-compatibility.cl new file mode 100644 index 0000000000000..b1920bddbf181 --- /dev/null +++ b/clang/test/SemaOpenCL/half-float16-vector-compatibility.cl @@ -0,0 +1,75 @@ +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -verify -S -o - %s +// REQUIRES: amdgpu-registered-target + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +typedef int int4 __attribute__((ext_vector_type(4))); +typedef float float4 __attribute__((ext_vector_type(4))); + +typedef half half2 __attribute__((ext_vector_type(2))); +typedef half half3 __attribute__((ext_vector_type(3))); +typedef half half4 __attribute__((ext_vector_type(4))); +typedef half half8 __attribute__((ext_vector_type(8))); +typedef half half16 __attribute__((ext_vector_type(16))); + +typedef _Float16 float16_2 __attribute__((ext_vector_type(2))); +typedef _Float16 float16_3 __attribute__((ext_vector_type(3))); +typedef _Float16 float16_4 __attribute__((ext_vector_type(4))); +typedef _Float16 float16_8 __attribute__((ext_vector_type(8))); +typedef _Float16 float16_16 __attribute__((ext_vector_type(16))); + +void test_half_vector_to_float16(float16_2 f16_2, float16_3 f16_3, float16_4 f16_4, float16_8 f16_8, float16_16 f16_16) { + half2 h2 = f16_2; // expected-no-error + half3 h3 = f16_3; // expected-no-error + half4 h4 = f16_4; // expected-no-error + half8 h8 = f16_8; // expected-no-error + half16 h16 = f16_16; // expected-no-error +} + +void test_float16_vector_to_half(half2 h2, half3 h3, half4 h4, half8 h8, half16 h16) { + float16_2 f16_2 = h2; // expected-no-error + float16_3 f16_3 = h3; // expected-no-error + float16_4 f16_4 = h4; // expected-no-error + float16_8 f16_8 = h8; // expected-no-error + float16_16 f16_16 = h16; // expected-no-error +} + +half4 test_return_half4_from_float16_vector(float16_4 f16_4) { + return f16_4; // expected-no-error +} + +float16_4 test_return_float16_4_from_half4(half4 h4) { + return h4; // expected-no-error +} + +half4 test_explicit_cast_half4_to_float16_vector(half4 h4) { + return (float16_4)h4; // expected-no-error +} + +float16_4 test_explicit_cast_float16_4_to_half4(float16_4 f16_4) { + return (half4)f16_4; // expected-no-error +} + +half4 test_builtin_image_load_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) { + return __builtin_amdgcn_image_load_2d_v4f16_i32(100, i32, i32, tex, 120, 110); // expected-no-error +} + +half4 test_builtin_amdgcn_image_sample_2d_v4f16_f32(half4 v4f16, int i32, float f32, __amdgpu_texture_t tex, int4 vec4i32) { + return __builtin_amdgcn_image_sample_2d_v4f16_f32(100, f32, f32, tex, vec4i32, 0, 120, 110); // expected-no-error +} + +void test_half_mismatch_vector_size_error(float16_2 f16_2, float16_3 f16_3, float16_4 f16_4, float16_8 f16_8, float16_16 f16_16) { + half2 h2 = f16_3 ; // expected-error{{initializing '__private half2' (vector of 2 'half' values) with an expression of incompatible type '__private float16_3' (vector of 3 '_Float16' values)}} + half3 h3 = f16_2; // expected-error{{initializing '__private half3' (vector of 3 'half' values) with an expression of incompatible type '__private float16_2' (vector of 2 '_Float16' values)}} + half4 h4 = f16_8; // expected-error{{initializing '__private half4' (vector of 4 'half' values) with an expression of incompatible type '__private float16_8' (vector of 8 '_Float16' values)}} + half8 h8 = f16_4; // expected-error{{initializing '__private half8' (vector of 8 'half' values) with an expression of incompatible type '__private float16_4' (vector of 4 '_Float16' values)}} + half16 h16 = f16_4; // expected-error{{initializing '__private half16' (vector of 16 'half' values) with an expression of incompatible type '__private float16_4' (vector of 4 '_Float16' values)}} +} + +void test_float16_mismatch_vector_size_error(half2 h2, half3 h3, half4 h4, half8 h8, half16 h16) { + float16_2 f16_2 = h3; // expected-error{{initializing '__private float16_2' (vector of 2 '_Float16' values) with an expression of incompatible type '__private half3' (vector of 3 'half' values)}} + float16_3 f16_3 = h2; // expected-error{{initializing '__private float16_3' (vector of 3 '_Float16' values) with an expression of incompatible type '__private half2' (vector of 2 'half' values)}} + float16_4 f16_4 = h8; // expected-error{{initializing '__private float16_4' (vector of 4 '_Float16' values) with an expression of incompatible type '__private half8' (vector of 8 'half' values)}} + float16_8 f16_8 = h4; // expected-error{{initializing '__private float16_8' (vector of 8 '_Float16' values) with an expression of incompatible type '__private half4' (vector of 4 'half' values)}} + float16_16 f16_16 = h4; // expected-error{{initializing '__private float16_16' (vector of 16 '_Float16' values) with an expression of incompatible type '__private half4' (vector of 4 'half' values)}} +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
