Skip to content

Commit b32a2f4

Browse files
ranapratap55ranapratap55
authored
[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.
1 parent fccb65e commit b32a2f4

File tree

5 files changed

+164
-31
lines changed

5 files changed

+164
-31
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 30 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -902,75 +902,75 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*",
902902
// Image builtins
903903
//===----------------------------------------------------------------------===//
904904
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiQtii", "nc", "image-insts")
905-
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiQtii", "nc", "image-insts")
905+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4xiiQtii", "nc", "image-insts")
906906
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
907-
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
907+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
908908
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiQtii", "nc", "image-insts")
909909
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
910-
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
910+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
911911
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiQtii", "nc", "image-insts")
912912
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
913-
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
913+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
914914
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
915-
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
915+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
916916
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
917-
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
917+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
918918
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
919-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
919+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
920920
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
921-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
921+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
922922
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiQtii", "nc", "image-insts")
923923
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
924-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
924+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
925925
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiQtii", "nc", "image-insts")
926926
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
927-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
927+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
928928
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
929-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
929+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
930930
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
931-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
931+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
932932
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiQtii", "nc", "image-insts")
933-
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiQtii", "nc", "image-insts")
933+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4xiiQtii", "nc", "image-insts")
934934
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
935-
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
935+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
936936
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiQtii", "nc", "image-insts")
937937
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
938-
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
938+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
939939
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiQtii", "nc", "image-insts")
940940
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
941-
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
941+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
942942
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
943-
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
943+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
944944
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
945-
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
945+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
946946
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
947-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
947+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
948948
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
949-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
949+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
950950
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiQtii", "nc", "image-insts")
951951
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
952-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
952+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
953953
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiQtii", "nc", "image-insts")
954954
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
955-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
955+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
956956
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
957-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
957+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
958958
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
959-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
959+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
960960
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "image-insts")
961-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifQtV4ibii", "nc", "image-insts")
961+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4xifQtV4ibii", "nc", "image-insts")
962962
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
963-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
963+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts")
964964
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffQtV4ibii", "nc", "image-insts")
965965
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
966-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
966+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts")
967967
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffQtV4ibii", "nc", "image-insts")
968968
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
969-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
969+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
970970
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
971-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
971+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
972972
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
973-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
973+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
974974
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "extended-image-insts")
975975
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f16_f32, "V4eifQtV4ibii", "nc", "extended-image-insts")
976976
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "extended-image-insts")

‎clang/lib/AST/ASTContext.cpp‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10527,6 +10527,21 @@ bool ASTContext::areCompatibleVectorTypes(QualType FirstVec,
1052710527
Second->getVectorKind() != VectorKind::RVVFixedLengthMask_4)
1052810528
return true;
1052910529

10530+
// In OpenCL, treat half and _Float16 vector types as compatible.
10531+
if (getLangOpts().OpenCL &&
10532+
First->getNumElements() == Second->getNumElements()) {
10533+
QualType FirstElt = First->getElementType();
10534+
QualType SecondElt = Second->getElementType();
10535+
10536+
if ((FirstElt->isFloat16Type() && SecondElt->isHalfType()) ||
10537+
(FirstElt->isHalfType() && SecondElt->isFloat16Type())) {
10538+
if (First->getVectorKind() != VectorKind::AltiVecPixel &&
10539+
First->getVectorKind() != VectorKind::AltiVecBool &&
10540+
Second->getVectorKind() != VectorKind::AltiVecPixel &&
10541+
Second->getVectorKind() != VectorKind::AltiVecBool)
10542+
return true;
10543+
}
10544+
}
1053010545
return false;
1053110546
}
1053210547

‎clang/lib/Sema/SemaExpr.cpp‎

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7819,7 +7819,8 @@ ExprResult Sema::CheckExtVectorCast(SourceRange R, QualType DestTy,
78197819
if (SrcTy->isVectorType()) {
78207820
if (!areLaxCompatibleVectorTypes(SrcTy, DestTy) ||
78217821
(getLangOpts().OpenCL &&
7822-
!Context.hasSameUnqualifiedType(DestTy, SrcTy))) {
7822+
!Context.hasSameUnqualifiedType(DestTy, SrcTy) &&
7823+
!Context.areCompatibleVectorTypes(DestTy, SrcTy))) {
78237824
Diag(R.getBegin(),diag::err_invalid_conversion_between_ext_vectors)
78247825
<< DestTy << SrcTy << R;
78257826
return ExprError();
@@ -9414,6 +9415,12 @@ AssignConvertType Sema::CheckAssignmentConstraints(QualType LHSType,
94149415
Kind = CK_IntegralToBoolean;
94159416
return AssignConvertType::Compatible;
94169417
}
9418+
// In OpenCL, allow compatible vector types (e.g. half to _Float16)
9419+
if (Context.getLangOpts().OpenCL &&
9420+
Context.areCompatibleVectorTypes(LHSType, RHSType)) {
9421+
Kind = CK_BitCast;
9422+
return AssignConvertType::Compatible;
9423+
}
94179424
return AssignConvertType::Incompatible;
94189425
}
94199426
if (RHSType->isArithmeticType()) {

‎clang/test/CodeGen/half-float16-vector-compatibility.cl‎

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2+
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 %s -emit-llvm -o - | FileCheck %s
3+
4+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
5+
6+
typedef int int4 __attribute__((ext_vector_type(4)));
7+
typedef float float4 __attribute__((ext_vector_type(4)));
8+
typedef _Float16 float16_4 __attribute__((ext_vector_type(4)));
9+
typedef half half4 __attribute__((ext_vector_type(4)));
10+
11+
// CHECK-LABEL: define dso_local noundef <4 x half> @test_assign_half4_to_float16_4(
12+
// CHECK-SAME: <4 x half> noundef returned [[F16_4:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
13+
// CHECK-NEXT: [[ENTRY:.*:]]
14+
// CHECK-NEXT: ret <4 x half> [[F16_4]]
15+
//
16+
half4 test_assign_half4_to_float16_4(float16_4 f16_4) {
17+
return f16_4;
18+
}
19+
20+
// CHECK-LABEL: define dso_local noundef <4 x half> @test_assign_float16_4_to_half4(
21+
// CHECK-SAME: <4 x half> noundef returned [[H4:%.*]]) local_unnamed_addr #[[ATTR0]] {
22+
// CHECK-NEXT: [[ENTRY:.*:]]
23+
// CHECK-NEXT: ret <4 x half> [[H4]]
24+
//
25+
float16_4 test_assign_float16_4_to_half4(half4 h4) {
26+
return h4;
27+
}
28+
29+
// CHECK-LABEL: define dso_local noundef <4 x half> @test_float16_4_to_half4(
30+
// CHECK-SAME: <4 x half> noundef returned [[F16_4:%.*]]) local_unnamed_addr #[[ATTR0]] {
31+
// CHECK-NEXT: [[ENTRY:.*:]]
32+
// CHECK-NEXT: ret <4 x half> [[F16_4]]
33+
//
34+
half4 test_float16_4_to_half4(float16_4 f16_4) {
35+
return (half4)f16_4;
36+
}

‎clang/test/SemaOpenCL/half-float16-vector-compatibility.cl‎

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -verify -S -o - %s
2+
// REQUIRES: amdgpu-registered-target
3+
4+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
5+
6+
typedef int int4 __attribute__((ext_vector_type(4)));
7+
typedef float float4 __attribute__((ext_vector_type(4)));
8+
9+
typedef half half2 __attribute__((ext_vector_type(2)));
10+
typedef half half3 __attribute__((ext_vector_type(3)));
11+
typedef half half4 __attribute__((ext_vector_type(4)));
12+
typedef half half8 __attribute__((ext_vector_type(8)));
13+
typedef half half16 __attribute__((ext_vector_type(16)));
14+
15+
typedef _Float16 float16_2 __attribute__((ext_vector_type(2)));
16+
typedef _Float16 float16_3 __attribute__((ext_vector_type(3)));
17+
typedef _Float16 float16_4 __attribute__((ext_vector_type(4)));
18+
typedef _Float16 float16_8 __attribute__((ext_vector_type(8)));
19+
typedef _Float16 float16_16 __attribute__((ext_vector_type(16)));
20+
21+
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) {
22+
half2 h2 = f16_2; // expected-no-error
23+
half3 h3 = f16_3; // expected-no-error
24+
half4 h4 = f16_4; // expected-no-error
25+
half8 h8 = f16_8; // expected-no-error
26+
half16 h16 = f16_16; // expected-no-error
27+
}
28+
29+
void test_float16_vector_to_half(half2 h2, half3 h3, half4 h4, half8 h8, half16 h16) {
30+
float16_2 f16_2 = h2; // expected-no-error
31+
float16_3 f16_3 = h3; // expected-no-error
32+
float16_4 f16_4 = h4; // expected-no-error
33+
float16_8 f16_8 = h8; // expected-no-error
34+
float16_16 f16_16 = h16; // expected-no-error
35+
}
36+
37+
half4 test_return_half4_from_float16_vector(float16_4 f16_4) {
38+
return f16_4; // expected-no-error
39+
}
40+
41+
float16_4 test_return_float16_4_from_half4(half4 h4) {
42+
return h4; // expected-no-error
43+
}
44+
45+
half4 test_explicit_cast_half4_to_float16_vector(half4 h4) {
46+
return (float16_4)h4; // expected-no-error
47+
}
48+
49+
float16_4 test_explicit_cast_float16_4_to_half4(float16_4 f16_4) {
50+
return (half4)f16_4; // expected-no-error
51+
}
52+
53+
half4 test_builtin_image_load_2d_2(half4 v4f16, int i32, __amdgpu_texture_t tex) {
54+
return __builtin_amdgcn_image_load_2d_v4f16_i32(100, i32, i32, tex, 120, 110); // expected-no-error
55+
}
56+
57+
half4 test_builtin_amdgcn_image_sample_2d_v4f16_f32(half4 v4f16, int i32, float f32, __amdgpu_texture_t tex, int4 vec4i32) {
58+
return __builtin_amdgcn_image_sample_2d_v4f16_f32(100, f32, f32, tex, vec4i32, 0, 120, 110); // expected-no-error
59+
}
60+
61+
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) {
62+
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)}}
63+
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)}}
64+
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)}}
65+
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)}}
66+
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)}}
67+
}
68+
69+
void test_float16_mismatch_vector_size_error(half2 h2, half3 h3, half4 h4, half8 h8, half16 h16) {
70+
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)}}
71+
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)}}
72+
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)}}
73+
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)}}
74+
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)}}
75+
}

0 commit comments

Comments
 (0)