From 80ccee5411e0dfee67a8ff8708f6e5f8e45ff2ac Mon Sep 17 00:00:00 2001
From: Abid Qadeer <haqadeer@amd.com>
Date: Mon, 11 Aug 2025 17:44:11 +0100
Subject: [PATCH] [OMPIRBuilder] Avoid invalid debug location.

Fixes #153043.

This is another case of debug location not getting updated when the
insert point is changed by the restoreIP. Fixed by using the wrapper
function that updates the debug location.
---
 clang/test/OpenMP/amdgcn_debug_nowait.c   | 16 ++++++++++++++++
 llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp |  2 +-
 2 files changed, 17 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/OpenMP/amdgcn_debug_nowait.c

diff --git a/clang/test/OpenMP/amdgcn_debug_nowait.c b/clang/test/OpenMP/amdgcn_debug_nowait.c
new file mode 100644
index 0000000000000..d691327512ff7
--- /dev/null
+++ b/clang/test/OpenMP/amdgcn_debug_nowait.c
@@ -0,0 +1,16 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -debug-info-kind=line-tables-only -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+
+int test() {
+  int c;
+
+#pragma omp target data map(tofrom: c)
+{
+  #pragma omp target nowait
+  {
+      c = 2;
+  }
+}
+  return c;
+}
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 2ac9a0d3fbd66..c16b0dde1a3da 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -7247,7 +7247,7 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createTargetData(
           BodyGenCB(Builder.saveIP(), BodyGenTy::NoPriv);
       if (!AfterIP)
         return AfterIP.takeError();
-      Builder.restoreIP(*AfterIP);
+      restoreIPandDebugLoc(Builder, *AfterIP);
 
       if (IfCond)
         return emitIfClause(IfCond, EndThenGen, EndElseGen, AllocaIP);