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

Signed-off-by: Steffen Larsen <[email protected]>

Author: steffenlarsen

llvm/lib/Target/NVPTX/SYCL/GlobalOffset.cpp

@@ -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");
@@ -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);
@@ -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) {

sycl/test/regression/implicit_offset_debug_info.cpp

@@ -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();
+}