Skip to content

Commit d27b674

Browse files
jcranmer-intelsys-ce-bb
jcranmer-intel
authored and
sys-ce-bb
committed
Use -emit-opaque-pointers for SPIR-V-friendly IR.
There are a couple of tests that use some magic to reuse the checks between opencl-flavored IR and SPIR-V-friendly IR; these tests are not yet ported to make diffs smaller. Original commit: KhronosGroup/SPIRV-LLVM-Translator@60ac579
1 parent 313ad0a commit d27b674

File tree

12 files changed

+101
-146
lines changed

12 files changed

+101
-146
lines changed

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_blocking_pipes/PipeBlocking.ll

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6,8 +6,8 @@
66
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc
77
; RUN: llvm-dis -opaque-pointers=0 < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

9-
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc --spirv-target-env=SPV-IR
10-
; RUN: llvm-dis -opaque-pointers=0 < %t.bc | FileCheck %s --check-prefix=CHECK-SPV-IR
9+
; RUN: llvm-spirv -r %t.spv -o %t.bc --spirv-target-env=SPV-IR
10+
; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-SPV-IR
1111

1212
; ModuleID = 'test/CodeGenOpenCL/pipe_builtin.cl'
1313
source_filename = "test/CodeGenOpenCL/pipe_builtin.cl"
@@ -40,8 +40,7 @@ target triple = "spir64-unknown-unknown"
4040

4141
; CHECK-LLVM: call spir_func void @__write_pipe_2_bl(%opencl.pipe_wo_t addrspace(1)*{{.*}}, i8 addrspace(4)* %{{[0-9]+}}, i32 2, i32 2)
4242

43-
; CHECK-SPV-IR: %spirv.Pipe._0 = type opaque
44-
; CHECK-SPV-IR: call spir_func void @_Z30__spirv_WritePipeBlockingINTEL{{.*}}(%spirv.Pipe._1 addrspace(1)*{{.*}}, i9 addrspace(4)*{{.*}}, i32 2, i32 2)
43+
; CHECK-SPV-IR: call spir_func void @_Z30__spirv_WritePipeBlockingINTEL{{.*}}(target("spirv.Pipe", 1){{.*}}, ptr addrspace(4){{.*}}, i32 2, i32 2)
4544

4645
; Function Attrs: convergent noinline nounwind optnone
4746
define spir_func void @foo(%opencl.pipe_ro_t addrspace(1)* %p, i32 addrspace(1)* %ptr) #0 {

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_media_block_io/SPV_INTEL_media_block_io.cl

Lines changed: 30 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -4,8 +4,8 @@
44
// RUN: spirv-val %t.spv
55
// RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
66
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
7-
// RUN: llvm-spirv -r -emit-opaque-pointers=0 --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
8-
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR
7+
// RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
8+
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR
99

1010
uchar __attribute__((overloadable)) intel_sub_group_media_block_read_uc(int2 src_offset, int width, int height, read_only image2d_t image);
1111
uchar2 __attribute__((overloadable)) intel_sub_group_media_block_read_uc2(int2 src_offset, int width, int height, read_only image2d_t image);
@@ -215,31 +215,31 @@ __kernel void intel_media_block_test(int2 edgeCoord, __read_only image2d_t src_l
215215
// CHECK-LLVM: call spir_func void @_Z37intel_sub_group_media_block_write_ui4Dv2_iiiDv4_j14ocl_image2d_wo(<2 x i32> %edgeCoord, i32 1, i32 16, <4 x i32> %{{.*}}, %opencl.image2d_wo_t addrspace(1)* %dst_luma_image)
216216
// CHECK-LLVM: call spir_func void @_Z37intel_sub_group_media_block_write_ui8Dv2_iiiDv8_j14ocl_image2d_wo(<2 x i32> %edgeCoord, i32 1, i32 16, <8 x i32> %{{.*}}, %opencl.image2d_wo_t addrspace(1)* %dst_luma_image)
217217

218-
// CHECK-SPV-IR: call spir_func i8 @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_RcharPU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
219-
// CHECK-SPV-IR: call spir_func <2 x i8> @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_Rchar2PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
220-
// CHECK-SPV-IR: call spir_func <4 x i8> @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_Rchar4PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
221-
// CHECK-SPV-IR: call spir_func <8 x i8> @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_Rchar8PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
222-
// CHECK-SPV-IR: call spir_func <16 x i8> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rchar16PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
223-
// CHECK-SPV-IR: call spir_func i16 @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_RshortPU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
224-
// CHECK-SPV-IR: call spir_func <2 x i16> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rshort2PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
225-
// CHECK-SPV-IR: call spir_func <4 x i16> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rshort4PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
226-
// CHECK-SPV-IR: call spir_func <8 x i16> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rshort8PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
227-
// CHECK-SPV-IR: call spir_func <16 x i16> @_Z49__spirv_SubgroupImageMediaBlockReadINTEL_Rshort16PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
228-
// CHECK-SPV-IR: call spir_func i32 @_Z45__spirv_SubgroupImageMediaBlockReadINTEL_RintPU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
229-
// CHECK-SPV-IR: call spir_func <2 x i32> @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_Rint2PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
230-
// CHECK-SPV-IR: call spir_func <4 x i32> @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_Rint4PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
231-
// CHECK-SPV-IR: call spir_func <8 x i32> @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_Rint8PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
232-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiic(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, i8 %{{.*}})
233-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv2_c(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <2 x i8> %{{.*}})
234-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv4_c(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <4 x i8> %{{.*}})
235-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv8_c(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <8 x i8> %{{.*}})
236-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv16_c(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <16 x i8> %{{.*}})
237-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiis(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, i16 %{{.*}})
238-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv2_s(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <2 x i16> %{{.*}})
239-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv4_s(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <4 x i16> %{{.*}})
240-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv8_s(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <8 x i16> %{{.*}})
241-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv16_s(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <16 x i16> %{{.*}})
242-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiii(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, i32 %{{.*}})
243-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiS2_(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <2 x i32> %{{.*}})
244-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv4_i(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <4 x i32> %{{.*}})
245-
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv8_i(%spirv.Image._void_1_0_0_0_0_0_1 addrspace(1)* %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <8 x i32> %{{.*}})
218+
// CHECK-SPV-IR: call spir_func i8 @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_RcharPU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
219+
// CHECK-SPV-IR: call spir_func <2 x i8> @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_Rchar2PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
220+
// CHECK-SPV-IR: call spir_func <4 x i8> @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_Rchar4PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
221+
// CHECK-SPV-IR: call spir_func <8 x i8> @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_Rchar8PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
222+
// CHECK-SPV-IR: call spir_func <16 x i8> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rchar16PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
223+
// CHECK-SPV-IR: call spir_func i16 @_Z47__spirv_SubgroupImageMediaBlockReadINTEL_RshortPU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
224+
// CHECK-SPV-IR: call spir_func <2 x i16> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rshort2PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
225+
// CHECK-SPV-IR: call spir_func <4 x i16> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rshort4PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
226+
// CHECK-SPV-IR: call spir_func <8 x i16> @_Z48__spirv_SubgroupImageMediaBlockReadINTEL_Rshort8PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
227+
// CHECK-SPV-IR: call spir_func <16 x i16> @_Z49__spirv_SubgroupImageMediaBlockReadINTEL_Rshort16PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
228+
// CHECK-SPV-IR: call spir_func i32 @_Z45__spirv_SubgroupImageMediaBlockReadINTEL_RintPU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
229+
// CHECK-SPV-IR: call spir_func <2 x i32> @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_Rint2PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
230+
// CHECK-SPV-IR: call spir_func <4 x i32> @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_Rint4PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
231+
// CHECK-SPV-IR: call spir_func <8 x i32> @_Z46__spirv_SubgroupImageMediaBlockReadINTEL_Rint8PU3AS133__spirv_Image__void_1_0_0_0_0_0_0Dv2_iii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16)
232+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiic(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, i8 %{{.*}})
233+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv2_c(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <2 x i8> %{{.*}})
234+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv4_c(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <4 x i8> %{{.*}})
235+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv8_c(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <8 x i8> %{{.*}})
236+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv16_c(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <16 x i8> %{{.*}})
237+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiis(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, i16 %{{.*}})
238+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv2_s(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <2 x i16> %{{.*}})
239+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv4_s(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <4 x i16> %{{.*}})
240+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv8_s(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <8 x i16> %{{.*}})
241+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv16_s(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <16 x i16> %{{.*}})
242+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiii(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, i32 %{{.*}})
243+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiS2_(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <2 x i32> %{{.*}})
244+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <4 x i32> %{{.*}})
245+
// CHECK-SPV-IR: call spir_func void @_Z41__spirv_SubgroupImageMediaBlockWriteINTELPU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iiiDv8_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %dst_luma_image, <2 x i32> %edgeCoord, i32 1, i32 16, <8 x i32> %{{.*}})

0 commit comments

Comments
 (0)