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

[via llvm-commits]

Author: Arseniy Obolenskiy
Date: 2026-09-14T12:07:49+02:00
New Revision: 37fe7024f2944f71d58b8608d209bda3254fa43a

URL: https://github.com/llvm/llvm-project/commit/37fe7024f2944f71d58b8608d209bda3254fa43a
DIFF: https://github.com/llvm/llvm-project/commit/37fe7024f2944f71d58b8608d209bda3254fa43a.diff

LOG: [SPIR-V] Add missing read_imageh builtin mapping (#222549)

Added: 
    

Modified: 
    llvm/lib/Target/SPIRV/SPIRVBuiltins.td
    llvm/test/CodeGen/SPIRV/read_image.ll

Removed: 
    


################################################################################
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..22613b6a0612b 100644
--- a/llvm/test/CodeGen/SPIRV/read_image.ll
+++ b/llvm/test/CodeGen/SPIRV/read_image.ll
@@ -9,13 +9,18 @@
 ; in the test case, and leave the check under CHECK-SPIRV-NOT to track
 ; the issue and follow-up its final resolution when ready.
 ; CHECK-SPIRV-NOT: OpCapability StorageImageReadWithoutFormat
+; CHECK-SPIRV: OpCapability Float16
+; CHECK-SPIRV-NOT: OpCapability StorageImageReadWithoutFormat
 
 ; CHECK-SPIRV: %[[#IntTy:]] = OpTypeInt
-; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]]
-; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat
-; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]]
+; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]] 4
+; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]] 4
+; CHECK-SPIRV: %[[#HalfTy:]] = OpTypeFloat 16
+; CHECK-SPIRV: %[[#HVecTy:]] = OpTypeVector %[[#HalfTy]] 4
 ; 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 +29,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 +65,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)