58 lines · plain
1; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV2 3; StorageImageReadWithoutFormat/StorageImageWriteWithoutFormat implicitly4; declare Shader, causing a SPIR-V module to be rejected by the OpenCL5; run-time. See https://github.com/KhronosGroup/SPIRV-Headers/issues/4876; De-facto, OpImageRead and OpImageWrite are allowed to use Unknown Image7; Formats when the Kernel capability is declared. We reflect this behavior8; in the test case, and leave the check under CHECK-SPIRV-NOT to track9; the issue and follow-up its final resolution when ready.10; CHECK-SPIRV-NOT: OpCapability StorageImageReadWithoutFormat11 12; CHECK-SPIRV: %[[#IntTy:]] = OpTypeInt13; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]]14; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat15; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]]16; CHECK-SPIRV: OpImageRead %[[#IVecTy]]17; CHECK-SPIRV: OpImageRead %[[#FVecTy]]18 19;; __kernel void kernelA(__read_only image3d_t input) {20;; uint4 c = read_imageui(input, (int4)(0, 0, 0, 0));21;; }22;;23;; __kernel void kernelB(__read_only image3d_t input) {24;; float4 f = read_imagef(input, (int4)(0, 0, 0, 0));25;; }26 27define dso_local spir_kernel void @kernelA(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {28entry:29 %input.addr = alloca target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), align 830 %c = alloca <4 x i32>, align 1631 %.compoundliteral = alloca <4 x i32>, align 1632 store target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input, ptr %input.addr, align 833 %0 = load target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), ptr %input.addr, align 834 store <4 x i32> zeroinitializer, ptr %.compoundliteral, align 1635 %1 = load <4 x i32>, ptr %.compoundliteral, align 1636 %call = call spir_func <4 x i32> @_Z12read_imageui14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)37 store <4 x i32> %call, ptr %c, align 1638 ret void39}40 41declare spir_func <4 x i32> @_Z12read_imageui14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)42 43define dso_local spir_kernel void @kernelB(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {44entry:45 %input.addr = alloca target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), align 846 %f = alloca <4 x float>, align 1647 %.compoundliteral = alloca <4 x i32>, align 1648 store target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input, ptr %input.addr, align 849 %0 = load target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), ptr %input.addr, align 850 store <4 x i32> zeroinitializer, ptr %.compoundliteral, align 1651 %1 = load <4 x i32>, ptr %.compoundliteral, align 1652 %call = call 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)53 store <4 x float> %call, ptr %f, align 1654 ret void55}56 57declare 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)58