[llvm] [SPIR-V] Add missing read_imageh builtin mapping (PR #222549)

[Arseniy Obolenskiy via llvm-commits]

https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/222549

None

>From 3958e67dd76dde2638792b25b5c96d15847259c7 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 10 Sep 2026 10:22:52 +0200
Subject: [PATCH] [SPIR-V] Add missing read_imageh builtin mapping

---
 llvm/lib/Target/SPIRV/SPIRVBuiltins.td |  1 +
 llvm/test/CodeGen/SPIRV/read_image.ll  | 25 ++++++++++++++++++++++++-
 2 files changed, 25 insertions(+), 1 deletion(-)

diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
index a382c69f47593..0f4bd12313c2c 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td
@@ -122,6 +122,7 @@ def : DemangledBuiltin<"__spirv_DotAccSat", OpenCL_std, IntegerDot, 3, 3>;
 def : DemangledBuiltin<"read_imagei", OpenCL_std, ReadImage, 2, 4>;
 def : DemangledBuiltin<"read_imageui", OpenCL_std, ReadImage, 2, 4>;
 def : DemangledBuiltin<"read_imagef", OpenCL_std, ReadImage, 2, 4>;
+def : DemangledBuiltin<"read_imageh", OpenCL_std, ReadImage, 2, 4>;
 def : DemangledBuiltin<"__spirv_ImageRead", OpenCL_std, ReadImage, 2, 0>;
 
 def : DemangledBuiltin<"write_imagef", OpenCL_std, WriteImage, 3, 4>;
diff --git a/llvm/test/CodeGen/SPIRV/read_image.ll b/llvm/test/CodeGen/SPIRV/read_image.ll
index 8fe8ac2a9628a..b53e0ceec8b60 100644
--- a/llvm/test/CodeGen/SPIRV/read_image.ll
+++ b/llvm/test/CodeGen/SPIRV/read_image.ll
@@ -12,10 +12,13 @@
 
 ; CHECK-SPIRV: %[[#IntTy:]] = OpTypeInt
 ; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]]
-; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat
+; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat 32
 ; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]]
+; CHECK-SPIRV: %[[#HalfTy:]] = OpTypeFloat 16
+; CHECK-SPIRV: %[[#HVecTy:]] = OpTypeVector %[[#HalfTy]]
 ; CHECK-SPIRV: OpImageRead %[[#IVecTy]]
 ; CHECK-SPIRV: OpImageRead %[[#FVecTy]]
+; CHECK-SPIRV: OpImageRead %[[#HVecTy]]
 
 ;; __kernel void kernelA(__read_only image3d_t input) {
 ;;   uint4 c = read_imageui(input, (int4)(0, 0, 0, 0));
@@ -24,6 +27,10 @@
 ;; __kernel void kernelB(__read_only image3d_t input) {
 ;;   float4 f = read_imagef(input, (int4)(0, 0, 0, 0));
 ;; }
+;;
+;; __kernel void kernelC(__read_only image3d_t input) {
+;;   half4 h = read_imageh(input, (int4)(0, 0, 0, 0));
+;; }
 
 define dso_local spir_kernel void @kernelA(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {
 entry:
@@ -56,3 +63,19 @@ entry:
 }
 
 declare spir_func <4 x float> @_Z11read_imagef14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)
+
+define dso_local spir_kernel void @kernelC(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {
+entry:
+  %input.addr = alloca target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), align 8
+  %h = alloca <4 x half>, align 8
+  %.compoundliteral = alloca <4 x i32>, align 16
+  store target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input, ptr %input.addr, align 8
+  %0 = load target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), ptr %input.addr, align 8
+  store <4 x i32> zeroinitializer, ptr %.compoundliteral, align 16
+  %1 = load <4 x i32>, ptr %.compoundliteral, align 16
+  %call = call spir_func <4 x half> @_Z11read_imageh14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)
+  store <4 x half> %call, ptr %h, align 8
+  ret void
+}
+
+declare spir_func <4 x half> @_Z11read_imageh14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)