[clang] [llvm] [mlir] [OpenMP] Remove the LocationDescription constructor that dropped the location. (PR #221949)
Abid Qadeer via cfe-commits
cfe-commits at lists.llvm.org
Tue Sep 8 03:19:40 PDT 2026
https://github.com/abidh created https://github.com/llvm/llvm-project/pull/221949
LocationDescription had an implicit conversion from a bare insertion point, so `ompBuilder->emitSomething(someInsertPoint, ...)` compiled happily and silently produced an empty debug location. Because updateToLocation() installs the location unconditionally, this was worse than a missing assignment: it cleared whatever the builder was carrying, and the emitted runtime call ended up with no !dbg. On the device those calls are inlinable, so the verifier rejects them once the runtime carries debug info -- which is how this kept turning up as bug reports rather than as anything visible at the callsite.
With the callers in OMPIRBuilder, OpenMPOpt, the MLIR translation and clang all converted, the constructor can go, and the compiler will now refuse the shape that caused the problem. Callers have to say which location they mean: pass the
IRBuilder to take its current one, or spell out the insertion point and location as a pair. The last two conversions are still in flight, so this depends on both #221842 (the device shared memory allocations) and #219548 (the clang hand-off to OMPIRBuilder) landing first.
>From 2edce14c053eb8775f028c39efa1f738d5ef8dbf Mon Sep 17 00:00:00 2001
From: Abid Qadeer <haqadeer at amd.com>
Date: Fri, 21 Aug 2026 20:38:07 +0100
Subject: [PATCH 1/4] [OpenMP] Give the device shared memory allocations a
debug location.
Fixes https://github.com/llvm/llvm-project/issues/221831.
In generic mode on the device, the buffers that carry values into an
outlined region come from device shared memory rather than the stack, so
they are emitted as calls to __kmpc_alloc_shared and __kmpc_free_shared.
Those are definitions in the DeviceRTL, which makes them inlinable calls,
and the verifier requires an inlinable call in a function with debug info
to carry a !dbg location. None of these calls had one, for two separate
reasons.
1. allocateVar() and deallocateVar() took a bare insertion point, so an
override had no debug location to set on the runtime calls it emits in
place of the alloca the base class would have created. Fixed by adding
a DebugLoc parameter that carries one. emitReplacerCall() passes the
extracted region's first location, the same one it puts on the call to
the outlined function. The default implementations ignore it because a
plain alloca needs no location.
2. The createOMPAllocShared() and createOMPFreeShared() calls in
createParallel() relied on the implicit conversion from an insertion
point to a LocationDescription, which selects the constructor that
leaves the debug location empty, and updateToLocation() then installs
that empty location over whatever the builder had. Fixed by passing the
Builder, which carries the debug location along with the insertion
point. This is the same fix as #218961.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
.../llvm/Transforms/Utils/CodeExtractor.h | 11 +++--
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 21 ++++++---
llvm/lib/Transforms/Utils/CodeExtractor.cpp | 35 ++++++++-------
.../omptarget-debug-shared-alloc-loc.mlir | 43 +++++++++++++++++++
4 files changed, 84 insertions(+), 26 deletions(-)
create mode 100644 mlir/test/Target/LLVMIR/omptarget-debug-shared-alloc-loc.mlir
diff --git a/llvm/include/llvm/Transforms/Utils/CodeExtractor.h b/llvm/include/llvm/Transforms/Utils/CodeExtractor.h
index 05f8287aebf6f..fa52f9264eaea 100644
--- a/llvm/include/llvm/Transforms/Utils/CodeExtractor.h
+++ b/llvm/include/llvm/Transforms/Utils/CodeExtractor.h
@@ -18,6 +18,7 @@
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SetVector.h"
#include "llvm/IR/BasicBlock.h"
+#include "llvm/IR/DebugLoc.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/Support/Compiler.h"
#include <limits>
@@ -260,15 +261,17 @@ class LLVM_ABI CodeExtractor {
void excludeArgFromAggregate(Value *Arg);
protected:
- /// Allocate an intermediate variable at the specified point.
+ /// Allocate an intermediate variable at the specified point. \p DL is a debug
+ /// location for anything an override emits that needs one.
virtual Instruction *allocateVar(IRBuilder<>::InsertPoint AllocaIP,
- Type *VarType, const Twine &Name = Twine(""),
+ DebugLoc DL, Type *VarType,
+ const Twine &Name = Twine(""),
AddrSpaceCastInst **CastedAlloc = nullptr);
/// Deallocate a previously-allocated intermediate variable at the specified
- /// point.
+ /// point. \p DL is as for allocateVar().
virtual Instruction *deallocateVar(IRBuilder<>::InsertPoint DeallocIP,
- Value *Var, Type *VarType);
+ DebugLoc DL, Value *Var, Type *VarType);
private:
struct LifetimeMarkerInfo {
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 3f80c858d033e..c6019278aed23 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -545,15 +545,16 @@ class DeviceSharedMemCodeExtractor : public OMPCodeExtractor {
protected:
virtual Instruction *
- allocateVar(IRBuilder<>::InsertPoint AllocaIP, Type *VarType,
+ allocateVar(IRBuilder<>::InsertPoint AllocaIP, DebugLoc DL, Type *VarType,
const Twine &Name = Twine(""),
AddrSpaceCastInst **CastedAlloc = nullptr) override {
- return OMPBuilder.createOMPAllocShared(AllocaIP, VarType, Name);
+ return OMPBuilder.createOMPAllocShared({AllocaIP, DL}, VarType, Name);
}
virtual Instruction *deallocateVar(IRBuilder<>::InsertPoint DeallocIP,
- Value *Var, Type *VarType) override {
- return OMPBuilder.createOMPFreeShared(DeallocIP, Var, VarType);
+ DebugLoc DL, Value *Var,
+ Type *VarType) override {
+ return OMPBuilder.createOMPFreeShared({DeallocIP, DL}, Var, VarType);
}
};
@@ -2176,12 +2177,18 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createParallel(
Value *Ptr;
if (UsesDeviceSharedMemory) {
// Use device shared memory instead, if needed.
- Ptr = createOMPAllocShared(OuterAllocIP, V.getType(),
+ Ptr = createOMPAllocShared(Builder, V.getType(),
V.getName() + ".reloaded");
- for (BasicBlock *DeallocBlock : OuterDeallocBlocks)
+ for (BasicBlock *DeallocBlock : OuterDeallocBlocks) {
+ assert(DeallocBlock->getParent() ==
+ OuterAllocIP.getBlock()->getParent() &&
+ "Dealloc block is not in the function holding the allocation, "
+ "so its debug location cannot be reused there");
createOMPFreeShared(
- InsertPointTy(DeallocBlock, DeallocBlock->getFirstInsertionPt()),
+ {InsertPointTy(DeallocBlock, DeallocBlock->getFirstInsertionPt()),
+ Builder.getCurrentDebugLocation()},
Ptr, V.getType());
+ }
} else {
Ptr = Builder.CreateAlloca(V.getType(), nullptr,
V.getName() + ".reloaded");
diff --git a/llvm/lib/Transforms/Utils/CodeExtractor.cpp b/llvm/lib/Transforms/Utils/CodeExtractor.cpp
index 4f224d0a18e48..6a53688fb7027 100644
--- a/llvm/lib/Transforms/Utils/CodeExtractor.cpp
+++ b/llvm/lib/Transforms/Utils/CodeExtractor.cpp
@@ -451,8 +451,10 @@ CodeExtractor::findOrCreateBlockForHoisting(BasicBlock *CommonExitBlock) {
}
Instruction *CodeExtractor::allocateVar(IRBuilder<>::InsertPoint AllocaIP,
- Type *VarType, const Twine &Name,
+ DebugLoc, Type *VarType,
+ const Twine &Name,
AddrSpaceCastInst **CastedAlloc) {
+ // An alloca needs no debug location, so the one passed in goes unused here.
const DataLayout &DL = AllocaIP.getBlock()->getModule()->getDataLayout();
Instruction *Alloca = new AllocaInst(VarType, DL.getAllocaAddrSpace(),
nullptr, Name, AllocaIP.getPoint());
@@ -466,8 +468,8 @@ Instruction *CodeExtractor::allocateVar(IRBuilder<>::InsertPoint AllocaIP,
return Alloca;
}
-Instruction *CodeExtractor::deallocateVar(IRBuilder<>::InsertPoint, Value *,
- Type *) {
+Instruction *CodeExtractor::deallocateVar(IRBuilder<>::InsertPoint, DebugLoc,
+ Value *, Type *) {
// Default alloca instructions created by allocateVar are released implicitly.
return nullptr;
}
@@ -1874,6 +1876,13 @@ CallInst *CodeExtractor::emitReplacerCall(
BasicBlock *AllocaBlock =
AllocationBlock ? AllocationBlock : &oldFunction->getEntryBlock();
+ // If the original function has debug info, the terminator of the entry block
+ // of the extracted function contains the first debug location of the
+ // extracted function, set in extractCodeRegion.
+ DebugLoc DL;
+ if (oldFunction->getSubprogram())
+ DL = newFunction->getEntryBlock().getTerminator()->getDebugLoc();
+
// Update the entry count of the function.
if (BFI)
BFI->setBlockFreq(codeReplacer, EntryFreq);
@@ -1897,7 +1906,7 @@ CallInst *CodeExtractor::emitReplacerCall(
Value *OutAlloc =
allocateVar(IRBuilder<>::InsertPoint(
AllocaBlock, AllocaBlock->getFirstInsertionPt()),
- output->getType(), output->getName() + ".loc");
+ DL, output->getType(), output->getName() + ".loc");
params.push_back(OutAlloc);
ReloadOutputs.push_back(OutAlloc);
}
@@ -1907,7 +1916,7 @@ CallInst *CodeExtractor::emitReplacerCall(
AddrSpaceCastInst *StructSpaceCast = nullptr;
Struct = allocateVar(IRBuilder<>::InsertPoint(
AllocaBlock, AllocaBlock->getFirstInsertionPt()),
- StructArgTy, "structArg", &StructSpaceCast);
+ DL, StructArgTy, "structArg", &StructSpaceCast);
if (StructSpaceCast)
params.push_back(StructSpaceCast);
else
@@ -1949,13 +1958,9 @@ CallInst *CodeExtractor::emitReplacerCall(
}
// Add debug location to the new call, if the original function has debug
- // info. In that case, the terminator of the entry block of the extracted
- // function contains the first debug location of the extracted function,
- // set in extractCodeRegion.
- if (codeReplacer->getParent()->getSubprogram()) {
- if (auto DL = newFunction->getEntryBlock().getTerminator()->getDebugLoc())
- call->setDebugLoc(DL);
- }
+ // info.
+ if (DL)
+ call->setDebugLoc(DL);
// Reload the outputs passed in by reference, use the struct if output is in
// the aggregate or reload from the scalar argument.
@@ -2060,13 +2065,13 @@ CallInst *CodeExtractor::emitReplacerCall(
int Index = 0;
for (Value *Output : outputs) {
if (!StructValues.contains(Output))
- deallocateVar(IRBuilder<>::InsertPoint(DeallocBlock, DeallocIP),
+ deallocateVar(IRBuilder<>::InsertPoint(DeallocBlock, DeallocIP), DL,
ReloadOutputs[Index++], Output->getType());
}
if (Struct)
- deallocateVar(IRBuilder<>::InsertPoint(DeallocBlock, DeallocIP), Struct,
- StructArgTy);
+ deallocateVar(IRBuilder<>::InsertPoint(DeallocBlock, DeallocIP), DL,
+ Struct, StructArgTy);
};
if (DeallocationBlocks.empty()) {
diff --git a/mlir/test/Target/LLVMIR/omptarget-debug-shared-alloc-loc.mlir b/mlir/test/Target/LLVMIR/omptarget-debug-shared-alloc-loc.mlir
new file mode 100644
index 0000000000000..19f7fe573b4be
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/omptarget-debug-shared-alloc-loc.mlir
@@ -0,0 +1,43 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32>>, llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_gpu = true, omp.is_target_device = true} {
+ llvm.func @_QQmain(%arg0: !llvm.ptr) {
+ %0 = omp.map.info var_ptr(%arg0 : !llvm.ptr, i32) map_clauses(tofrom) capture(ByRef) name("") -> !llvm.ptr
+ omp.target kernel_type(generic) map_entries(%0 -> %arg1 : !llvm.ptr) {
+ %1 = llvm.load %arg1 : !llvm.ptr -> i32 loc(#loc1)
+ omp.parallel {
+ %2 = llvm.add %1, %1 : i32 loc(#loc1)
+ llvm.store %2, %arg1 : i32, !llvm.ptr loc(#loc1)
+ omp.terminator
+ } loc(#loc1)
+ omp.terminator
+ } loc(#loc3)
+ llvm.return
+ } loc(#loc2)
+}
+#file = #llvm.di_file<"target.f90" in "">
+#cu = #llvm.di_compile_unit<id = distinct[0]<>,
+ sourceLanguage = DW_LANG_Fortran95, file = #file, isOptimized = false,
+ emissionKind = LineTablesOnly>
+#sp_ty = #llvm.di_subroutine_type<callingConvention = DW_CC_normal>
+#sp = #llvm.di_subprogram<id = distinct[1]<>, compileUnit = #cu, scope = #file,
+ name = "_QQmain", file = #file, subprogramFlags = "Definition", type = #sp_ty>
+#sp1 = #llvm.di_subprogram<id = distinct[2]<>, compileUnit = #cu, scope = #file,
+ name = "__omp_offloading_target", file = #file, subprogramFlags = "Definition",
+ type = #sp_ty>
+#loc1 = loc("target.f90":12:5)
+#loc2 = loc(fused<#sp>[#loc1])
+#loc3 = loc(fused<#sp1>[#loc1])
+
+// Both the aggregate holding the outlined region's arguments and the buffer
+// forwarding the non-pointer value are allocated in device shared memory
+// rather than on the stack, so the runtime calls that allocate and free them
+// must carry a debug location, scoped to the correct function.
+
+// CHECK: define {{.*}}@__omp_offloading_{{.*}} !dbg ![[SP:[0-9]+]] {
+// CHECK: call {{.*}}@__kmpc_alloc_shared(i64 16), !dbg ![[LOC:[0-9]+]]
+// CHECK: call {{.*}}@__kmpc_alloc_shared(i64 4), !dbg ![[LOC]]
+// CHECK: call void @__kmpc_free_shared(ptr {{.*}}, i64 16), !dbg ![[LOC]]
+// CHECK: call void @__kmpc_free_shared(ptr {{.*}}, i64 4), !dbg ![[LOC]]
+// CHECK-DAG: ![[SP]] = distinct !DISubprogram(name: "__omp_offloading_target"
+// CHECK-DAG: ![[LOC]] = !DILocation(line: 12, column: 5, scope: ![[SP]])
>From 55b642ae8f68273fefa5e11cc8bcf5bcef14f3cb Mon Sep 17 00:00:00 2001
From: Abid Qadeer <haqadeer at amd.com>
Date: Fri, 21 Aug 2026 20:38:30 +0100
Subject: [PATCH 2/4] [CGOpenMPRuntime] Keep the current debug location when
handing off to OMPIRBuilder.
Both callsites took clang's insertion point but not its debug location, which
selected the LocationDescription constructor that leaves the location empty.
Because updateToLocation() installs the location unconditionally, that actively
cleared the location clang had established, and the __kmpc_global_thread_num
and target data calls emitted from there lost their !dbg. Passing the builder
carries the location across too.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 1a1479b4b9b7e..4222a860d0b95 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1417,7 +1417,7 @@ llvm::Value *CGOpenMPRuntime::getThreadID(CodeGenFunction &CGF,
// the clang invariants used below might be broken.
if (CGM.getLangOpts().OpenMPIRBuilder) {
SmallString<128> Buffer;
- OMPBuilder.updateToLocation(CGF.Builder.saveIP());
+ OMPBuilder.updateToLocation(CGF.Builder);
uint32_t SrcLocStrSize;
auto *SrcLocStr = OMPBuilder.getOrCreateSrcLocStr(
getIdentStringFromSourceLocation(CGF, Loc, Buffer), SrcLocStrSize);
@@ -11802,7 +11802,7 @@ void CGOpenMPRuntime::emitTargetDataCalls(
CGF.AllocaInsertPt->getIterator());
InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
CGF.Builder.GetInsertPoint());
- llvm::OpenMPIRBuilder::LocationDescription OmpLoc(CodeGenIP);
+ llvm::OpenMPIRBuilder::LocationDescription OmpLoc(CGF.Builder);
llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
cantFail(OMPBuilder.createTargetData(
OmpLoc, AllocaIP, CodeGenIP, /*DeallocBlocks=*/{}, DeviceID,
>From d0535f240807e6334cb93cbfece56c86280f7dd3 Mon Sep 17 00:00:00 2001
From: Abid Qadeer <haqadeer at amd.com>
Date: Fri, 28 Aug 2026 18:32:26 +0100
Subject: [PATCH 3/4] [CGOpenMPRuntime] Add a test for the debug location
hand-off to OMPIRBuilder.
Cover both places that lost the location: the thread-num call getThreadID()
delegates to the builder when -fopenmp-enable-irbuilder is on, and the branch
createTargetData() emits for the 'if' clause of a target data region. The
mapper calls in that region are not useful here because restoreIP() reinstalls
a location from the instruction at the insertion point, so they keep their
!dbg either way.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
.../OpenMP/debug-info-ompirbuilder-handoff.c | 35 +++++++++++++++++++
1 file changed, 35 insertions(+)
create mode 100644 clang/test/OpenMP/debug-info-ompirbuilder-handoff.c
diff --git a/clang/test/OpenMP/debug-info-ompirbuilder-handoff.c b/clang/test/OpenMP/debug-info-ompirbuilder-handoff.c
new file mode 100644
index 0000000000000..fe29b0e26dbd7
--- /dev/null
+++ b/clang/test/OpenMP/debug-info-ompirbuilder-handoff.c
@@ -0,0 +1,35 @@
+// Check that the debug location clang has established survives the hand-off to
+// the OpenMPIRBuilder, so that the IR the builder emits on clang's behalf still
+// carries a !dbg attachment.
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-enable-irbuilder -triple x86_64-unknown-unknown \
+// RUN: -debug-info-kind=limited -emit-llvm %s -o - | FileCheck %s --check-prefix=GTID
+
+// RUN: %clang_cc1 -fopenmp -triple x86_64-unknown-unknown \
+// RUN: -fopenmp-targets=x86_64-unknown-linux-gnu -debug-info-kind=limited \
+// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=TDATA
+
+int cond;
+void use(int);
+
+// CGOpenMPRuntime::getThreadID() defers to the OpenMPIRBuilder when it is
+// enabled, so the thread-num call is emitted by the builder.
+
+// GTID-LABEL: define {{.*}}@single_region
+// GTID: entry:
+// GTID-NEXT: call i32 @__kmpc_global_thread_num({{.*}}), !dbg
+void single_region(void) {
+#pragma omp single
+ use(1);
+}
+
+// CGOpenMPRuntime::emitTargetDataCalls() passes the 'if' condition down to
+// OpenMPIRBuilder::createTargetData(), which emits the branch on it.
+
+// TDATA-LABEL: define {{.*}}@target_data_if
+// TDATA: %[[TOBOOL:.+]] = icmp ne i32 %{{.+}}, 0, !dbg
+// TDATA-NEXT: br i1 %[[TOBOOL]], label %{{.+}}, label %{{.+}}, !dbg
+void target_data_if(int *p) {
+#pragma omp target data map(tofrom : p[0 : 4]) if (cond)
+ use(2);
+}
>From ccc8ee5bafa3071bbfd223b57d7d7a5b12ab6e7b Mon Sep 17 00:00:00 2001
From: Abid Qadeer <haqadeer at amd.com>
Date: Fri, 21 Aug 2026 20:38:42 +0100
Subject: [PATCH 4/4] [OpenMP] Remove the LocationDescription constructor that
dropped the location.
LocationDescription had an implicit conversion from a bare insertion point, so
`ompBuilder->emitSomething(someInsertPoint, ...)` compiled happily and silently
produced an empty debug location. Because updateToLocation() installs the
location unconditionally, this was worse than a missing assignment: it cleared
whatever the builder was carrying, and the emitted runtime call ended up with
no !dbg. On the device those calls are inlinable, so the verifier rejects them
once the runtime carries debug info -- which is how this kept turning up as
bug reports rather than as anything visible at the callsite.
With the callers in OMPIRBuilder, OpenMPOpt, the MLIR translation and clang all
converted, the constructor can go, and the compiler will now refuse the shape
that caused the problem. Callers have to say which location they mean: pass the
IRBuilder to take its current one, or spell out the insertion point and
location as a pair. The last two conversions are still in flight, so this
depends on both #221842 (the device shared memory allocations) and #219548 (the
clang hand-off to OMPIRBuilder) landing first.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h | 1 -
1 file changed, 1 deletion(-)
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 5f3da79cf53b1..8825e71e18c5c 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -744,7 +744,6 @@ class OpenMPIRBuilder {
struct LocationDescription {
LocationDescription(const IRBuilderBase &IRB)
: IP(IRB.saveIP()), DL(IRB.getCurrentDebugLocation()) {}
- LocationDescription(const InsertPointTy &IP) : IP(IP) {}
LocationDescription(const InsertPointTy &IP, const DebugLoc &DL)
: IP(IP), DL(DL) {}
InsertPointTy IP;
More information about the cfe-commits
mailing list