Skip to content

Add spirv work-item builtin debug-info patch #196

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Oct 9, 2020
Merged

Add spirv work-item builtin debug-info patch #196

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
115 changes: 115 additions & 0 deletions patches/spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,115 @@
From 2dfb0384d51a9bb82fa985ebbd7b8b40aecbfd01 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he@intel.com>
Date: Fri, 9 Oct 2020 15:09:57 +0800
Subject: [PATCH] Fix debug info of work-item builtin translation (#745)

debug info of work-item builtins are lost in both llvm IR -> spirv and
spirv -> llvm IR translations. See #744
---
lib/SPIRV/OCL20ToSPIRV.cpp | 5 ++-
lib/SPIRV/SPIRVReader.cpp | 1 +
test/DebugInfo/builtin-get-global-id.ll | 60 +++++++++++++++++++++++++
3 files changed, 65 insertions(+), 1 deletion(-)
create mode 100644 test/DebugInfo/builtin-get-global-id.ll

diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp
index a5e1df7..c876128 100644
--- a/lib/SPIRV/OCL20ToSPIRV.cpp
+++ b/lib/SPIRV/OCL20ToSPIRV.cpp
@@ -1262,11 +1262,14 @@ void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() {
for (auto UI = I.user_begin(), UE = I.user_end(); UI != UE; ++UI) {
auto CI = dyn_cast<CallInst>(*UI);
assert(CI && "invalid instruction");
- Value *NewValue = new LoadInst(GVType, BV, "", CI);
+ const DebugLoc &DLoc = CI->getDebugLoc();
+ Instruction *NewValue = new LoadInst(GVType, BV, "", CI);
+ NewValue->setDebugLoc(DLoc);
LLVM_DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n');
if (IsVec) {
NewValue =
ExtractElementInst::Create(NewValue, CI->getArgOperand(0), "", CI);
+ NewValue->setDebugLoc(DLoc);
LLVM_DEBUG(dbgs() << *NewValue << '\n');
}
NewValue->takeName(CI);
diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp
index f02a746..bdb9d34 100644
--- a/lib/SPIRV/SPIRVReader.cpp
+++ b/lib/SPIRV/SPIRVReader.cpp
@@ -315,6 +315,7 @@ bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV,
auto Replace = [&](std::vector<Value *> Arg, Instruction *I) {
auto Call = CallInst::Create(Func, Arg, "", I);
Call->takeName(I);
+ Call->setDebugLoc(I->getDebugLoc());
setAttrByCalledFunc(Call);
SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << *Call
<< '\n';)
diff --git a/test/DebugInfo/builtin-get-global-id.ll b/test/DebugInfo/builtin-get-global-id.ll
new file mode 100644
index 0000000..a4a00e6
--- /dev/null
+++ b/test/DebugInfo/builtin-get-global-id.ll
@@ -0,0 +1,60 @@
+; Check debug info of builtin get_global_id is preserved from LLVM IR to spirv
+; and spirv to LLVM IR translation.
+
+; Original .cl source:
+; kernel void test() {
+; size_t gid = get_global_id(0);
+; }
+
+; Command line:
+; ./clang -cc1 1.cl -triple spir64 -cl-std=cl2.0 -emit-llvm -finclude-default-header -debug-info-kind=line-tables-only -O0
+
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix CHECK-SPIRV
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64"
+
+; CHECK-SPIRV: ExtInst {{.*}} DebugScope
+; CHECK-SPIRV-NEXT: Line {{[0-9]+}} 2 16
+; CHECK-SPIRV-NEXT: Load {{[0-9]+}} [[LoadRes:[0-9]+]]
+; CHECK-SPIRV-NEXT: CompositeExtract {{[0-9]+}} {{[0-9]+}} [[LoadRes]] 0
+
+; Function Attrs: convergent noinline norecurse nounwind optnone
+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 {
+entry:
+ %gid = alloca i64, align 8
+ %call = call spir_func i64 @_Z13get_global_idj(i32 0) #2, !dbg !10
+; CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0) #1, !dbg [[DBG:![0-9]+]]
+ store i64 %call, i64* %gid, align 8, !dbg !11
+ ret void, !dbg !12
+}
+
+; Function Attrs: convergent nounwind readnone
+declare spir_func i64 @_Z13get_global_idj(i32) #1
+
+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" }
+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" }
+attributes #2 = { convergent nounwind readnone }
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!3, !4}
+!opencl.ocl.version = !{!5}
+!opencl.spir.version = !{!5}
+!llvm.ident = !{!6}
+
+!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)
+!1 = !DIFile(filename: "<stdin>", directory: "")
+!2 = !{}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !{i32 1, !"wchar_size", i32 4}
+!5 = !{i32 2, i32 0}
+!6 = !{!"clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)"}
+!7 = distinct !DISubprogram(name: "test", scope: !8, file: !8, line: 1, type: !9, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0, retainedNodes: !2)
+!8 = !DIFile(filename: "1.cl", directory: "")
+!9 = !DISubroutineType(types: !2)
+!10 = !DILocation(line: 2, column: 16, scope: !7)
+!11 = !DILocation(line: 2, column: 10, scope: !7)
+!12 = !DILocation(line: 3, column: 1, scope: !7)
--
2.18.1