Skip to content

Commit 96e49dc

Browse files
fzou1fzou1
authored
Merge pull request #195 from hewj03/ocl-open-100
Add spirv work-item builtin debug-info patch
2 parents 3e9401e + 746fbe9 commit 96e49dc

File tree

1 file changed

+115
-0
lines changed

1 file changed

+115
-0
lines changed

‎patches/spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch

Lines changed: 115 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,115 @@
1+
From 64f58776c28c7842b9b771bc3118da0b8cbd12b3 Mon Sep 17 00:00:00 2001
2+
From: Wenju He <[email protected]>
3+
Date: Fri, 9 Oct 2020 14:21:09 +0800
4+
Subject: [PATCH] Fix debug info of work-item builtin translation (#745)
5+
6+
debug info of work-item builtins are lost in both llvm IR -> spirv and
7+
spirv -> llvm IR translations. See #744
8+
---
9+
lib/SPIRV/OCL20ToSPIRV.cpp | 5 ++-
10+
lib/SPIRV/SPIRVReader.cpp | 1 +
11+
test/DebugInfo/builtin-get-global-id.ll | 60 +++++++++++++++++++++++++
12+
3 files changed, 65 insertions(+), 1 deletion(-)
13+
create mode 100644 test/DebugInfo/builtin-get-global-id.ll
14+
15+
diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp
16+
index 1262c48..a742c8c 100644
17+
--- a/lib/SPIRV/OCL20ToSPIRV.cpp
18+
+++ b/lib/SPIRV/OCL20ToSPIRV.cpp
19+
@@ -1297,11 +1297,14 @@ void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() {
20+
for (auto UI = I.user_begin(), UE = I.user_end(); UI != UE; ++UI) {
21+
auto CI = dyn_cast<CallInst>(*UI);
22+
assert(CI && "invalid instruction");
23+
- Value *NewValue = new LoadInst(BV, "", CI);
24+
+ const DebugLoc &DLoc = CI->getDebugLoc();
25+
+ Instruction *NewValue = new LoadInst(BV, "", CI);
26+
+ NewValue->setDebugLoc(DLoc);
27+
LLVM_DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n');
28+
if (IsVec) {
29+
NewValue =
30+
ExtractElementInst::Create(NewValue, CI->getArgOperand(0), "", CI);
31+
+ NewValue->setDebugLoc(DLoc);
32+
LLVM_DEBUG(dbgs() << *NewValue << '\n');
33+
}
34+
NewValue->takeName(CI);
35+
diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp
36+
index 0003e0a..03b4ea1 100644
37+
--- a/lib/SPIRV/SPIRVReader.cpp
38+
+++ b/lib/SPIRV/SPIRVReader.cpp
39+
@@ -299,6 +299,7 @@ bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV,
40+
Arg.push_back(EEI->getIndexOperand());
41+
auto Call = CallInst::Create(Func, Arg, "", I);
42+
Call->takeName(I);
43+
+ Call->setDebugLoc(I->getDebugLoc());
44+
setAttrByCalledFunc(Call);
45+
SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << *Call
46+
<< '\n';)
47+
diff --git a/test/DebugInfo/builtin-get-global-id.ll b/test/DebugInfo/builtin-get-global-id.ll
48+
new file mode 100644
49+
index 0000000..a4a00e6
50+
--- /dev/null
51+
+++ b/test/DebugInfo/builtin-get-global-id.ll
52+
@@ -0,0 +1,60 @@
53+
+; Check debug info of builtin get_global_id is preserved from LLVM IR to spirv
54+
+; and spirv to LLVM IR translation.
55+
+
56+
+; Original .cl source:
57+
+; kernel void test() {
58+
+; size_t gid = get_global_id(0);
59+
+; }
60+
+
61+
+; Command line:
62+
+; ./clang -cc1 1.cl -triple spir64 -cl-std=cl2.0 -emit-llvm -finclude-default-header -debug-info-kind=line-tables-only -O0
63+
+
64+
+; RUN: llvm-as %s -o %t.bc
65+
+; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix CHECK-SPIRV
66+
+; RUN: llvm-spirv %t.bc -o %t.spv
67+
+; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s
68+
+
69+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
70+
+target triple = "spir64"
71+
+
72+
+; CHECK-SPIRV: ExtInst {{.*}} DebugScope
73+
+; CHECK-SPIRV-NEXT: Line {{[0-9]+}} 2 16
74+
+; CHECK-SPIRV-NEXT: Load {{[0-9]+}} [[LoadRes:[0-9]+]]
75+
+; CHECK-SPIRV-NEXT: CompositeExtract {{[0-9]+}} {{[0-9]+}} [[LoadRes]] 0
76+
+
77+
+; Function Attrs: convergent noinline norecurse nounwind optnone
78+
+define spir_kernel void @test() #0 !dbg !7 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 {
79+
+entry:
80+
+ %gid = alloca i64, align 8
81+
+ %call = call spir_func i64 @_Z13get_global_idj(i32 0) #2, !dbg !10
82+
+; CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0) #1, !dbg [[DBG:![0-9]+]]
83+
+ store i64 %call, i64* %gid, align 8, !dbg !11
84+
+ ret void, !dbg !12
85+
+}
86+
+
87+
+; Function Attrs: convergent nounwind readnone
88+
+declare spir_func i64 @_Z13get_global_idj(i32) #1
89+
+
90+
+attributes #0 = { convergent noinline norecurse nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" "unsafe-fp-math"="false" "use-soft-float"="false" }
91+
+attributes #1 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
92+
+attributes #2 = { convergent nounwind readnone }
93+
+
94+
+!llvm.dbg.cu = !{!0}
95+
+!llvm.module.flags = !{!3, !4}
96+
+!opencl.ocl.version = !{!5}
97+
+!opencl.spir.version = !{!5}
98+
+!llvm.ident = !{!6}
99+
+
100+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2, nameTableKind: None)
101+
+!1 = !DIFile(filename: "<stdin>", directory: "")
102+
+!2 = !{}
103+
+!3 = !{i32 2, !"Debug Info Version", i32 3}
104+
+!4 = !{i32 1, !"wchar_size", i32 4}
105+
+!5 = !{i32 2, i32 0}
106+
+!6 = !{!"clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)"}
107+
+!7 = distinct !DISubprogram(name: "test", scope: !8, file: !8, line: 1, type: !9, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0, retainedNodes: !2)
108+
+!8 = !DIFile(filename: "1.cl", directory: "")
109+
+!9 = !DISubroutineType(types: !2)
110+
+!10 = !DILocation(line: 2, column: 16, scope: !7)
111+
+!11 = !DILocation(line: 2, column: 10, scope: !7)
112+
+!12 = !DILocation(line: 3, column: 1, scope: !7)
113+
--
114+
2.18.1
115+

0 commit comments

Comments
 (0)