Skip to content

Commit 2c12e65

Browse files
smilczekigcbot
smilczek
authored and
igcbot
committed
Treat geps with indices zero as bitcast in value tracker
Bitcasting a struct pointer to zero'th field of the struct is essentially the same as doing a gep with zero as every indice. Because of that the value tracker becomes confused when a kernel utilizes this property by using gepping for zero'th elt and bitcasting for them in other places. This commit fixes the issue by making the value tracker treat such geps as bitcasts.
1 parent f45f621 commit 2c12e65

File tree

3 files changed

+208
-2
lines changed

3 files changed

+208
-2
lines changed

‎IGC/Compiler/Optimizer/ValueTracker.cpp‎

Lines changed: 16 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -268,6 +268,13 @@ Value* ValueTracker::handleConstExpr(ConstantExpr* CE)
268268
}
269269
}
270270

271+
// returns true if all GEP indices are 0 and constant.
272+
static bool GEP_is_bitcast(const GetElementPtrInst* GEP)
273+
{
274+
// if all indices are zero, the gep is essentially a bitcast.
275+
return GEP->hasAllZeroIndices();
276+
}
277+
271278
// This function represents the second step of the overall algorithm. It goes
272279
// down through the tree and looks for the value stored in alloca. In most cases
273280
// it returns the final value (image, sampler or constant). For more complex cases,
@@ -291,6 +298,11 @@ Value* ValueTracker::findAllocaValue(Value* V, const uint depth)
291298
continue;
292299
}
293300

301+
if (GEP_is_bitcast(GEP)) {
302+
if (auto leaf = findAllocaValue(GEP, depth); isValidLeaf(leaf))
303+
return leaf;
304+
}
305+
294306
unsigned numIndices = GEP->getNumIndices();
295307
if (numIndices > depth + 1)
296308
continue;
@@ -455,8 +467,10 @@ Value* ValueTracker::trackValue(Value* I)
455467
}
456468
else if (auto* I = dyn_cast<GetElementPtrInst>(baseValue))
457469
{
458-
for (unsigned int i = I->getNumIndices(); i > 1; --i)
459-
gepIndices.push_back(cast<ConstantInt>(I->getOperand(i)));
470+
if (!GEP_is_bitcast(I)) {
471+
for (unsigned int i = I->getNumIndices(); i > 1; --i)
472+
gepIndices.push_back(cast<ConstantInt>(I->getOperand(i)));
473+
}
460474

461475
baseValue = I->getOperand(0);
462476
}

‎IGC/Compiler/tests/OCLBIConverter/value-tracker-find-alloca-val-indirect-direct-bitcast-reverse.ll‎

Lines changed: 96 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,96 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2024 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; Check kernel argument is successfully tracked when store address operand is
10+
; reached through other operations than load address operand. Such a thing is
11+
; possible because getelementptr instructions with all zero indices act like
12+
; bitcasts.
13+
14+
; RUN: igc_opt --opaque-pointers -igc-conv-ocl-to-common -S < %s -o - | FileCheck %s
15+
16+
; CHECK-NOT: assertion failed
17+
; CHECK: addrspacecast ptr addrspace(1) %a to ptr addrspace(393218)
18+
19+
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-n8:16:32"
20+
target triple = "spir64-unknown-unknown"
21+
22+
%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
23+
%"class.sycl::_V1::detail::array" = type { [3 x i64] }
24+
%spirv.Image._void_2_0_0_0_0_0_0 = type opaque
25+
%spirv.Sampler = type opaque
26+
%"class.sycl::_V1::detail::RoundedRangeIDGenerator" = type <{ %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", i8, [7 x i8] }>
27+
%"class.sycl::_V1::detail::RoundedRangeKernel" = type { %"class.sycl::_V1::range", %class.accessor_sampler }
28+
%class.accessor_sampler = type { %"class.sycl::_V1::accessor", %"class.sycl::_V1::sampler" }
29+
%"class.sycl::_V1::accessor" = type { %"class.sycl::_V1::detail::image_accessor" }
30+
%"class.sycl::_V1::detail::image_accessor" = type { %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, [24 x i8] }
31+
%"class.sycl::_V1::sampler" = type { %"class.sycl::_V1::detail::sampler_impl", [8 x i8] }
32+
%"class.sycl::_V1::detail::sampler_impl" = type { %spirv.Sampler addrspace(2)* }
33+
%spirv.SampledImage._void_2_0_0_0_0_0_0 = type opaque
34+
35+
; Function Attrs: convergent nounwind
36+
define spir_kernel void @indirect_vs_direct_bc(%spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)* %a, %spirv.Sampler addrspace(2)* %b) #0 {
37+
entry:
38+
%alloc = alloca %"class.sycl::_V1::detail::RoundedRangeKernel", align 8
39+
%store_addr = getelementptr inbounds %"class.sycl::_V1::detail::RoundedRangeKernel", %"class.sycl::_V1::detail::RoundedRangeKernel"* %alloc, i64 0, i32 1
40+
br label %bitcast_with_bc_and_gep
41+
42+
bitcast_with_bc_and_gep:
43+
%bc_store_addr_direct = bitcast %class.accessor_sampler* %store_addr to %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)**
44+
store %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)* %a, %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)** %bc_store_addr_direct, align 8
45+
br label %bitcast_straight_to_desired_type
46+
47+
bitcast_straight_to_desired_type:
48+
%bc_store_addr = bitcast %class.accessor_sampler* %store_addr to %"class.sycl::_V1::detail::image_accessor"*
49+
%gep_as_bc = getelementptr inbounds %"class.sycl::_V1::detail::image_accessor", %"class.sycl::_V1::detail::image_accessor"* %bc_store_addr, i64 0, i32 0
50+
%load_img_ptr = load %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)** %gep_as_bc, align 8
51+
br label %exit
52+
53+
exit:
54+
%to_i64 = ptrtoint %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)* %load_img_ptr to i64
55+
%arg0 = trunc i64 %to_i64 to i32
56+
%sampler = ptrtoint %spirv.Sampler addrspace(2)* %b to i64
57+
%or_1 = or i64 %sampler, 1
58+
%arg1 = trunc i64 %or_1 to i32
59+
%bi_call = call spir_func <4 x float> @__builtin_IB_OCL_3d_sample_l(i32 noundef %arg0, i32 noundef %arg1, <3 x float> noundef zeroinitializer, float noundef 0.000000e+00)
60+
ret void
61+
}
62+
63+
declare spir_func <4 x float> @__builtin_IB_OCL_3d_sample_l(i32 noundef, i32 noundef, <3 x float> noundef, float noundef) local_unnamed_addr
64+
65+
66+
attributes #0 = { convergent nounwind }
67+
68+
!spirv.MemoryModel = !{!0}
69+
!spirv.Source = !{!1}
70+
!spirv.Generator = !{!2}
71+
!igc.functions = !{!3}
72+
!opencl.ocl.version = !{!5}
73+
!opencl.spir.version = !{!5}
74+
!IGCMetadata = !{!6}
75+
76+
!0 = !{i32 2, i32 2}
77+
!1 = !{i32 4, i32 100000}
78+
!2 = !{i16 6, i16 14}
79+
!3 = !{void (%spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @indirect_vs_direct_bc, !4}
80+
!4 = !{}
81+
!5 = !{i32 2, i32 0}
82+
!6 = !{!"ModuleMD", !7, !20}
83+
!7 = !{!"FuncMD", !8, !9}
84+
!8 = !{!"FuncMDMap[71]", void (%spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @indirect_vs_direct_bc}
85+
!9 = !{!"FuncMDValue[71]", !10}
86+
!10 = !{!"resAllocMD", !11}
87+
!11 = !{!"argAllocMDList", !12, !16}
88+
!12 = !{!"argAllocMDListVec[0]", !13, !14, !15}
89+
!13 = !{!"type", i32 4}
90+
!14 = !{!"extensionType", i32 0}
91+
!15 = !{!"indexType", i32 2}
92+
!16 = !{!"argAllocMDListVec[1]", !17, !18, !19}
93+
!17 = !{!"type", i32 5}
94+
!18 = !{!"extensionType", i32 -1}
95+
!19 = !{!"indexType", i32 0}
96+
!20 = !{!"UseBindlessImage", i1 true}

‎IGC/Compiler/tests/OCLBIConverter/value-tracker-find-alloca-val-indirect-direct-bitcast.ll‎

Lines changed: 96 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,96 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2024 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; Check kernel argument is successfully tracked when store address operand is
10+
; reached through other operations than load address operand. Such a thing is
11+
; possible because getelementptr instructions with all zero indices act like
12+
; bitcasts.
13+
14+
; RUN: igc_opt --opaque-pointers -igc-conv-ocl-to-common -S < %s -o - | FileCheck %s
15+
16+
; CHECK-NOT: assertion failed
17+
; CHECK: addrspacecast ptr addrspace(1) %a to ptr addrspace(393218)
18+
19+
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-n8:16:32"
20+
target triple = "spir64-unknown-unknown"
21+
22+
%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
23+
%"class.sycl::_V1::detail::array" = type { [3 x i64] }
24+
%spirv.Image._void_2_0_0_0_0_0_0 = type opaque
25+
%spirv.Sampler = type opaque
26+
%"class.sycl::_V1::detail::RoundedRangeIDGenerator" = type <{ %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", i8, [7 x i8] }>
27+
%"class.sycl::_V1::detail::RoundedRangeKernel" = type { %"class.sycl::_V1::range", %class.accessor_sampler }
28+
%class.accessor_sampler = type { %"class.sycl::_V1::accessor", %"class.sycl::_V1::sampler" }
29+
%"class.sycl::_V1::accessor" = type { %"class.sycl::_V1::detail::image_accessor" }
30+
%"class.sycl::_V1::detail::image_accessor" = type { %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, [24 x i8] }
31+
%"class.sycl::_V1::sampler" = type { %"class.sycl::_V1::detail::sampler_impl", [8 x i8] }
32+
%"class.sycl::_V1::detail::sampler_impl" = type { %spirv.Sampler addrspace(2)* }
33+
%spirv.SampledImage._void_2_0_0_0_0_0_0 = type opaque
34+
35+
; Function Attrs: convergent nounwind
36+
define spir_kernel void @indirect_vs_direct_bc(%spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)* %a, %spirv.Sampler addrspace(2)* %b) #0 {
37+
entry:
38+
%alloc = alloca %"class.sycl::_V1::detail::RoundedRangeKernel", align 8
39+
%store_addr = getelementptr inbounds %"class.sycl::_V1::detail::RoundedRangeKernel", %"class.sycl::_V1::detail::RoundedRangeKernel"* %alloc, i64 0, i32 1
40+
br label %bitcast_with_bc_and_gep
41+
42+
bitcast_with_bc_and_gep:
43+
%bc_store_addr = bitcast %class.accessor_sampler* %store_addr to %"class.sycl::_V1::detail::image_accessor"*
44+
%gep_as_bc = getelementptr inbounds %"class.sycl::_V1::detail::image_accessor", %"class.sycl::_V1::detail::image_accessor"* %bc_store_addr, i64 0, i32 0
45+
store %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)* %a, %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)** %gep_as_bc, align 8
46+
br label %bitcast_straight_to_desired_type
47+
48+
bitcast_straight_to_desired_type:
49+
%bc_store_addr_direct = bitcast %class.accessor_sampler* %store_addr to %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)**
50+
%load_img_ptr = load %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)** %bc_store_addr_direct, align 8
51+
br label %exit
52+
53+
exit:
54+
%to_i64 = ptrtoint %spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)* %load_img_ptr to i64
55+
%arg0 = trunc i64 %to_i64 to i32
56+
%sampler = ptrtoint %spirv.Sampler addrspace(2)* %b to i64
57+
%or_1 = or i64 %sampler, 1
58+
%arg1 = trunc i64 %or_1 to i32
59+
%bi_call = call spir_func <4 x float> @__builtin_IB_OCL_3d_sample_l(i32 noundef %arg0, i32 noundef %arg1, <3 x float> noundef zeroinitializer, float noundef 0.000000e+00)
60+
ret void
61+
}
62+
63+
declare spir_func <4 x float> @__builtin_IB_OCL_3d_sample_l(i32 noundef, i32 noundef, <3 x float> noundef, float noundef) local_unnamed_addr
64+
65+
66+
attributes #0 = { convergent nounwind }
67+
68+
!spirv.MemoryModel = !{!0}
69+
!spirv.Source = !{!1}
70+
!spirv.Generator = !{!2}
71+
!igc.functions = !{!3}
72+
!opencl.ocl.version = !{!5}
73+
!opencl.spir.version = !{!5}
74+
!IGCMetadata = !{!6}
75+
76+
!0 = !{i32 2, i32 2}
77+
!1 = !{i32 4, i32 100000}
78+
!2 = !{i16 6, i16 14}
79+
!3 = !{void (%spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @indirect_vs_direct_bc, !4}
80+
!4 = !{}
81+
!5 = !{i32 2, i32 0}
82+
!6 = !{!"ModuleMD", !7, !20}
83+
!7 = !{!"FuncMD", !8, !9}
84+
!8 = !{!"FuncMDMap[71]", void (%spirv.Image._void_2_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @indirect_vs_direct_bc}
85+
!9 = !{!"FuncMDValue[71]", !10}
86+
!10 = !{!"resAllocMD", !11}
87+
!11 = !{!"argAllocMDList", !12, !16}
88+
!12 = !{!"argAllocMDListVec[0]", !13, !14, !15}
89+
!13 = !{!"type", i32 4}
90+
!14 = !{!"extensionType", i32 0}
91+
!15 = !{!"indexType", i32 2}
92+
!16 = !{!"argAllocMDListVec[1]", !17, !18, !19}
93+
!17 = !{!"type", i32 5}
94+
!18 = !{!"extensionType", i32 -1}
95+
!19 = !{!"indexType", i32 0}
96+
!20 = !{!"UseBindlessImage", i1 true}

0 commit comments

Comments
 (0)