Skip to content

Commit d7a41cf

Browse files
michalpaszkowskiigcbot
michalpaszkowski
authored and
igcbot
committed
Rewrite TargetExtTy retyper using ValueMapTypeRemapper
This change rewrites the TargetExtTy retyper to use the ValueMapTypeRemapper infrastructure, significantly improving the overall design and maintainability of the code. The change also removes unused cases added for additional safety if earlier retyping logic fails. Two additional test cases are added, covering more complex retyping scenarios.
1 parent ba8538b commit d7a41cf

File tree

6 files changed

+407
-559
lines changed

6 files changed

+407
-559
lines changed

‎IGC/Compiler/tests/PreprocessSPVIR/retype-function-declaration.ll‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2025 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; REQUIRES: llvm-16-plus
10+
; RUN: igc_opt --opaque-pointers -igc-preprocess-spvir -S < %s | FileCheck %s
11+
12+
; This test verifies that PreprocessSPVIR retypes TargetExtTys in function
13+
; declarations.
14+
15+
define spir_func void @foo(i64 %in) {
16+
%img = call spir_func target("spirv.SampledImage", void, 1, 0, 0, 0, 0, 0, 0) @_Z90__spirv_ConvertHandleToSampledImageINTEL_RPU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0m(i64 %in)
17+
ret void
18+
}
19+
20+
declare spir_func target("spirv.SampledImage", void, 1, 0, 0, 0, 0, 0, 0) @_Z90__spirv_ConvertHandleToSampledImageINTEL_RPU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0m(i64)
21+
22+
; CHECK: declare spir_func ptr @_Z90__spirv_ConvertHandleToSampledImageINTEL_RPU3AS140__spirv_SampledImage__void_1_0_0_0_0_0_0m(i64)

‎IGC/Compiler/tests/PreprocessSPVIR/retype-targetextty-constant.ll‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2025 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; REQUIRES: llvm-16-plus
10+
; RUN: igc_opt --opaque-pointers -igc-preprocess-spvir -S < %s | FileCheck %s
11+
12+
; This test verifies PreprocessSPVIR retypes TargetExtTy constant
13+
; zeroinitializer arguments to pointer null.
14+
15+
declare spir_func void @UseEvent(target("spirv.Event") %A)
16+
17+
define spir_kernel void @TestKernel() {
18+
entry:
19+
call spir_func void @UseEvent(target("spirv.Event") zeroinitializer)
20+
ret void
21+
}
22+
23+
; CHECK-LABEL: define spir_kernel void @TestKernel(
24+
; CHECK: call spir_func void @UseEvent(
25+
; CHECK-SAME: ptr null
26+
; CHECK-NOT: target("spirv.Event")
27+
; CHECK: ret void

‎IGC/Compiler/tests/PreprocessSPVIR/retype-targetextty-inside-struct-with-gep-use.ll‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2025 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; REQUIRES: llvm-16-plus
10+
; RUN: igc_opt --opaque-pointers -igc-preprocess-spvir -S < %s | FileCheck %s
11+
12+
; This test verifies PreprocessSPVIR retypes TargetExtTy inside structs used by
13+
; GEPs.
14+
15+
; CHECK-NOT: target("spirv.Event")
16+
17+
%"class.sycl::_V1::device_event" = type { target("spirv.Event") }
18+
; CHECK: %"class.sycl::_V1::device_event" = type { ptr }
19+
20+
define internal spir_func void @_ZN4sycl3_V112device_event4waitEv(ptr addrspace(4) align 8 %arg) {
21+
%a = alloca ptr addrspace(4), align 8
22+
%b = addrspacecast ptr %a to ptr addrspace(4)
23+
store ptr addrspace(4) %arg, ptr addrspace(4) %b, align 8
24+
%c = load ptr addrspace(4), ptr addrspace(4) %b, align 8
25+
; CHECK: %d = getelementptr inbounds %"class.sycl::_V1::device_event", ptr addrspace(4) %c, i32 0, i32 0
26+
%d = getelementptr inbounds %"class.sycl::_V1::device_event", ptr addrspace(4) %c, i32 0, i32 0
27+
; CHECK: %e = getelementptr inbounds %"class.sycl::_V1::device_event", ptr addrspace(4) %c, i32 0, i32 0
28+
%e = getelementptr inbounds %"class.sycl::_V1::device_event", ptr addrspace(4) %c, i32 0, i32 0
29+
ret void
30+
}

‎IGC/Compiler/tests/PreprocessSPVIR/retype-targetextty-inside-structs.ll‎

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2025 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; REQUIRES: llvm-16-plus
10+
; RUN: igc_opt --opaque-pointers -igc-preprocess-spvir -S < %s | FileCheck %s
11+
12+
; Verify that the PreprocessSPVIR pass retypes struct members of TargetExtTy
13+
; type as pointers. This is needed so that the following passes can perform
14+
; pointer optimizations correctly and the IR can be linked with builtins module
15+
; coming from Clang (LLVM 16 Clang does not support TargetExtTy).
16+
17+
%struct.SamplerHolder = type { target("spirv.Sampler"), i32 }
18+
%struct.Wrapper = type { %struct.SamplerHolder, i64 }
19+
20+
define spir_kernel void @Kernel(target("spirv.Sampler") %arg) {
21+
entry:
22+
%holder = alloca %struct.SamplerHolder, align 8
23+
%wrapper = alloca %struct.Wrapper, align 8
24+
%holder.sampler.field.gep = getelementptr inbounds %struct.SamplerHolder, ptr %holder, i32 0, i32 0
25+
store target("spirv.Sampler") %arg, ptr %holder.sampler.field.gep, align 8
26+
%wrapper.holder.gep = getelementptr inbounds %struct.Wrapper, ptr %wrapper, i32 0, i32 0
27+
%wrapper.holder.sampler.gep = getelementptr inbounds %struct.SamplerHolder, ptr %wrapper.holder.gep, i32 0, i32 0
28+
%loaded.sampler = load target("spirv.Sampler"), ptr %holder.sampler.field.gep, align 8
29+
store target("spirv.Sampler") %loaded.sampler, ptr %wrapper.holder.sampler.gep, align 8
30+
call spir_func void @Helper(target("spirv.Sampler") %loaded.sampler)
31+
ret void
32+
}
33+
34+
define internal spir_func void @Helper(target("spirv.Sampler") %S) {
35+
entry:
36+
ret void
37+
}
38+
39+
; Check that struct fields are retyped, including nested/wrapped structs.
40+
; CHECK: %struct.SamplerHolder = type { ptr addrspace(2), i32 }
41+
; CHECK: %struct.Wrapper = type { %struct.SamplerHolder, i64 }
42+
43+
; CHECK-LABEL: define spir_kernel void @Kernel(
44+
; CHECK-SAME: ptr addrspace(2) %arg
45+
; CHECK: %holder = alloca %struct.SamplerHolder
46+
; CHECK: %wrapper = alloca %struct.Wrapper
47+
48+
; CHECK-NOT: target("spirv.Sampler")
49+

‎IGC/Compiler/tests/PreprocessSPVIR/retype-targetextty-sret.ll‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2025 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; REQUIRES: llvm-16-plus
10+
; RUN: igc_opt --opaque-pointers -igc-preprocess-spvir -S < %s | FileCheck %s
11+
12+
; This test verifies PreprocessSPVIR retypes TargetExtTy inside structs used in
13+
; function argument sret attributes.
14+
15+
%union.anon = type { ptr addrspace(1) }
16+
%"class.sycl::_V1::multi_ptr" = type { ptr addrspace(3) }
17+
%"class.sycl::_V1::device_event" = type { target("spirv.Event") }
18+
19+
; CHECK-NOT: target("spirv.Event")
20+
; CHECK: %"class.sycl::_V1::device_event" = type { ptr }
21+
22+
define void @f(ptr addrspace(4) noalias sret(%"class.sycl::_V1::device_event") align 8 %arg1, ptr addrspace(4) align 1 %arg2, ptr byval(%"class.sycl::_V1::multi_ptr") align 8 %arg3, ptr byval(%union.anon) align 8 %arg4, i64 %arg5) {
23+
call spir_func void @g(ptr addrspace(4) noalias sret(%"class.sycl::_V1::device_event") align 8 %arg1, ptr addrspace(4) align 1 %arg2, ptr byval(%"class.sycl::_V1::multi_ptr") align 8 %arg3, ptr byval(%union.anon) align 8 %arg4, i64 1, i64 1)
24+
; CHECK: call spir_func void @g(ptr addrspace(4) noalias sret(%"class.sycl::_V1::device_event") align 8 %arg1, ptr addrspace(4) align 1 %arg2, ptr byval(%"class.sycl::_V1::multi_ptr") align 8 %arg3, ptr byval(%union.anon) align 8 %arg4, i64 1, i64 1)
25+
ret void
26+
}
27+
28+
define void @g(ptr addrspace(4) noalias sret(%"class.sycl::_V1::device_event") align 8 %arg1, ptr addrspace(4) align 1 %arg2, ptr byval(%"class.sycl::_V1::multi_ptr") align 8 %arg3, ptr byval(%union.anon) align 8 %arg4, i64 %arg5, i64 %arg6) {
29+
ret void
30+
}

0 commit comments

Comments
 (0)