Skip to content

[SYCL][PTX][CUDA] Fixes debug info cloning for implicit global offset #2129

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 2 commits into from
Jul 17, 2020
Merged

[SYCL][PTX][CUDA] Fixes debug info cloning for implicit global offset #2129

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
f46e101
[SYCL][PTX][CUDA] Fixes debug info cloning for implicit global offset
Jul 16, 2020
63fec4e
Adds regression test
Jul 16, 2020
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
8 changes: 4 additions & 4 deletions llvm/lib/Target/NVPTX/SYCL/GlobalOffset.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -259,6 +259,9 @@ class GlobalOffset : public ModulePass {
Function *NewFunc = Function::Create(NewFuncTy, Func->getLinkage(),
Func->getAddressSpace());

// Keep original function ordering.
M.getFunctionList().insertAfter(Func->getIterator(), NewFunc);

if (KeepOriginal) {
// TODO: Are there better naming alternatives that allow for unmangling?
NewFunc->setName(Func->getName() + "_with_offset");
Expand All @@ -272,7 +275,7 @@ class GlobalOffset : public ModulePass {
}

SmallVector<ReturnInst *, 8> Returns;
CloneFunctionInto(NewFunc, Func, VMap, /*ModuleLevelChanges=*/false,
CloneFunctionInto(NewFunc, Func, VMap, /*ModuleLevelChanges=*/true,
Returns);
} else {
NewFunc->copyAttributesFrom(Func);
Expand All @@ -298,9 +301,6 @@ class GlobalOffset : public ModulePass {
NewFunc->addMetadata(MD.first, *MD.second);
}

// Keep original function ordering.
M.getFunctionList().insertAfter(Func->getIterator(), NewFunc);

Value *ImplicitOffset = NewFunc->arg_begin() + (NewFunc->arg_size() - 1);
// Add bitcast to match the return type of the intrinsic if needed.
if (ImplicitArgumentType != ImplicitOffsetPtrType) {
Expand Down
20 changes: 20 additions & 0 deletions sycl/test/regression/implicit_offset_debug_info.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
// RUN: %clangxx -g -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// REQUIRES: cuda

// NOTE: Tests that the implicit global offset pass copies debug information

#include <CL/sycl.hpp>
using namespace cl::sycl;

int main() {
queue q;
buffer<uint64_t, 1> t1(10);
q.submit([&](handler &cgh) {
auto table = t1.get_access<access::mode::write>(cgh);
cgh.parallel_for<class kernel>(10, [=](id<1> gtid) {
table[gtid] = gtid[0];
});
});
q.wait();
}