diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 2d5bf9b20802b..6f3031161bfff 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -2256,6 +2256,169 @@ void CGOpenMPRuntime::emitTaskyieldCall(CodeGenFunction &CGF,
     Region->emitUntiedSwitch(CGF);
 }
 
+/// Emit a helper with the runtime relocation signature (kmp_task_relocate_t):
+///   void relocate(kmp_task_t *task, void *outer_captures);
+///
+/// On taskgraph replay the runtime invokes this helper to refresh the task's
+/// shared-pointer table. Each capture (a shared-by-ref variable or \c this)
+/// that the task body actually dereferences at execution time is
+/// re-projected from the freshly reconstructed outer record passed as
+/// \p outer_captures and stored back into \c task->shareds.
+///
+/// Captures that the body cannot observe a changed address for across
+/// replays are skipped here:
+///
+///   * captures of a variable that appears as a firstprivate list item
+///     -- the body sources the value from the per-task '.kmp_privates.t'
+///     snapshot rather than from the shareds slot, so the (potentially
+///     stale) original address in the shareds entry is harmless;
+///
+///   * captures of a variable with static (global / namespace-scope /
+///     static-local / static-data-member) storage duration -- the
+///     captured pointer is the variable's link-time-fixed address, which
+///     is identical at recording and on every replay, so no re-projection
+///     is meaningful.
+///
+/// The relocate helper is therefore only ever called upon to refresh
+/// shareds slots that the body genuinely depends on at execution time
+/// (shared-by-ref to a local variable, captured \c this on a heap or
+/// stack object, etc.).  When every capture falls in one of the
+/// skip-eligible categories the helper is emitted as a (still non-null)
+/// no-op: today's runtime only inspects null-vs-non-null, and a non-null
+/// no-op is the right signal that there is nothing the body actually
+/// needs the shareds table refreshed for.
+///
+/// Reduction captures of a local-stack variable still keep the existing
+/// null-relocate-and-abort behaviour: the taskred runtime state is keyed
+/// off the recording-time taskgroup hierarchy and is not currently usable
+/// on replay, so it is preferable to fail loudly (#302) than to silently
+/// misbehave.  Reduction captures of a static-storage variable do not run
+/// into this hazard at the relocate layer -- the captured pointer is
+/// stable -- and are no-op-skipped via the static-storage rule above;
+/// whether the reduction body itself then succeeds on replay is a
+/// separate concern.
+///
+/// Returns null only when at least one capture is genuinely shared (none
+/// of the skip-eligible categories apply) AND cannot be resolved in
+/// \p OuterCSI; in that case the caller passes a null relocation function
+/// to the runtime and the runtime fails fast at replay.
+static llvm::Function *
+emitTaskRelocationFunction(CodeGenModule &CGM, SourceLocation Loc,
+                           const CapturedStmt &CS,
+                           const CodeGenFunction::CGCapturedStmtInfo *OuterCSI,
+                           const OMPTaskDataTy &Data) {
+  ASTContext &C = CGM.getContext();
+
+  // Variables that don't need their shareds slot refreshed across replays
+  // because the body sources them from the per-task '.kmp_privates.t'
+  // snapshot.  Today this is the set of firstprivate list items (snapshot
+  // is taken at task allocation and reused unchanged by every replay).
+  llvm::SmallPtrSet<const VarDecl *, 8> NoRelocateFirstprivateVars;
+  for (const Expr *E : Data.FirstprivateVars) {
+    if (!E)
+      continue;
+    if (const auto *DRE = dyn_cast<DeclRefExpr>(E->IgnoreParenImpCasts()))
+      if (const auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
+        NoRelocateFirstprivateVars.insert(VD->getCanonicalDecl());
+  }
+
+  // A capture is "no-op-safe" with respect to taskgraph replay when
+  // refreshing its shareds slot is provably unnecessary - either because
+  // the body never reads from that slot (firstprivate) or because the
+  // captured pointer is a link-time-fixed address and is therefore
+  // identical at every replay (static storage duration).
+  auto IsNoOpRelocate = [&](const CapturedStmt::Capture &Cap) {
+    if (Cap.capturesThis() || !Cap.capturesVariable())
+      return false;
+    const VarDecl *VD = Cap.getCapturedVar();
+    if (VD->hasGlobalStorage())
+      return true;
+    return NoRelocateFirstprivateVars.contains(VD->getCanonicalDecl());
+  };
+
+  auto LookupOuterField =
+      [&](const CapturedStmt::Capture &Cap) -> const FieldDecl * {
+    if (!OuterCSI)
+      return nullptr;
+    return Cap.capturesThis() ? OuterCSI->getThisFieldDecl()
+                              : OuterCSI->lookup(Cap.getCapturedVar());
+  };
+
+  // Bail out before emitting any IR if a genuinely-shared capture cannot
+  // be resolved in the containing context.  No-op-safe captures (see the
+  // function-level comment) don't participate in this preflight; they
+  // simply cause the helper to skip their slot below.
+  if (llvm::any_of(CS.captures(), [&](const CapturedStmt::Capture &Cap) {
+        assert((Cap.capturesThis() || Cap.capturesVariable()) &&
+               "OpenMP task capture must be shared-by-ref or 'this'");
+        return !IsNoOpRelocate(Cap) && !LookupOuterField(Cap);
+      }))
+    return nullptr;
+
+  // void relocate(void *task, void *outer_captures)
+  auto *TaskArg =
+      ImplicitParamDecl::Create(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
+                                C.VoidPtrTy, ImplicitParamKind::Other);
+  auto *OuterArg =
+      ImplicitParamDecl::Create(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
+                                C.VoidPtrTy, ImplicitParamKind::Other);
+  FunctionArgList Args{TaskArg, OuterArg};
+  const CGFunctionInfo &FnInfo =
+      CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
+
+  std::string Name =
+      CGM.getOpenMPRuntime().getName({"omp", "taskgraph", "relocate", ""});
+  auto *Fn = llvm::Function::Create(CGM.getTypes().GetFunctionType(FnInfo),
+                                    llvm::GlobalValue::InternalLinkage, Name,
+                                    &CGM.getModule());
+  CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo);
+  if (!CGM.getCodeGenOpts().SampleProfileFile.empty())
+    Fn->addFnAttr("sample-profile-suffix-elision-policy", "selected");
+  Fn->setDoesNotRecurse();
+
+  CodeGenFunction CGF(CGM);
+  CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, FnInfo, Args, Loc, Loc);
+
+  CGBuilderTy &Bld = CGF.Builder;
+  CharUnits PtrAlign = CGF.getPointerAlign();
+
+  // Base of the reconstructed outer record for this replay.
+  llvm::Value *OuterRaw = Bld.CreateLoad(CGF.GetAddrOfLocalVar(OuterArg));
+
+  // kmp_task_t::shareds is the first field of the runtime task descriptor;
+  // load it to obtain the void* shared table that we will refresh in place.
+  // The table holds one void* per by-ref capture.
+  llvm::Value *TaskRaw = Bld.CreateLoad(CGF.GetAddrOfLocalVar(TaskArg));
+  llvm::Value *SharedRaw =
+      Bld.CreateLoad(Address(TaskRaw, CGF.VoidPtrTy, PtrAlign));
+  Address SharedTable(SharedRaw, CGF.VoidPtrTy, PtrAlign);
+
+  unsigned Index = 0;
+  for (const CapturedStmt::Capture &Cap : CS.captures()) {
+    // Always advance the slot index so that we stay aligned with the
+    // shareds-table layout established at task allocation.
+    unsigned ThisIndex = Index++;
+    if (IsNoOpRelocate(Cap))
+      continue;
+    // Project the capture's referent from the freshly reconstructed outer
+    // record. EmitLValueForField auto-loads the outer reference field, so
+    // the resulting pointer is the live referent address (not the slot).
+    const FieldDecl *OuterField = LookupOuterField(Cap);
+    assert(OuterField && "preflight should have rejected this capture");
+    QualType OuterTy =
+        C.getCanonicalTagType(cast<RecordDecl>(OuterField->getDeclContext()));
+    LValue OuterBase = CGF.MakeAddrLValue(
+        Address(OuterRaw, CGF.ConvertTypeForMem(OuterTy), PtrAlign), OuterTy);
+    llvm::Value *Mapped =
+        CGF.EmitLValueForField(OuterBase, OuterField).getPointer(CGF);
+    Mapped = Bld.CreatePointerBitCastOrAddrSpaceCast(Mapped, CGM.VoidPtrTy);
+    Bld.CreateStore(Mapped, Bld.CreateConstGEP(SharedTable, ThisIndex));
+  }
+
+  CGF.FinishFunction();
+  return Fn;
+}
+
 void CGOpenMPRuntime::emitTaskgraphCall(CodeGenFunction &CGF,
                                         SourceLocation Loc,
                                         const OMPExecutableDirective &D,
@@ -4822,22 +4985,27 @@ void CGOpenMPRuntime::emitTaskCall(
     TGTaskArgs[2] = Result.NewTask;
     TGTaskArgs[3] = TaskAllocArgs[0]; // TaskFlags
     TGTaskArgs[4] = TaskAllocArgs[1]; // KmpTaskTWithPrivatesTySize
-    TGTaskArgs[5] = Shareds.emitRawPointer(CGF);
-    TGTaskArgs[6] = TaskAllocArgs[2]; // SharedsSize
+    TGTaskArgs[5] = TaskAllocArgs[2]; // SharedsSize
     if (auto RecType = dyn_cast<RecordType>(SharedsTy)) {
       auto *RD = RecType->getAsRecordDecl();
       if (RD->fields().empty()) {
         // FIXME: The condition might not be precisely correct here.
-        TGTaskArgs[6] = CGF.Builder.getSize(0);
+        TGTaskArgs[5] = CGF.Builder.getSize(0);
       }
     }
     if (Data.Dependences.size() == 0) {
-      TGTaskArgs[7] = CGF.Builder.getInt32(0);
-      TGTaskArgs[8] = llvm::ConstantPointerNull::get(CGF.VoidPtrTy);
+      TGTaskArgs[6] = CGF.Builder.getInt32(0);
+      TGTaskArgs[7] = llvm::ConstantPointerNull::get(CGF.VoidPtrTy);
     } else {
-      TGTaskArgs[7] = NumOfElements;
-      TGTaskArgs[8] = DependenciesArray.emitRawPointer(CGF);
-    }
+      TGTaskArgs[6] = NumOfElements;
+      TGTaskArgs[7] = DependenciesArray.emitRawPointer(CGF);
+    }
+    const auto *CS = cast<CapturedStmt>(D.getAssociatedStmt());
+    llvm::Function *RelocFn =
+        emitTaskRelocationFunction(CGM, Loc, *CS, CGF.CapturedStmtInfo, Data);
+    TGTaskArgs[8] = RelocFn ? CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+                                  RelocFn, CGM.VoidPtrTy)
+                            : llvm::ConstantPointerNull::get(CGF.VoidPtrTy);
     CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                             CGM.getModule(), OMPRTL___kmpc_taskgraph_task),
                         TGTaskArgs);
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 9f2c331a092e7..9b951d7b480b7 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -360,7 +360,7 @@ __OMP_RTL(__kmpc_taskgroup, false, Void, IdentPtr, Int32)
 __OMP_RTL(__kmpc_taskgraph, false, Void, IdentPtr, Int32, VoidPtrPtr, SizeTy,
           Int32, Int32, VoidPtr, VoidPtr)
 __OMP_RTL(__kmpc_taskgraph_task, false, Int32, IdentPtr, Int32, VoidPtr, Int32,
-          SizeTy, VoidPtr, SizeTy, Int32, VoidPtr)
+          SizeTy, SizeTy, Int32, VoidPtr, VoidPtr)
 __OMP_RTL(__kmpc_taskgraph_taskloop, false, Int32, IdentPtr, Int32, VoidPtr,
           Int32, SizeTy, VoidPtr, SizeTy, Int32, Int64Ptr, Int64Ptr, Int64,
           Int32, Int32, Int64, Int32, VoidPtr)
diff --git a/openmp/runtime/src/i18n/en_US.txt b/openmp/runtime/src/i18n/en_US.txt
index 46f5e794e5a27..5fc34123f4776 100644
--- a/openmp/runtime/src/i18n/en_US.txt
+++ b/openmp/runtime/src/i18n/en_US.txt
@@ -484,6 +484,7 @@ AffIgnoringNonHybrid         "%1$s ignored: This machine is not a hybrid archite
 AffIgnoringNotAvailable      "%1$s ignored: %2$s is not available. Using \"%3$s\" instead."
 SetNumTeamsInParOrTeamsRegion "%1$s() called inside a parallel or teams region; call ignored."
 SetTeamsThreadLimitInParOrTeamsRegion "%1$s() called inside a parallel or teams region; call ignored."
+OmpTaskgraphBadCapture       "Cannot locate captured shared variable reference for taskgraph replay"
 OmpTaskgraphConcurrentRecord "%1$s: multiple threads attempting to re-record taskgraph concurrently for graph_id %2$d"
 
 # --------------------------------------------------------------------------------------------------
diff --git a/openmp/runtime/src/kmp.h b/openmp/runtime/src/kmp.h
index f0366e36af899..5ad6b88db6d35 100644
--- a/openmp/runtime/src/kmp.h
+++ b/openmp/runtime/src/kmp.h
@@ -2469,6 +2469,7 @@ extern kmp_uint64 __kmp_taskloop_min_tasks;
 /*!
  */
 typedef kmp_int32 (*kmp_routine_entry_t)(kmp_int32, void *);
+typedef void (*kmp_task_relocate_t)(struct kmp_task *, void *);
 
 typedef union kmp_cmplrdata {
   kmp_int32 priority; /**< priority specified by user for the task */
@@ -2690,6 +2691,7 @@ enum kmp_taskgraph_region_type {
 typedef struct kmp_taskgraph_node {
   kmp_task_t *task;
   bool taskloop_task;
+  kmp_task_relocate_t relocate;
   kmp_taskgraph_reduce_input_data_t *reduce_input;
   union {
     // Valid when KMP_TDG_RECORDING in parent taskgraph record.
@@ -2773,6 +2775,7 @@ typedef struct kmp_taskgraph_record {
   struct kmp_taskgraph_exec_descr *exec_descrs = nullptr;
   kmp_size_t num_exec_descrs = 0;
   kmp_lock_t replay_lock;
+  void *taskgraph_args = nullptr;
   // We need a taskgroup structure to keep track of recorded tasks.  This is
   // set to TRUE if the user requested "nogroup" on the taskgraph directive
   // (then we can avoid blocking at the end of the taskgraph region on replay,
@@ -4530,8 +4533,8 @@ KMP_EXPORT void __kmpc_taskgraph(ident_t *loc_ref, kmp_int32 gtid,
                                  void *args);
 KMP_EXPORT kmp_uint32 __kmpc_taskgraph_task(
     ident_t *loc_ref, kmp_int32 gtid, kmp_task_t *new_task, kmp_int32 flags,
-    size_t sizeof_kmp_task_t, void *shareds, size_t sizeof_shareds,
-    kmp_int32 ndeps, kmp_depend_info_t *dep_list);
+    size_t sizeof_kmp_task_t, size_t sizeof_shareds, kmp_int32 ndeps,
+    kmp_depend_info_t *dep_list, kmp_task_relocate_t reloc);
 KMP_EXPORT kmp_uint32 __kmpc_taskgraph_taskloop(
     ident_t *loc_ref, kmp_int32 gtid, kmp_task_t *new_task, kmp_int32 flags,
     size_t sizeof_kmp_task_t, void *shareds, size_t sizeof_shareds,
diff --git a/openmp/runtime/src/kmp_tasking.cpp b/openmp/runtime/src/kmp_tasking.cpp
index 2a04d9bb58812..b500652f5f10f 100644
--- a/openmp/runtime/src/kmp_tasking.cpp
+++ b/openmp/runtime/src/kmp_tasking.cpp
@@ -2495,10 +2495,11 @@ __kmp_build_exec_descrs(kmp_info_t *thread, kmp_taskgraph_region_t *region,
 
 /// Reset, reparent and regroup the recorded task TASK and re-invoke it.
 
-static void __kmp_omp_tg_task(kmp_int32 gtid, kmp_task_t *task,
+static void __kmp_omp_tg_task(kmp_int32 gtid, kmp_taskgraph_node_t *node,
                               kmp_taskgroup_t *taskgroup,
                               kmp_taskdata_t *parent_taskdata,
                               bool serialize_immediate) {
+  kmp_task_t *task = node->task;
   kmp_taskdata_t *taskdata = KMP_TASK_TO_TASKDATA(task);
   taskdata->td_parent = parent_taskdata;
 
@@ -2521,6 +2522,18 @@ static void __kmp_omp_tg_task(kmp_int32 gtid, kmp_task_t *task,
   if (parent_taskdata->td_flags.tasktype == TASK_EXPLICIT)
     KMP_ATOMIC_INC(&parent_taskdata->td_allocated_child_tasks);
 
+  if (node->relocate) {
+    // Call the task's relocation function with the incoming args from the
+    // owning taskgraph.  This rewrites capture-by-reference variables to point
+    // to the correct location on the replayed taskgraph's stack (which may not
+    // be the same as the location from the initial recorded taskgraph).
+    node->relocate(task, taskdata->owning_taskgraph->taskgraph_args);
+  } else if (task->shareds != NULL) {
+    // A missing relocation callback is only fatal when there is a non-empty
+    // shareds payload that may contain by-reference captures needing remap.
+    KMP_FATAL(OmpTaskgraphBadCapture);
+  }
+
   __kmp_omp_task(gtid, task, false);
 }
 
@@ -2543,7 +2556,7 @@ static void __kmp_taskgraph_exec_descr_start(kmp_int32 gtid, kmp_info_t *thread,
     // __kmp_taskgraph_exec_descr_finish.
     kmp_taskgraph_node_t *node = descr->region->task.node;
     kmp_taskdata_t *current_taskdata = thread->th.th_current_task;
-    __kmp_omp_tg_task(gtid, node->task, taskgroup, current_taskdata, false);
+    __kmp_omp_tg_task(gtid, node, taskgroup, current_taskdata, false);
     break;
   }
   case TASKGRAPH_REGION_PARALLEL:
@@ -5104,6 +5117,7 @@ __kmp_taskgraph_node_alloc(kmp_taskgraph_record_t *rec, kmp_task_t *task,
 
   new_task->task = task;
   new_task->taskloop_task = false;
+  new_task->relocate = nullptr;
   new_task->reduce_input = nullptr;
   new_task->u.unresolved.ndeps = 0;
   new_task->u.unresolved.dep_list = nullptr;
@@ -5946,10 +5960,6 @@ static kmp_task_t *__kmp_taskgraph_clone_task(kmp_info_t *thread,
   // FIXME: This should use a "taskdup" function like taskloops in cases where
   // private variables are not trivially copyable.  For now, do it by plain
   // bitwise copy.
-  // FIXME 2: It's intended that this copy be persistent, and can be
-  // re-executed on taskgraph replay.  Make sure that works (for shared
-  // variables) if stack addresses change (i.e. a task-generating function is
-  // called from different call stack depths).
   kmp_taskdata_t *taskdata = KMP_TASK_TO_TASKDATA(orig);
   size_t shareds_offset = sizeof(kmp_taskdata_t) + sizeof_kmp_task_t;
   shareds_offset = __kmp_round_up_to_val(shareds_offset, sizeof(kmp_uint64));
@@ -5958,6 +5968,11 @@ static kmp_task_t *__kmp_taskgraph_clone_task(kmp_info_t *thread,
   KMP_MEMCPY(copy_td, taskdata, shareds_offset + sizeof_shareds);
   // Tasks cloned for a taskgraph always have this field set.
   copy_td->owning_taskgraph = taskgraph;
+  kmp_task_t *copy_task = KMP_TASKDATA_TO_TASK(copy_td);
+  if (orig->shareds) {
+    // New task's shared data has now moved.  Update the pointer.
+    copy_task->shareds = (void *)((char *)copy_td + shareds_offset);
+  }
   KMP_ATOMIC_ST_RLX(&copy_td->td_incomplete_child_tasks, 0);
   return KMP_TASKDATA_TO_TASK(copy_td);
 }
@@ -6073,6 +6088,9 @@ void __kmpc_taskgraph(ident_t *loc_ref, kmp_int32 gtid,
     // taskgroup.
     KMP_ATOMIC_ST_REL(&taskgroup->taskgraph.recording, record);
   }
+  // Keep the current taskgraph invocation's outlined-entry args for
+  // replay-time relocation of by-reference captures.
+  record->taskgraph_args = args;
   __kmp_release_lock(&header->header_lock, gtid);
 
   kmp_taskgraph_status_t status = KMP_ATOMIC_LD_ACQ(&record->status);
@@ -6108,9 +6126,10 @@ void __kmpc_taskgraph(ident_t *loc_ref, kmp_int32 gtid,
 
 kmp_uint32 __kmpc_taskgraph_task(ident_t *loc_ref, kmp_int32 gtid,
                                  kmp_task_t *new_task, kmp_int32 flags,
-                                 size_t sizeof_kmp_task_t, void *shareds,
+                                 size_t sizeof_kmp_task_t,
                                  size_t sizeof_shareds, kmp_int32 ndeps,
-                                 kmp_depend_info_t *dep_list) {
+                                 kmp_depend_info_t *dep_list,
+                                 kmp_task_relocate_t relocate) {
   kmp_info_t *thread = __kmp_threads[gtid];
   kmp_taskgroup_t *taskgroup = thread->th.th_current_task->td_taskgroup;
   kmp_taskgraph_record_t *rec = __kmp_taskgraph_or_parent_recording(taskgroup);
@@ -6146,6 +6165,7 @@ kmp_uint32 __kmpc_taskgraph_task(ident_t *loc_ref, kmp_int32 gtid,
           thread, ndeps * sizeof(kmp_depend_info_t));
       KMP_MEMCPY(node->u.unresolved.dep_list, dep_list,
                  ndeps * sizeof(kmp_depend_info_t));
+      node->relocate = relocate;
     } else if (status == KMP_TDG_READY) {
 #ifdef DEBUG_TASKGRAPH
       fprintf(stderr,
diff --git a/openmp/runtime/test/taskgraph/taskgraph_firstprivate_stack_depth.cpp b/openmp/runtime/test/taskgraph/taskgraph_firstprivate_stack_depth.cpp
new file mode 100644
index 0000000000000..e3f8976bc1ae0
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_firstprivate_stack_depth.cpp
@@ -0,0 +1,112 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t && env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+static volatile int StackSink = 0;
+// Keep the observable result in stable storage so this test isolates
+// firstprivate replay across different stack depths. There is a separate bug
+// test for replayed tasks writing through a stack-local shared pointer cached
+// at record time.
+static volatile int ReplayResult = -1;
+
+struct Payload {
+  int values[6];
+  int bias;
+};
+
+__attribute__((noinline)) static int evaluate_payload(const Payload &payload,
+                                                      int seed) {
+  return seed * payload.values[0] - payload.values[1] +
+         payload.values[2] * payload.values[3] - payload.bias +
+         payload.values[4] - payload.values[5];
+}
+
+__attribute__((noinline)) static void clobber_stack(int base) {
+  volatile int scratch[4096];
+
+  for (int i = 0; i < 4096; ++i)
+    scratch[i] = base + i;
+
+  StackSink += scratch[base & 63];
+}
+
+__attribute__((noinline)) static int run_taskgraph(int seed) {
+  Payload payload{
+      {seed + 1, seed + 3, seed + 5, seed + 7, seed + 11, seed + 13},
+      seed * 17 + 19};
+  ReplayResult = -1;
+
+#pragma omp taskgraph graph_id(91)
+  {
+#pragma omp task firstprivate(payload, seed) shared(ReplayResult)
+    {
+      ReplayResult = evaluate_payload(payload, seed);
+    }
+  }
+
+  return ReplayResult;
+}
+
+__attribute__((noinline)) static int call_with_depth(int seed, int depth) {
+  volatile int padding[128];
+
+  for (int i = 0; i < 128; ++i)
+    padding[i] = seed + depth + i;
+
+  StackSink += padding[(seed + depth) & 127];
+
+  if (depth == 0)
+    return run_taskgraph(seed);
+  return call_with_depth(seed, depth - 1);
+}
+
+__attribute__((noinline)) static int expected_result(int seed) {
+  Payload payload{
+      {seed + 1, seed + 3, seed + 5, seed + 7, seed + 11, seed + 13},
+      seed * 17 + 19};
+  return evaluate_payload(payload, seed);
+}
+
+int main() {
+  constexpr int NumCalls = 4;
+  constexpr int Seeds[NumCalls] = {3, 17, 29, 41};
+  constexpr int Depths[NumCalls] = {0, 3, 1, 5};
+
+  int recorded = -1;
+  bool failed = false;
+
+#pragma omp parallel num_threads(4)
+  {
+#pragma omp single
+    {
+      recorded = call_with_depth(Seeds[0], Depths[0]);
+      if (recorded != expected_result(Seeds[0])) {
+        std::fprintf(stderr, "FAIL initial record got=%d expected=%d\n",
+                     recorded, expected_result(Seeds[0]));
+        failed = true;
+      }
+
+      for (int i = 1; i < NumCalls; ++i) {
+        clobber_stack(Seeds[i] * 1000);
+        const int replayed = call_with_depth(Seeds[i], Depths[i]);
+        if (replayed != recorded) {
+          std::fprintf(stderr,
+                       "FAIL replay %d depth=%d seed=%d got=%d expected=%d\n",
+                       i, Depths[i], Seeds[i], replayed, recorded);
+          failed = true;
+        }
+      }
+    }
+  }
+
+  if (failed)
+    return 1;
+
+  std::fprintf(stderr, "PASS firstprivate stack result=%d\n", recorded);
+  return 0;
+}
+
+// CHECK: PASS firstprivate stack result=
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_mixed_capture.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_mixed_capture.cpp
new file mode 100644
index 0000000000000..0ea8e4bdadbb3
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_mixed_capture.cpp
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+__attribute__((noinline)) static int run_taskgraph_mixed_capture(int seed) {
+  int x = seed;
+  int y = seed * 2;
+  int out = -1;
+  int fp = 7;
+
+#pragma omp taskgraph graph_id(401)
+  {
+#pragma omp task replayable(1) shared(x, y, out) firstprivate(fp)              \
+    depend(inout : x, y)
+    {
+      x += fp;
+      y += x;
+      out = y + fp;
+    }
+  }
+
+  return out;
+}
+
+int main() {
+  const int first = run_taskgraph_mixed_capture(1);
+  const int second = run_taskgraph_mixed_capture(100);
+
+  if (first != 17 || second != 314) {
+    std::fprintf(stderr,
+                 "FAIL lexical mixed capture replay first=%d second=%d "
+                 "expected=17/314\n",
+                 first, second);
+    return 1;
+  }
+
+  std::fprintf(stderr, "PASS lexical mixed capture replay first=%d second=%d\n",
+               first, second);
+  return 0;
+}
+
+// CHECK: PASS lexical mixed capture replay first=17 second=314
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_nontrivial_type.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_nontrivial_type.cpp
new file mode 100644
index 0000000000000..f3a50a27176d4
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_nontrivial_type.cpp
@@ -0,0 +1,60 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+struct Tracker {
+  static int Ctors;
+  static int Dtors;
+
+  int Value;
+
+  explicit Tracker(int V) : Value(V) { ++Ctors; }
+  ~Tracker() { ++Dtors; }
+
+  void bump(int Delta) { Value += Delta; }
+};
+
+int Tracker::Ctors = 0;
+int Tracker::Dtors = 0;
+
+__attribute__((noinline)) static int run_taskgraph_nontrivial(int seed) {
+  Tracker Obj(seed);
+  int out = -1;
+
+#pragma omp taskgraph graph_id(403)
+  {
+#pragma omp task replayable(1) shared(Obj, out)
+    {
+      Obj.bump(11);
+      out = Obj.Value;
+    }
+  }
+
+  return out;
+}
+
+int main() {
+  const int first = run_taskgraph_nontrivial(1);
+  const int second = run_taskgraph_nontrivial(100);
+
+  if (first != 12 || second != 111 || Tracker::Ctors < 2 ||
+      Tracker::Dtors < 2 || Tracker::Ctors != Tracker::Dtors) {
+    std::fprintf(
+        stderr,
+        "FAIL lexical nontrivial replay first=%d second=%d ctors=%d dtors=%d\n",
+        first, second, Tracker::Ctors, Tracker::Dtors);
+    return 1;
+  }
+
+  std::fprintf(
+      stderr,
+      "PASS lexical nontrivial replay first=%d second=%d ctors=%d dtors=%d\n",
+      first, second, Tracker::Ctors, Tracker::Dtors);
+  return 0;
+}
+
+// CHECK: PASS lexical nontrivial replay first=12 second=111
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_nontrivial_type_recursive.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_nontrivial_type_recursive.cpp
new file mode 100644
index 0000000000000..6012b6194e56b
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_nontrivial_type_recursive.cpp
@@ -0,0 +1,88 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+struct Tracker {
+  static int Ctors;
+  static int Dtors;
+
+  int Value;
+
+  explicit Tracker(int V) : Value(V) { ++Ctors; }
+  ~Tracker() { ++Dtors; }
+
+  void bump(int Delta) { Value += Delta; }
+};
+
+int Tracker::Ctors = 0;
+int Tracker::Dtors = 0;
+
+__attribute__((noinline)) static int expected_recursive(int depth, int seed,
+                                                        int run_tag) {
+  int local = seed + (depth + 1) * 5 + run_tag;
+  if (depth == 0)
+    return local;
+  return local + expected_recursive(depth - 1, seed + 9, run_tag);
+}
+
+__attribute__((noinline)) static int
+run_recursive_nontrivial(int depth, int seed, int run_tag) {
+  Tracker Obj(seed);
+  int out = -1;
+
+  int gid = 500 + depth;
+#pragma omp taskgraph graph_id(gid)
+  {
+#pragma omp task replayable(1) shared(Obj, out, depth, run_tag)
+    {
+      Obj.bump((depth + 1) * 5 + run_tag);
+      out = Obj.Value;
+    }
+  }
+
+  if (depth == 0)
+    return out;
+  return out + run_recursive_nontrivial(depth - 1, seed + 9, run_tag);
+}
+
+int main() {
+  const int depth = 3;
+  int total_actual = 0;
+  int total_expected = 0;
+
+  for (int run = 0; run < 3; ++run) {
+    const int seed = 100 * run + 1;
+    const int actual = run_recursive_nontrivial(depth, seed, run);
+    const int expected = expected_recursive(depth, seed, run);
+
+    if (actual != expected) {
+      std::fprintf(stderr,
+                   "FAIL recursive nontrivial run=%d actual=%d expected=%d\n",
+                   run, actual, expected);
+      return 1;
+    }
+
+    total_actual += actual;
+    total_expected += expected;
+  }
+
+  if (Tracker::Ctors != Tracker::Dtors || Tracker::Ctors < 12) {
+    std::fprintf(stderr,
+                 "FAIL recursive nontrivial lifetime ctors=%d dtors=%d "
+                 "total=%d expected=%d\n",
+                 Tracker::Ctors, Tracker::Dtors, total_actual, total_expected);
+    return 1;
+  }
+
+  std::fprintf(
+      stderr,
+      "PASS recursive nontrivial total=%d expected=%d ctors=%d dtors=%d\n",
+      total_actual, total_expected, Tracker::Ctors, Tracker::Dtors);
+  return 0;
+}
+
+// CHECK: PASS recursive nontrivial total=
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_pointer.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_pointer.cpp
new file mode 100644
index 0000000000000..7690aff37853c
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_pointer.cpp
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+__attribute__((noinline)) static int run_taskgraph_pointer_shared(int seed) {
+  int value = seed;
+  int *ptr = &value;
+  int out = -1;
+
+#pragma omp taskgraph graph_id(402)
+  {
+#pragma omp task replayable(1) shared(ptr, out) depend(inout : value)
+    {
+      *ptr += 3;
+      out = *ptr;
+    }
+  }
+
+  return out;
+}
+
+int main() {
+  const int first = run_taskgraph_pointer_shared(1);
+  const int second = run_taskgraph_pointer_shared(100);
+
+  if (first != 4 || second != 103) {
+    std::fprintf(stderr,
+                 "FAIL lexical pointer shared replay first=%d second=%d "
+                 "expected=4/103\n",
+                 first, second);
+    return 1;
+  }
+
+  std::fprintf(stderr,
+               "PASS lexical pointer shared replay first=%d second=%d\n", first,
+               second);
+  return 0;
+}
+
+// CHECK: PASS lexical pointer shared replay first=4 second=103
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_pointer_recursive_frameid.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_pointer_recursive_frameid.cpp
new file mode 100644
index 0000000000000..142a5941f97ae
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_pointer_recursive_frameid.cpp
@@ -0,0 +1,76 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdint>
+#include <cstdio>
+
+__attribute__((noinline)) static int expected_recursive(int depth, int seed,
+                                                        int run_tag) {
+  int value = seed;
+  value += (depth + 1) * 3 + run_tag;
+  if (depth == 0)
+    return value;
+  return value + expected_recursive(depth - 1, seed + 7, run_tag);
+}
+
+__attribute__((noinline)) static int run_recursive_frameid(int depth, int seed,
+                                                           int run_tag) {
+  int value = seed;
+  int *ptr = &value;
+  int *&ptr_ref = ptr;
+  int out = -1;
+
+  // Typically, if captured pointers refer to locations on the stack, that
+  // would not be safe for taskgraph record/replay because we in general we
+  // cannot rewrite such pointers to point to the current (live) stack frame.
+  //
+  // This is one possible way around that though: we keep a taskgraph record
+  // per stack-depth, each of which may refer to the local stack frame.
+  //
+  // I probably wouldn't recommend use of this technique in production code.
+  uintptr_t frame_gid = reinterpret_cast<uintptr_t>(__builtin_frame_address(0));
+
+#pragma omp taskgraph graph_id(frame_gid)
+  {
+#pragma omp task shared(ptr_ref, out, depth, run_tag) depend(inout : value)
+    {
+      *ptr_ref += (depth + 1) * 3 + run_tag;
+      out = *ptr_ref;
+    }
+  }
+
+  if (depth == 0)
+    return out;
+  return out + run_recursive_frameid(depth - 1, seed + 7, run_tag);
+}
+
+int main() {
+  const int depth = 3;
+  int actual_sum = 0;
+  int expected_sum = 0;
+
+  for (int run = 0; run < 3; ++run) {
+    int seed = 100 * run + 1;
+    int actual = run_recursive_frameid(depth, seed, run);
+    int expected = expected_recursive(depth, seed, run);
+    if (actual != expected) {
+      std::fprintf(
+          stderr,
+          "FAIL recursive pointer frameid run=%d actual=%d expected=%d\n", run,
+          actual, expected);
+      return 1;
+    }
+    actual_sum += actual;
+    expected_sum += expected;
+  }
+
+  std::fprintf(stderr,
+               "PASS recursive pointer frameid runs=3 total=%d expected=%d\n",
+               actual_sum, expected_sum);
+  return 0;
+}
+
+// CHECK: PASS recursive pointer frameid runs=3 total=
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_recursive.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_recursive.cpp
new file mode 100644
index 0000000000000..4d0d7b5613108
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_recursive.cpp
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+__attribute__((noinline)) static int run_taskgraph_recursive(int depth,
+                                                             int seed) {
+  int x = seed;
+  int out = -1;
+
+#pragma omp taskgraph graph_id(450)
+  {
+#pragma omp task replayable(1) shared(x, out, depth) depend(inout : x)
+    {
+      x += depth + 1;
+      out = x;
+    }
+  }
+
+  if (depth == 0)
+    return out;
+
+  return out + run_taskgraph_recursive(depth - 1, seed + 10);
+}
+
+int main() {
+  const int first = run_taskgraph_recursive(3, 1);
+  const int second = run_taskgraph_recursive(3, 100);
+
+  if (first != 74 || second != 470) {
+    std::fprintf(
+        stderr,
+        "FAIL lexical recursive replay first=%d second=%d expected=74/470\n",
+        first, second);
+    return 1;
+  }
+
+  std::fprintf(stderr, "PASS lexical recursive replay first=%d second=%d\n",
+               first, second);
+  return 0;
+}
+
+// CHECK: PASS lexical recursive replay first=74 second=470
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_works.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_works.cpp
new file mode 100644
index 0000000000000..82be9c82d3e7c
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_lexical_shared_works.cpp
@@ -0,0 +1,42 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+__attribute__((noinline)) static int run_taskgraph_lexical(int seed) {
+  int x = seed;
+  int out = -1;
+
+#pragma omp taskgraph graph_id(311)
+  {
+#pragma omp task replayable(1) shared(x, out) depend(inout : x)
+    {
+      x += 5;
+      out = x;
+    }
+  }
+
+  return out;
+}
+
+int main() {
+  const int first = run_taskgraph_lexical(1);
+  const int second = run_taskgraph_lexical(100);
+
+  if (first != 6 || second != 105) {
+    std::fprintf(
+        stderr,
+        "FAIL lexical shared replay first=%d second=%d expected=6/105\n", first,
+        second);
+    return 1;
+  }
+
+  std::fprintf(stderr, "PASS lexical shared replay first=%d second=%d\n", first,
+               second);
+  return 0;
+}
+
+// CHECK: PASS lexical shared replay first=6 second=105
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_nonlexical_shared_fails_1.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_nonlexical_shared_fails_1.cpp
new file mode 100644
index 0000000000000..e5483ae5e06bf
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_nonlexical_shared_fails_1.cpp
@@ -0,0 +1,52 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %not --crash %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+// This seems like it could work in principle, but in general we don't know
+// where the targets of the referenced variables are when the task is replayed.
+__attribute__((noinline)) static void emit_nonlexical_task(int &x, int &out) {
+#pragma omp task replayable(1) shared(x, out) depend(inout : x)
+  {
+    x += 5;
+    out = x;
+  }
+}
+
+__attribute__((noinline)) static int run_taskgraph_nonlexical(int seed) {
+  int x = seed;
+  int out = -1;
+
+#pragma omp taskgraph graph_id(312)
+  {
+    emit_nonlexical_task(x, out);
+  }
+
+  return out;
+}
+
+int main() {
+  const int recorded = run_taskgraph_nonlexical(1);
+  const int replayed = run_taskgraph_nonlexical(100);
+
+  // The "non-lexical" replayable task is emitted in a helper function outside
+  // the taskgraph lexical scope.  We expect this to raise a runtime error.
+  if (recorded == replayed) {
+    std::fprintf(
+        stderr,
+        "UNEXPECTED SUCCESS nonlexical replay recorded=%d replayed=%d\n",
+        recorded, replayed);
+    return 0;
+  }
+
+  std::fprintf(stderr,
+               "EXPECTED FAILURE nonlexical replay recorded=%d replayed=%d\n",
+               recorded, replayed);
+  return 1;
+}
+
+// CHECK: OMP: Error #{{[0-9]+}}: Cannot locate captured shared variable reference for
+// taskgraph replay
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_nonlexical_shared_fails_2.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_nonlexical_shared_fails_2.cpp
new file mode 100644
index 0000000000000..ee593b28e3c8e
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_nonlexical_shared_fails_2.cpp
@@ -0,0 +1,70 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %not --crash %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// clang-format on
+
+#include <cstdio>
+
+__attribute__((noinline)) static int emit_nonlexical_task(int seed) {
+  int x = seed;
+  int out = -1;
+
+// This is syntactically valid, but a taskgraph replay that includes this
+// task cannot possibly succeed, because the stack frame containing 'x' and
+// 'out' doesn't exist at replay time.  We can raise a runtime error in that
+// case.
+// This isn't a compile error because the code is still valid if no taskgraph
+// record/replay is in progress.
+#pragma omp task replayable(1) shared(x, out) depend(inout : x)
+  {
+    x += 5;
+    out = x;
+  }
+
+  return out;
+}
+
+__attribute__((noinline)) static int run_taskgraph_nonlexical(int seed) {
+  int out;
+
+#pragma omp taskgraph graph_id(312)
+  {
+#pragma omp task shared(out)
+    {
+      out = emit_nonlexical_task(seed);
+    }
+  }
+
+  return out;
+}
+
+int main() {
+  int out = emit_nonlexical_task(50);
+  if (out != 55) {
+    std::fprintf(stderr,
+                 "UNEXPECTED FAILURE: task outside taskgraph returned %d\n",
+                 out);
+  }
+
+  const int recorded = run_taskgraph_nonlexical(1);
+  const int replayed = run_taskgraph_nonlexical(100);
+
+  // The non-lexical replayable task is emitted in a helper function outside
+  // the taskgraph lexical scope.
+  if (recorded == replayed) {
+    std::fprintf(
+        stderr,
+        "UNEXPECTED SUCCESS nonlexical replay recorded=%d replayed=%d\n",
+        recorded, replayed);
+    return 0;
+  }
+
+  std::fprintf(stderr,
+               "EXPECTED FAILURE nonlexical replay recorded=%d replayed=%d\n",
+               recorded, replayed);
+  return 1;
+}
+
+// CHECK: OMP: Error #{{[0-9]+}}: Cannot locate captured shared variable reference for
+// taskgraph replay
diff --git a/openmp/runtime/test/taskgraph/taskgraph_replayable_saved_stack_depth.cpp b/openmp/runtime/test/taskgraph/taskgraph_replayable_saved_stack_depth.cpp
new file mode 100644
index 0000000000000..922cb85a53eec
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_replayable_saved_stack_depth.cpp
@@ -0,0 +1,119 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t && env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// XFAIL: *
+// clang-format on
+
+#include <cstdio>
+
+static volatile int StackSink = 0;
+static volatile int ReplayResult = -1;
+
+struct Payload {
+  int values[6];
+  int bias;
+};
+
+__attribute__((noinline)) static int evaluate_payload(const Payload &payload,
+                                                      int seed) {
+  return seed * payload.values[0] - payload.values[1] +
+         payload.values[2] * payload.values[3] - payload.bias +
+         payload.values[4] - payload.values[5];
+}
+
+__attribute__((noinline)) static void clobber_stack(int base) {
+  volatile int scratch[4096];
+
+  for (int i = 0; i < 4096; ++i)
+    scratch[i] = base + i;
+
+  StackSink += scratch[base & 63];
+}
+
+// Intended future usage of firstprivate(saved: ...): this replayable task is
+// not lexically nested within the taskgraph directive. It is created from a
+// helper function called inside the taskgraph region, and needs closure-like
+// capture of that helper's stack locals.
+__attribute__((noinline)) static void emit_replayable_task(int seed) {
+  Payload payload{
+      {seed + 1, seed + 3, seed + 5, seed + 7, seed + 11, seed + 13},
+      seed * 17 + 19};
+
+#pragma omp task replayable(1) firstprivate(saved : payload, seed)             \
+    shared(ReplayResult)
+  {
+    ReplayResult = evaluate_payload(payload, seed);
+  }
+}
+
+__attribute__((noinline)) static int run_taskgraph(int seed) {
+  ReplayResult = -1;
+
+#pragma omp taskgraph graph_id(93)
+  {
+    emit_replayable_task(seed);
+  }
+
+  return ReplayResult;
+}
+
+__attribute__((noinline)) static int call_with_depth(int seed, int depth) {
+  volatile int padding[128];
+
+  for (int i = 0; i < 128; ++i)
+    padding[i] = seed + depth + i;
+
+  StackSink += padding[(seed + depth) & 127];
+
+  if (depth == 0)
+    return run_taskgraph(seed);
+  return call_with_depth(seed, depth - 1);
+}
+
+__attribute__((noinline)) static int expected_result(int seed) {
+  Payload payload{
+      {seed + 1, seed + 3, seed + 5, seed + 7, seed + 11, seed + 13},
+      seed * 17 + 19};
+  return evaluate_payload(payload, seed);
+}
+
+int main() {
+  constexpr int NumCalls = 4;
+  constexpr int Seeds[NumCalls] = {3, 17, 29, 41};
+  constexpr int Depths[NumCalls] = {0, 3, 1, 5};
+
+  int recorded = -1;
+  bool failed = false;
+
+#pragma omp parallel num_threads(4)
+  {
+#pragma omp single
+    {
+      recorded = call_with_depth(Seeds[0], Depths[0]);
+      if (recorded != expected_result(Seeds[0])) {
+        std::fprintf(stderr, "FAIL initial record got=%d expected=%d\n",
+                     recorded, expected_result(Seeds[0]));
+        failed = true;
+      }
+
+      for (int i = 1; i < NumCalls; ++i) {
+        clobber_stack(Seeds[i] * 1000);
+        const int replayed = call_with_depth(Seeds[i], Depths[i]);
+        if (replayed != recorded) {
+          std::fprintf(stderr,
+                       "FAIL replay %d depth=%d seed=%d got=%d expected=%d\n",
+                       i, Depths[i], Seeds[i], replayed, recorded);
+          failed = true;
+        }
+      }
+    }
+  }
+
+  if (failed)
+    return 1;
+
+  std::fprintf(stderr, "PASS replayable saved stack result=%d\n", recorded);
+  return 0;
+}
+
+// CHECK: PASS replayable saved stack result=
diff --git a/openmp/runtime/test/taskgraph/taskgraph_shared_stack_depth.cpp b/openmp/runtime/test/taskgraph/taskgraph_shared_stack_depth.cpp
new file mode 100644
index 0000000000000..ad579e8ed1b3d
--- /dev/null
+++ b/openmp/runtime/test/taskgraph/taskgraph_shared_stack_depth.cpp
@@ -0,0 +1,94 @@
+// clang-format off
+// RUN: %clangXX %flags %openmp_flags -fopenmp-version=60 %s -o %t
+// RUN: env OMP_NUM_THREADS=4 %libomp-run 2>&1 | FileCheck %s
+// REQUIRES: omp_taskgraph_experimental
+// XFAIL: *
+// clang-format on
+
+#include <cstdio>
+
+static volatile int StackSink = 0;
+
+struct Payload {
+  int values[6];
+  int bias;
+};
+
+__attribute__((noinline)) static int evaluate_payload(const Payload &payload,
+                                                      int seed) {
+  return seed * payload.values[0] - payload.values[1] +
+         payload.values[2] * payload.values[3] - payload.bias +
+         payload.values[4] - payload.values[5];
+}
+
+__attribute__((noinline)) static void clobber_stack(int base) {
+  volatile int scratch[4096];
+
+  for (int i = 0; i < 4096; ++i)
+    scratch[i] = base + i;
+
+  StackSink += scratch[base & 63];
+}
+
+__attribute__((noinline)) static int run_taskgraph(int seed) {
+  Payload payload{
+      {seed + 1, seed + 3, seed + 5, seed + 7, seed + 11, seed + 13},
+      seed * 17 + 19};
+  int result = -1;
+
+#pragma omp taskgraph graph_id(92)
+  {
+#pragma omp task firstprivate(payload, seed) shared(result)
+    {
+      result = evaluate_payload(payload, seed);
+    }
+  }
+
+  return result;
+}
+
+__attribute__((noinline)) static int call_with_depth(int seed, int depth) {
+  volatile int padding[128];
+
+  for (int i = 0; i < 128; ++i)
+    padding[i] = seed + depth + i;
+
+  StackSink += padding[(seed + depth) & 127];
+
+  if (depth == 0)
+    return run_taskgraph(seed);
+  return call_with_depth(seed, depth - 1);
+}
+
+__attribute__((noinline)) static int expected_result(int seed) {
+  Payload payload{
+      {seed + 1, seed + 3, seed + 5, seed + 7, seed + 11, seed + 13},
+      seed * 17 + 19};
+  return evaluate_payload(payload, seed);
+}
+
+int main() {
+  constexpr int RecordSeed = 3;
+  constexpr int ReplaySeed = 17;
+
+  const int recorded = call_with_depth(RecordSeed, 0);
+  if (recorded != expected_result(RecordSeed)) {
+    std::fprintf(stderr, "FAIL initial record got=%d expected=%d\n", recorded,
+                 expected_result(RecordSeed));
+    return 1;
+  }
+
+  clobber_stack(ReplaySeed * 1000);
+  const int replayed = call_with_depth(ReplaySeed, 3);
+  if (replayed != recorded) {
+    std::fprintf(
+        stderr, "BUG shared stack replay depth=%d seed=%d got=%d expected=%d\n",
+        3, ReplaySeed, replayed, recorded);
+    return 1;
+  }
+
+  std::fprintf(stderr, "PASS shared stack replay=%d\n", replayed);
+  return 0;
+}
+
+// CHECK: PASS shared stack replay=14