https://github.com/lfmeadow created https://github.com/llvm/llvm-project/pull/217193
The SizeEmitter lambda in emitCommonOMPTargetDirective constructed OMPLoopScope as an unnamed temporary, so it was destroyed at the end of its own statement rather than at the end of the enclosing block. OMPLoopScope is a RunCleanupsScope, and destroying it emits the cleanups for the '.capture_expr.' variables that hold the loop bounds, including their llvm.lifetime.end calls. Those cleanups therefore ran before EmitScalarExpr(D.getNumIterations()) loaded the bounds, so Clang emitted a load from an alloca whose lifetime had already ended: store i32 %sub, ptr %.capture_expr.1 call void @llvm.lifetime.end.p0(ptr %.capture_expr.1) %0 = load i32, ptr %.capture_expr.1 ; read after lifetime ended This was harmless while PromoteMemToReg simply discarded lifetime intrinsics. Since #191909, promotion inserts a 'store undef' in place of each marker, so the load now correctly resolves to undef and the trip count folds to zero. KernelArgsTy::Tripcount then reaches the runtime as 0, which can no longer size the launch: a 'target teams distribute parallel for' over 100000 iterations with OMP_NUM_TEAMS=50 launches one team instead of 50. Both the misplaced cleanup and #191909 are present on main, so this currently breaks OpenMP target offloading at head: every target loop directive whose launch geometry or scheduling depends on the trip count is affected. It was caught by the AOMP smoke suite, where it took out all of the trip-count-sensitive tests at once. Name the scope so that it outlives the trip count computation, matching every other use of OMPLoopScope in this file. cc @isoard-amd @ronlieb >From 73cb9002f8cfcf2cf180d29669d61e1c4d764fd9 Mon Sep 17 00:00:00 2001 From: Larry Meadows <[email protected]> Date: Tue, 18 Aug 2026 21:00:57 -0500 Subject: [PATCH] [Clang][OpenMP] Keep OMPLoopScope alive while computing target trip count The SizeEmitter lambda in emitCommonOMPTargetDirective constructed OMPLoopScope as an unnamed temporary, so it was destroyed at the end of its own statement rather than at the end of the enclosing block. OMPLoopScope is a RunCleanupsScope, and destroying it emits the cleanups for the '.capture_expr.' variables that hold the loop bounds, including their llvm.lifetime.end calls. Those cleanups therefore ran before EmitScalarExpr(D.getNumIterations()) loaded the bounds, so Clang emitted a load from an alloca whose lifetime had already ended: store i32 %sub, ptr %.capture_expr.1 call void @llvm.lifetime.end.p0(ptr %.capture_expr.1) %0 = load i32, ptr %.capture_expr.1 ; read after lifetime ended This was harmless while PromoteMemToReg simply discarded lifetime intrinsics. Since #191909, promotion inserts a 'store undef' in place of each marker, so the load now correctly resolves to undef and the trip count folds to zero. KernelArgsTy::Tripcount then reaches the runtime as 0, which can no longer size the launch: a 'target teams distribute parallel for' over 100000 iterations with OMP_NUM_TEAMS=50 launches one team instead of 50. Both the misplaced cleanup and #191909 are present on main, so this currently breaks OpenMP target offloading at head: every target loop directive whose launch geometry or scheduling depends on the trip count is affected. It was caught by the AOMP smoke suite, where it took out all of the trip-count-sensitive tests at once. Name the scope so that it outlives the trip count computation, matching every other use of OMPLoopScope in this file. --- clang/lib/CodeGen/CGStmtOpenMP.cpp | 2 +- .../target_tripcount_lifetime_codegen.c | 20 +++++++++++++++++++ 2 files changed, 21 insertions(+), 1 deletion(-) create mode 100644 clang/test/OpenMP/target_tripcount_lifetime_codegen.c diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 122e5db698548..9c341a4fc2641 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -7315,7 +7315,7 @@ static void emitCommonOMPTargetDirective(CodeGenFunction &CGF, [IsOffloadEntry](CodeGenFunction &CGF, const OMPLoopDirective &D) -> llvm::Value * { if (IsOffloadEntry) { - OMPLoopScope(CGF, D); + OMPLoopScope PreInitScope(CGF, D); // Emit calculation of the iterations count. llvm::Value *NumIterations = CGF.EmitScalarExpr(D.getNumIterations()); NumIterations = CGF.Builder.CreateIntCast(NumIterations, CGF.Int64Ty, diff --git a/clang/test/OpenMP/target_tripcount_lifetime_codegen.c b/clang/test/OpenMP/target_tripcount_lifetime_codegen.c new file mode 100644 index 0000000000000..56c47e37952c9 --- /dev/null +++ b/clang/test/OpenMP/target_tripcount_lifetime_codegen.c @@ -0,0 +1,20 @@ +// RUN: %clang_cc1 -verify -fopenmp -O1 -disable-llvm-passes -x c -triple x86_64-unknown-linux-gnu -fopenmp-targets=x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +// The loop bounds used to compute a target region's trip count live in +// '.capture_expr.' allocas created by OMPLoopScope. The scope has to outlive +// the trip count computation, otherwise its cleanups emit lifetime.end before +// the loads and the trip count reads dead memory. + +void f(int n, int *a) { +#pragma omp target teams distribute parallel for + for (int i = 0; i < n; ++i) + a[i] = i; +} + +// CHECK-LABEL: define {{.*}}@f( +// CHECK: store i32 %{{.+}}, ptr %[[CE:\.capture_expr\.[0-9]+]], align 4 +// CHECK-NEXT: [[TC:%.+]] = load i32, ptr %[[CE]], align 4 +// CHECK-NEXT: [[ADD:%.+]] = add nsw i32 [[TC]], 1 +// CHECK-NEXT: [[EXT:%.+]] = zext i32 [[ADD]] to i64 +// CHECK-NEXT: call void @llvm.lifetime.end.p0(ptr %[[CE]]) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
