diff --git a/docs/investigations/2026-08-host-orch-phase-tail-is-page-faults.md b/docs/investigations/2026-08-host-orch-phase-tail-is-page-faults.md
index 33ec020804..f8dc6ab44b 100644
--- a/docs/investigations/2026-08-host-orch-phase-tail-is-page-faults.md
+++ b/docs/investigations/2026-08-host-orch-phase-tail-is-page-faults.md
@@ -96,7 +96,7 @@ where the fault landed. Over the 447 faulting calls above 10 µs (867 faults, 19
 | `recording.tensor_sources`, doubles | 24 B | 70 | 174 | **16.1%** |
 | `recording.internal_fanins` (8 B) and `.predicates` (192 B), double | — | 77 | 195 | **14.7%** |
 | `recording.tensor_map` entry pool, initialized on write | 128 B | 36 | 43 | 4.5% |
-| `recording.scalars` (8 B) / `.scalar_sources` (16 B), double | — | 16 | 25 | 2.7% |
+| `recording.scalars` (8 B) / `.scalar_sources` (16 B, since renamed `.scalar_inheritance` and narrowed to 4 B), double | — | 16 | 25 | 2.7% |
 
 165 of the 447 faulting calls reallocated no recording-owned container at all, which
 leaves the per-node vector as the allocation: it is the single largest source.
diff --git a/simpler_setup/kernel_compiler.py b/simpler_setup/kernel_compiler.py
index 9cb364aa3c..45cfb99f98 100644
--- a/simpler_setup/kernel_compiler.py
+++ b/simpler_setup/kernel_compiler.py
@@ -13,6 +13,7 @@
 import subprocess
 import sys
 import tempfile
+import warnings
 from functools import cache
 from pathlib import Path
 from typing import Optional, Union
@@ -363,9 +364,24 @@ def incore_compile_cache_token(self, core_type: str) -> dict[str, object]:
         }
 
     def _run_subprocess(
-        self, cmd: list[str], label: str, error_hint: str = "Compiler not found"
+        self,
+        cmd: list[str],
+        label: str,
+        error_hint: str = "Compiler not found",
+        surface_diagnostics: bool = False,
     ) -> subprocess.CompletedProcess:
-        """Run a subprocess command with standardized logging and error handling."""
+        """Run a subprocess command with standardized logging and error handling.
+
+        surface_diagnostics reports a successful compile's stderr as a warning. It goes
+        through `warnings` rather than `logger`: pytest shows its warnings summary with no
+        flag at all, while logger output below ERROR is swallowed unless the run passes
+        --log-cli-level, and child pytest processes do not inherit that option (conftest's
+        _resource_child_command forwards only --manual). A diagnostic nobody sees by
+        default is a diagnostic that does not exist.
+
+        It is opt-in per call site rather than global because the kernel toolchains carry
+        pre-existing warnings that would bury the ones a caller turned this on to see.
+        """
         logger.debug(f"[{label}] Command: {' '.join(cmd)}")
         try:
             result = subprocess.run(cmd, check=False, capture_output=True, text=True, cwd=self.project_root)
@@ -374,6 +390,8 @@ def _run_subprocess(
                 logger.debug(f"[{label}] stdout:\n{result.stdout}")
             if result.stderr and logger.isEnabledFor(10):
                 logger.debug(f"[{label}] stderr:\n{result.stderr}")
+            if result.stderr and result.returncode == 0 and surface_diagnostics:
+                warnings.warn(f"[{label}] compiler diagnostics:\n{result.stderr}", stacklevel=2)
 
             if result.returncode != 0:
                 logger.error(f"[{label}] Compilation failed: {result.stderr}")
@@ -391,6 +409,7 @@ def _compile_to_bytes(
         label: str,
         error_hint: str = "Compiler not found",
         delete_output: bool = True,
+        surface_diagnostics: bool = False,
     ) -> bytes:
         """Run compilation command, read output file, clean up, return bytes.
 
@@ -399,6 +418,8 @@ def _compile_to_bytes(
             output_path: Path to expected output file
             label: Label for log messages
             error_hint: Message for FileNotFoundError
+            surface_diagnostics: Report a successful compile's stderr as a warning
+                (see _run_subprocess)
 
         Returns:
             Binary contents of the compiled output file
@@ -406,16 +427,20 @@ def _compile_to_bytes(
         Raises:
             RuntimeError: If compilation fails or output file not found
         """
-        self._run_subprocess(cmd, label, error_hint)
-
-        if not os.path.isfile(output_path):
-            raise RuntimeError(f"Compilation succeeded but output file not found: {output_path}")
+        # The cleanup is a finally because _run_subprocess reports diagnostics through
+        # warnings.warn, which raises under an error-level warning filter -- a compile
+        # that produced its output would otherwise leave the file behind.
+        try:
+            self._run_subprocess(cmd, label, error_hint, surface_diagnostics=surface_diagnostics)
 
-        with open(output_path, "rb") as f:
-            binary_data = f.read()
+            if not os.path.isfile(output_path):
+                raise RuntimeError(f"Compilation succeeded but output file not found: {output_path}")
 
-        if delete_output:
-            os.remove(output_path)
+            with open(output_path, "rb") as f:
+                binary_data = f.read()
+        finally:
+            if delete_output and os.path.isfile(output_path):
+                os.remove(output_path)
         logger.info(f"[{label}] Compilation {output_path} successful: {len(binary_data)} bytes")
         return binary_data
 
diff --git a/src/a2a3/runtime/host_build_graph/orchestration/arg_with_deps.h b/src/a2a3/runtime/host_build_graph/orchestration/arg_with_deps.h
index fd251ddf26..3daa25c872 100644
--- a/src/a2a3/runtime/host_build_graph/orchestration/arg_with_deps.h
+++ b/src/a2a3/runtime/host_build_graph/orchestration/arg_with_deps.h
@@ -53,9 +53,9 @@ class CoreTaskArgsWithDeps : private CoreTaskArgs {
     using CoreTaskArgs::add_output;
     using CoreTaskArgs::add_scalar;
     using CoreTaskArgs::add_scalars;
-    using CoreTaskArgs::add_scalars_i32;
-    using CoreTaskArgs::allow_early_resolve;  // speculative early-dispatch hint (getter)
-    using CoreTaskArgs::copy_scalars_from;
+    using CoreTaskArgs::add_static_scalar;
+    using CoreTaskArgs::add_static_scalars;
+    using CoreTaskArgs::allow_early_resolve;      // speculative early-dispatch hint (getter)
     using CoreTaskArgs::set_allow_early_resolve;  // speculative early-dispatch hint (setter)
     using CoreTaskArgs::set_task_timing_slot;     // selective task-timing slot (setter)
     using CoreTaskArgs::task_timing_slot;         // selective task-timing slot (getter)
diff --git a/src/a2a3/runtime/host_build_graph/orchestration/orchestration_api.h b/src/a2a3/runtime/host_build_graph/orchestration/orchestration_api.h
index 1514664172..3a0cc0c9a2 100644
--- a/src/a2a3/runtime/host_build_graph/orchestration/orchestration_api.h
+++ b/src/a2a3/runtime/host_build_graph/orchestration/orchestration_api.h
@@ -122,10 +122,10 @@ static inline TaskOutputTensors alloc_tensors(const TensorCreateInfo create_info
     for (uint32_t i = 0; i < count; i++) {
         args.add_output(create_infos[i]);
     }
-    if (args.has_error) {
+    if (args.has_error()) {
         rt->ops->report_fatal(
             rt, SIMPLER_ERROR_INVALID_ARGS, __FUNCTION__, "%s",
-            args.error_msg ? args.error_msg : "alloc_tensors failed to construct output-only Arg"
+            args.error_msg() ? args.error_msg() : "alloc_tensors failed to construct output-only Arg"
         );
         return TaskOutputTensors{};
     }
@@ -145,10 +145,10 @@ static inline TaskOutputTensors alloc_tensors(const CIs &...cis) {
     }
     CoreTaskArgs args;
     (args.add_output(cis), ...);
-    if (args.has_error) {
+    if (args.has_error()) {
         rt->ops->report_fatal(
             rt, SIMPLER_ERROR_INVALID_ARGS, __FUNCTION__, "%s",
-            args.error_msg ? args.error_msg : "alloc_tensors failed to construct output-only Arg"
+            args.error_msg() ? args.error_msg() : "alloc_tensors failed to construct output-only Arg"
         );
         return TaskOutputTensors{};
     }
@@ -431,7 +431,7 @@ static inline uint64_t rt_graph_function_id(Function function) {
 // its own boundary copy as a parameter for exactly that reason.
 template <typename Invoke>
 static inline GraphSubmitResult rt_submit_graph_impl(uint64_t graph_key, const GraphTaskArgs &args, Invoke invoke) {
-    debug_assert(!args.has_error && "Graph boundary GraphTaskArgs construction failed");
+    debug_assert(!args.has_error() && "Graph boundary GraphTaskArgs construction failed");
     debug_assert(args.tensor_count() <= GRAPH_MAX_TENSOR_ARGS && "Graph boundary exceeds the tensor limit");
     debug_assert(
         args.explicit_dep_count() == 0 && "Explicit dependencies crossing the Graph boundary are not supported"
@@ -467,9 +467,13 @@ static inline GraphSubmitResult rt_submit_graph_impl(uint64_t graph_key, const G
     const uint64_t _begun_ns = rt_orch_phase_now_ns();
     if (result.recording) {
         void *handle = result.recording_handle;
+        // The formal parameters the body reads: the entry's own copy, not `args` -- the
+        // caller only lends those for the duration of this call, and recording resolves a
+        // task slot's origin against this object's slot array.
+        const GraphTaskArgs &params = *result.params;
         // A std::function rather than a bare lambda because the pool takes it as one
         // through the ops table's void *.
-        std::function<void(GraphTaskArgs &)> job = [invoke, handle](GraphTaskArgs &record_args) mutable {
+        std::function<void(const GraphTaskArgs &)> job = [invoke, handle](const GraphTaskArgs &record_args) mutable {
             try {
                 if (!rt_graph_prepare(handle, record_args)) {
                     rt_graph_abort(handle);
@@ -491,15 +495,15 @@ static inline GraphSubmitResult rt_submit_graph_impl(uint64_t graph_key, const G
         // costs nothing here, because the fallback below re-runs `invoke` -- captured by
         // value, so unaffected -- rather than the job.
         RuntimeContext *record_rt = current_runtime();
-        const bool queued =
-            record_rt->ops->graph_record_start != nullptr && record_rt->ops->graph_record_start(record_rt, args, &job);
+        const bool queued = record_rt->ops->graph_record_start != nullptr &&
+                            record_rt->ops->graph_record_start(record_rt, params, &job);
         if (!queued) {
             try {
-                if (!rt_graph_prepare(handle, args)) {
+                if (!rt_graph_prepare(handle, params)) {
                     rt_graph_abort(handle);
                     return result;
                 }
-                invoke(args);
+                invoke(params);
                 (void)rt_graph_end();
             } catch (...) {
                 rt_graph_abort(handle);
diff --git a/src/a5/runtime/host_build_graph/orchestration/arg_with_deps.h b/src/a5/runtime/host_build_graph/orchestration/arg_with_deps.h
index fd251ddf26..3daa25c872 100644
--- a/src/a5/runtime/host_build_graph/orchestration/arg_with_deps.h
+++ b/src/a5/runtime/host_build_graph/orchestration/arg_with_deps.h
@@ -53,9 +53,9 @@ class CoreTaskArgsWithDeps : private CoreTaskArgs {
     using CoreTaskArgs::add_output;
     using CoreTaskArgs::add_scalar;
     using CoreTaskArgs::add_scalars;
-    using CoreTaskArgs::add_scalars_i32;
-    using CoreTaskArgs::allow_early_resolve;  // speculative early-dispatch hint (getter)
-    using CoreTaskArgs::copy_scalars_from;
+    using CoreTaskArgs::add_static_scalar;
+    using CoreTaskArgs::add_static_scalars;
+    using CoreTaskArgs::allow_early_resolve;      // speculative early-dispatch hint (getter)
     using CoreTaskArgs::set_allow_early_resolve;  // speculative early-dispatch hint (setter)
     using CoreTaskArgs::set_task_timing_slot;     // selective task-timing slot (setter)
     using CoreTaskArgs::task_timing_slot;         // selective task-timing slot (getter)
diff --git a/src/a5/runtime/host_build_graph/orchestration/orchestration_api.h b/src/a5/runtime/host_build_graph/orchestration/orchestration_api.h
index 1514664172..3a0cc0c9a2 100644
--- a/src/a5/runtime/host_build_graph/orchestration/orchestration_api.h
+++ b/src/a5/runtime/host_build_graph/orchestration/orchestration_api.h
@@ -122,10 +122,10 @@ static inline TaskOutputTensors alloc_tensors(const TensorCreateInfo create_info
     for (uint32_t i = 0; i < count; i++) {
         args.add_output(create_infos[i]);
     }
-    if (args.has_error) {
+    if (args.has_error()) {
         rt->ops->report_fatal(
             rt, SIMPLER_ERROR_INVALID_ARGS, __FUNCTION__, "%s",
-            args.error_msg ? args.error_msg : "alloc_tensors failed to construct output-only Arg"
+            args.error_msg() ? args.error_msg() : "alloc_tensors failed to construct output-only Arg"
         );
         return TaskOutputTensors{};
     }
@@ -145,10 +145,10 @@ static inline TaskOutputTensors alloc_tensors(const CIs &...cis) {
     }
     CoreTaskArgs args;
     (args.add_output(cis), ...);
-    if (args.has_error) {
+    if (args.has_error()) {
         rt->ops->report_fatal(
             rt, SIMPLER_ERROR_INVALID_ARGS, __FUNCTION__, "%s",
-            args.error_msg ? args.error_msg : "alloc_tensors failed to construct output-only Arg"
+            args.error_msg() ? args.error_msg() : "alloc_tensors failed to construct output-only Arg"
         );
         return TaskOutputTensors{};
     }
@@ -431,7 +431,7 @@ static inline uint64_t rt_graph_function_id(Function function) {
 // its own boundary copy as a parameter for exactly that reason.
 template <typename Invoke>
 static inline GraphSubmitResult rt_submit_graph_impl(uint64_t graph_key, const GraphTaskArgs &args, Invoke invoke) {
-    debug_assert(!args.has_error && "Graph boundary GraphTaskArgs construction failed");
+    debug_assert(!args.has_error() && "Graph boundary GraphTaskArgs construction failed");
     debug_assert(args.tensor_count() <= GRAPH_MAX_TENSOR_ARGS && "Graph boundary exceeds the tensor limit");
     debug_assert(
         args.explicit_dep_count() == 0 && "Explicit dependencies crossing the Graph boundary are not supported"
@@ -467,9 +467,13 @@ static inline GraphSubmitResult rt_submit_graph_impl(uint64_t graph_key, const G
     const uint64_t _begun_ns = rt_orch_phase_now_ns();
     if (result.recording) {
         void *handle = result.recording_handle;
+        // The formal parameters the body reads: the entry's own copy, not `args` -- the
+        // caller only lends those for the duration of this call, and recording resolves a
+        // task slot's origin against this object's slot array.
+        const GraphTaskArgs &params = *result.params;
         // A std::function rather than a bare lambda because the pool takes it as one
         // through the ops table's void *.
-        std::function<void(GraphTaskArgs &)> job = [invoke, handle](GraphTaskArgs &record_args) mutable {
+        std::function<void(const GraphTaskArgs &)> job = [invoke, handle](const GraphTaskArgs &record_args) mutable {
             try {
                 if (!rt_graph_prepare(handle, record_args)) {
                     rt_graph_abort(handle);
@@ -491,15 +495,15 @@ static inline GraphSubmitResult rt_submit_graph_impl(uint64_t graph_key, const G
         // costs nothing here, because the fallback below re-runs `invoke` -- captured by
         // value, so unaffected -- rather than the job.
         RuntimeContext *record_rt = current_runtime();
-        const bool queued =
-            record_rt->ops->graph_record_start != nullptr && record_rt->ops->graph_record_start(record_rt, args, &job);
+        const bool queued = record_rt->ops->graph_record_start != nullptr &&
+                            record_rt->ops->graph_record_start(record_rt, params, &job);
         if (!queued) {
             try {
-                if (!rt_graph_prepare(handle, args)) {
+                if (!rt_graph_prepare(handle, params)) {
                     rt_graph_abort(handle);
                     return result;
                 }
-                invoke(args);
+                invoke(params);
                 (void)rt_graph_end();
             } catch (...) {
                 rt_graph_abort(handle);
diff --git a/src/common/host_build_graph/device/graph_execution.cpp b/src/common/host_build_graph/device/graph_execution.cpp
index 7fd7169e18..aa4f138b63 100644
--- a/src/common/host_build_graph/device/graph_execution.cpp
+++ b/src/common/host_build_graph/device/graph_execution.cpp
@@ -62,8 +62,7 @@ bool bind_graph_topology(GraphExecution &execution) {
     // range overflows the increment before any bound check can see it.
     if (definition.task_count <= 0 || definition.task_count > MAX_IN_GRAPH_TASKS) return false;
     // GRAPH_MAX_SCALAR_ARGS, not MAX_SCALAR_ARGS: this counts the scalars the
-    // Graph BOUNDARY carries, which the recorder sizes with
-    // GraphTaskArgs = Arg<GRAPH_MAX_TENSOR_ARGS, GRAPH_MAX_SCALAR_ARGS> and the
+    // Graph BOUNDARY carries, which the recorder sizes with GraphTaskArgs and the
     // outer Graph payload hands it to GraphExecution, never through an in-graph task
     // payload. MAX_SCALAR_ARGS is the per-AICore-task cap (16) and applies to
     // InGraphTaskDefinition::scalar_count below, which is checked separately; using
@@ -373,18 +372,19 @@ GraphMaterializeResult graph_execution_materialize_slice(
         definition.scalar_arg_count == 0 ?
             nullptr :
             graph_definition_array<uint64_t>(definition, definition.off_scalars, definition.scalar_arg_count);
-    const GraphScalarSourceRef *scalar_sources =
-        definition.scalar_arg_count == 0 ? nullptr :
-                                           graph_definition_array<GraphScalarSourceRef>(
-                                               definition, definition.off_scalar_sources, definition.scalar_arg_count
-                                           );
+    const GraphScalarInheritance *scalar_inheritance =
+        definition.scalar_arg_count == 0 ?
+            nullptr :
+            graph_definition_array<GraphScalarInheritance>(
+                definition, definition.off_scalar_inheritance, definition.scalar_arg_count
+            );
     const GraphPredicate *predicates =
         definition.predicate_count == 0 ?
             nullptr :
             graph_definition_array<GraphPredicate>(definition, definition.off_predicates, definition.predicate_count);
     if (tasks == nullptr || in_graph_task_offsets == nullptr ||
         (definition.tensor_arg_count != 0 && (definition_tensors == nullptr || tensor_sources == nullptr)) ||
-        (definition.scalar_arg_count != 0 && (definition_scalars == nullptr || scalar_sources == nullptr)) ||
+        (definition.scalar_arg_count != 0 && (definition_scalars == nullptr || scalar_inheritance == nullptr)) ||
         (definition.predicate_count != 0 && predicates == nullptr)) {
         execution.materialize_busy.store(0, std::memory_order_release);
         return GraphMaterializeResult::INVALID;
@@ -462,18 +462,15 @@ GraphMaterializeResult graph_execution_materialize_slice(
         uint64_t *task_scalars = payload.scalar_data();
         for (int32_t j = 0; j < source.scalar_count; ++j) {
             const int32_t scalar_index = source.scalar_offset + j;
-            const GraphScalarSourceRef &ref = scalar_sources[scalar_index];
-            if (ref.source_kind == static_cast<uint8_t>(GraphScalarSourceKind::STATIC_VALUE)) {
+            const GraphScalarInheritance &ref = scalar_inheritance[scalar_index];
+            if (!ref.inherited()) {
                 task_scalars[j] = definition_scalars[scalar_index];
-            } else if (ref.source_kind == static_cast<uint8_t>(GraphScalarSourceKind::BOUNDARY)) {
-                if (ref.source_index >= execution.boundary_scalar_count || execution.boundary_scalars == nullptr) {
+            } else {
+                if (ref.boundary_index() >= execution.boundary_scalar_count || execution.boundary_scalars == nullptr) {
                     execution.materialize_busy.store(0, std::memory_order_release);
                     return GraphMaterializeResult::INVALID;
                 }
-                task_scalars[j] = execution.boundary_scalars[ref.source_index];
-            } else {
-                execution.materialize_busy.store(0, std::memory_order_release);
-                return GraphMaterializeResult::INVALID;
+                task_scalars[j] = execution.boundary_scalars[ref.boundary_index()];
             }
         }
         reset_graph_payload(payload);
diff --git a/src/common/host_build_graph/docs/GRAPH_EXECUTION.md b/src/common/host_build_graph/docs/GRAPH_EXECUTION.md
index a237c04c0f..22b917f7ed 100644
--- a/src/common/host_build_graph/docs/GRAPH_EXECUTION.md
+++ b/src/common/host_build_graph/docs/GRAPH_EXECUTION.md
@@ -29,10 +29,11 @@ TensorMap producers.
 
 ## API
 
-A Graph uses `CoreTaskArgs`, the existing incore argument type:
+A Graph boundary uses `GraphTaskArgs`; an in-graph task's arguments use
+`CoreTaskArgs`, the existing incore argument type:
 
 ```cpp
-void graph_function(const CoreTaskArgs &args, int variant) {
+void graph_function(const GraphTaskArgs &args, int variant) {
     const ChipTensor &input = args.tensor(0).ref();
     const ChipTensor &weight = args.tensor(1).ref();
     const ChipTensor &output = args.tensor(2).ref();
@@ -45,7 +46,7 @@ void graph_function(const CoreTaskArgs &args, int variant) {
     CoreTaskArgs matmul_args;
     matmul_args.add_input(input, weight);
     matmul_args.add_output(intermediate);
-    matmul_args.copy_scalars_from(args, 0, 1);  // current invocation's value
+    matmul_args.add_scalar(args.scalar(0));  // forwarded boundary parameter
     TaskOutputTensors matmul = rt_submit_aic_task(
         variant == 0 ? FUNC_MATMUL : FUNC_MATMUL_TRANSPOSED,
         matmul_args
@@ -57,7 +58,7 @@ void graph_function(const CoreTaskArgs &args, int variant) {
     rt_submit_aiv_task(FUNC_ACTIVATION, activation_args);
 }
 
-void submit_layer(const CoreTaskArgs &args) {
+void submit_layer(const GraphTaskArgs &args) {
     rt_submit_graph(&graph_function, args, /*variant=*/0);
 }
 ```
@@ -65,7 +66,7 @@ void submit_layer(const CoreTaskArgs &args) {
 The function pointer is the default Graph identity. Trailing integral,
 `float`, `double`, and `bool` construction parameters are forwarded to the
 Graph function and hashed by value into the cache key. They are separate from
-execution scalars in `CoreTaskArgs`: changing a construction parameter selects a
+execution scalars in `GraphTaskArgs`: changing a construction parameter selects a
 different Definition rather than patching an existing one.
 
 An explicit identity is available for call sites that need a stable name:
@@ -85,28 +86,70 @@ Graph function pointer from the cache identity so the key remains stable; using
 the same key for different functions can select the wrong recorded topology.
 
 There are no public `GraphArgs`, `GraphBindings`, `Patch`, or `ScalarRef`
-types. The boundary is represented by `CoreTaskArgs`.
-
-Boundary scalars are pass-through bindings. Forward them directly with
-`task_args.add_scalar(args.scalar(i))` or `copy_scalars_from(args, i, count)`
-so recording can retain their source indices.
-
-Ordinary C++ value transformations do not retain boundary provenance. Both
-`task_args.add_scalar(args.scalar(i) + 1)` and copying `args.scalar(i)` into a
-local arithmetic variable before calling `add_scalar` produce an ordinary
-static in-graph task scalar. That value is stored in the Definition, and later cache
-hits reuse the first invocation's value without a warning. The runtime cannot
-distinguish such a derived value from an intentional static literal after the
-C++ expression has produced a plain arithmetic value. Compute the derived value
-before constructing the Graph boundary and pass it as another boundary scalar,
-perform the transformation in a kernel, or use a construction parameter when
-the value changes the Graph structure.
-
-Access through a non-const `scalar()` invalidates inherited boundary provenance
-conservatively, because returning a mutable reference cannot distinguish a read
-from a later write. A Graph containing such an invalidated binding is not
-cached, which prevents replay from silently replacing the transformed value
-with the unmodified boundary value.
+types. The boundary is represented by `GraphTaskArgs`, which a Graph function
+receives as `const GraphTaskArgs &`. It is sized independently of
+`CoreTaskArgs`, and forwarding a boundary scalar into a task's `CoreTaskArgs`
+crosses those two capacities without either naming the other: what
+`args.scalar(i)` hands out identifies a parameter, not the `Arg` holding it.
+
+Boundary scalars are formal parameters. `args.scalar(i)` answers parameter `i`
+itself rather than its value, so forwarding it —
+`task_args.add_scalar(args.scalar(i))` — makes the destination slot follow that
+parameter on every replay. A slot names the parameter it came from, not the
+`Arg` it was copied through, so provenance survives any number of intermediate
+copies.
+
+Reading a parameter as a value freezes it. `uint64_t v = args.scalar(i)` and
+`static_cast<int32_t>(args.scalar(i))` both convert through a **deprecated**
+operator, so the compiler names the file and line: the destination slot becomes
+static Definition data holding the recording invocation's number, and later
+cache hits replay that number. Forwarding never reaches that operator, which is
+what keeps a correct pass-through silent.
+
+That diagnostic has one limit. GCC suppresses a deprecation instantiated inside
+a system header, so a value read whose conversion happens in third-party
+template code stays silent — `EXPECT_EQ(args.scalar(i), v)` is the case found so
+far. The warning is an inventory of the value reads written here, not a proof
+that none exists.
+
+When a value read is what you meant, say so with `args.scalar(i).to<T>()`. It
+applies `to_u64`'s actual inverse, which `static_cast` is not — a float slot
+holds a bit pattern, so `static_cast<float>` of `1.0f`'s pattern yields
+`1065353216.0`. It is also the only spelling that reaches an enum:
+`static_cast<DataType>(args.scalar(i))` does not compile, because a conversion
+to an enumeration does not accept a user-defined one on the way.
+
+Freezing on purpose has a second spelling, and the two do different things.
+`args.scalar(i).to<T>()` hands the body a `T` to compute with, and whatever the
+body does with it afterwards is ordinary host code.
+`task_args.add_static_scalar(args.scalar(i))` instead forwards the parameter
+into a slot and drops its origin: the slot carries the same bit pattern a
+forward would have, but is recorded as static Definition data rather than
+following the parameter. Reach for the first when the body needs the number, the
+second when a destination — typically a nested Graph's boundary — should hold
+the value the enclosing parameter had at record time.
+
+A derived value (`args.scalar(i) + 1`) is a value read and freezes the same way.
+Compute it before constructing the boundary and pass it as its own parameter,
+perform the transformation in a kernel, or use a construction parameter when the
+value changes the Graph's structure.
+
+Boundary scalar slots are read-only: `scalar()` hands out the parameter, not a
+mutable reference, so a binding cannot be overwritten after it is forwarded.
+
+**Only a parameter of the Graph's own boundary is refreshed on replay.** The
+Definition's scalar source refs index that boundary and nothing else, so a slot
+that inherits anything else — a slot of some other `GraphTaskArgs`, or one built
+inside the body — is static Definition data holding the value it resolved to at
+record time. The runtime does not reject that; which slot a body inherits from
+is the author's declaration, and this is the declared consequence. Two notes on
+why it cannot be diagnosed instead:
+
+- An address cannot tell "created inside this body" from "created outside it".
+  The body's `Arg`s are stack locals while the boundary is pool storage, so
+  their relative addresses are a platform accident, not a guarantee.
+- Whether the outside slot's value changes between invocations is invisible
+  here. If it does, the Definition keeps replaying the recorded one.
 
 ## Supported dynamic and static data
 
@@ -143,11 +186,11 @@ path remains the defensive release-build behavior.
 
 ## Qwen decoder-layer example
 
-The upper layer packages all ChipTensor I/O in `CoreTaskArgs`; the wrapper has no
+The upper layer packages all ChipTensor I/O in `GraphTaskArgs`; the wrapper has no
 separate `hidden`, `weight`, or `output` parameters:
 
 ```cpp
-void qwen_decoder_layer(const CoreTaskArgs &args) {
+void qwen_decoder_layer(const GraphTaskArgs &args) {
     const ChipTensor &hidden = args.tensor(0).ref();
     const ChipTensor &attention_weight = args.tensor(1).ref();
     const ChipTensor &mlp_weight = args.tensor(2).ref();
@@ -161,7 +204,7 @@ void qwen_decoder_layer(const CoreTaskArgs &args) {
     CoreTaskArgs attention_args;
     attention_args.add_input(hidden, attention_weight);
     attention_args.add_output(attention_out);
-    attention_args.copy_scalars_from(args, 0, 1);  // dynamic token position
+    attention_args.add_scalar(args.scalar(0));  // dynamic token position
     TaskOutputTensors attention =
         rt_submit_aic_task(FUNC_ATTENTION, attention_args);
 
@@ -175,7 +218,7 @@ void qwen_decoder_layer(const CoreTaskArgs &args) {
     rt_submit_task(mlp, mlp_args);
 }
 
-void submit_qwen_decoder_layer(const CoreTaskArgs &args) {
+void submit_qwen_decoder_layer(const GraphTaskArgs &args) {
     rt_submit_graph(&qwen_decoder_layer, args);
 }
 
@@ -187,7 +230,7 @@ void decode_three_layers(
     const std::array<uint32_t, 3> &token_position
 ) {
     for (std::size_t layer = 0; layer < hidden.size(); ++layer) {
-        CoreTaskArgs args;
+        GraphTaskArgs args;
         args.add_input(
             hidden[layer],
             attention_weight[layer],
@@ -508,7 +551,6 @@ builds:
   predicate's operand tensor;
 - a dispatch predicate whose operand is the predicated in-graph task's own output;
 - a dispatch predicate whose index vector leaves the operand tensor's extent;
-- a boundary-derived scalar accessed through mutable `scalar()`;
 - runtime allocation inside the Graph body;
 - more than 1024 in-graph tasks;
 - insufficient heap capacity while deferred shells are finalized.
diff --git a/src/common/host_build_graph/graph_cache.h b/src/common/host_build_graph/graph_cache.h
index b02da42db9..c20edaa318 100644
--- a/src/common/host_build_graph/graph_cache.h
+++ b/src/common/host_build_graph/graph_cache.h
@@ -29,6 +29,20 @@ struct GraphScopeResult {
     // mean traversing the in-flight map under its mutex on a path whose whole
     // purpose is to not contend with the submitting thread.
     void *recording_handle{nullptr};
+    // The formal parameters the recorded body must read: the entry's own deep copy, not
+    // the caller's arguments, which are only lent for the duration of the submit call.
+    // Set exactly when `recording` is, and valid until graph_commit drains the entry --
+    // which is after every job that could read it has finished.
+    //
+    // Const because the boundary is written once, by the submitting thread in graph_begin,
+    // and only read after that -- by the recorder, and by later same-key submissions
+    // comparing against it. Nothing may write it once it is handed out here.
+    //
+    // The body must read these and not the caller's arguments even on the synchronous
+    // fallback path: recording resolves a task slot's origin against this object's slot
+    // array, so an origin from any other object falls outside it and the parameter is
+    // recorded as static.
+    const GraphTaskArgs *params{nullptr};
 };
 
 using GraphSubmitResult = GraphScopeResult;
@@ -66,14 +80,9 @@ constexpr uint64_t graph_const_hash_impl(const char *s, uint64_t h) {
 constexpr uint64_t GRAPH_KEY(const char *s) { return graph_const_hash_impl(s, 1469598103934665603ULL); }
 
 inline bool rt_graph_args_cacheable(const GraphTaskArgs &args) {
-    if (args.has_error || args.tensor_count() <= 0 || args.tensor_count() > GRAPH_MAX_TENSOR_ARGS) {
+    if (args.has_error() || args.tensor_count() <= 0 || args.tensor_count() > GRAPH_MAX_TENSOR_ARGS) {
         return false;
     }
-    for (int32_t i = 0; i < args.tensor_count(); ++i) {
-        // A Graph boundary is caller-owned storage. Runtime-allocated
-        // TensorCreateInfo outputs remain on the ordinary submit path.
-        if (args.tag(i) == TensorArgType::OUTPUT) return false;
-    }
     return true;
 }
 
diff --git a/src/common/host_build_graph/graph_execution.h b/src/common/host_build_graph/graph_execution.h
index 3ef220e56f..2700af4a37 100644
--- a/src/common/host_build_graph/graph_execution.h
+++ b/src/common/host_build_graph/graph_execution.h
@@ -93,15 +93,41 @@ struct GraphTensorSourceRef {
     uint64_t packed_offset;
 };
 
-enum class GraphScalarSourceKind : uint8_t {
-    STATIC_VALUE = 0,
-    BOUNDARY = 1,
-};
-
-struct GraphScalarSourceRef {
-    uint16_t source_index;
-    uint8_t source_kind;
-    uint8_t reserved;
+// Where one in-graph task scalar slot takes its value from: the Definition's own
+// scalars[] entry, or the boundary parameter named by boundary_index(). It is the wire
+// form of the two things recording knows about a slot -- whether it inherits, and which
+// parameter it inherits -- so inherited() is the same predicate as Arg::scalar_inherited.
+//
+// Only a parameter of the replaying Graph's boundary can be refreshed; the index reaches
+// that boundary and nothing else, and means nothing while inherited() is false. The
+// fields are private so the pair can only be set together, through a factory that decides
+// both: an entry carrying an index while claiming not to inherit, or the reverse, cannot
+// be spelled. What the index means is still a claim about a boundary this entry cannot
+// see, so the readers bound it against the boundary they do have.
+class GraphScalarInheritance {
+public:
+    // The image's scalar section is allocated as an array, so a default-constructible
+    // slot is required; the factories below are what a caller fills one with.
+    GraphScalarInheritance() = default;
+
+    static GraphScalarInheritance self_value() { return {0, false}; }
+    static GraphScalarInheritance from_boundary(uint16_t index) { return {index, true}; }
+
+    bool inherited() const { return inherited_ != 0; }
+    uint16_t boundary_index() const { return boundary_index_; }
+
+private:
+    GraphScalarInheritance(uint16_t index, bool inherits) :
+        boundary_index_(index),
+        inherited_(inherits ? 1 : 0) {}
+
+    // One access level for every field, which is what keeps this standard-layout and so
+    // safe to memcpy to the device. inherited_ is a uint16_t rather than a bool so the
+    // two fields fill the size alignof(uint16_t) rounds this type up to: there is no
+    // padding byte, and so no indeterminate byte in the image's scalar section, which the
+    // static_assert below pins the size of because that is what the section indexes by.
+    uint16_t boundary_index_;
+    uint16_t inherited_;
 };
 
 // Wire representation of an in-graph task's dispatch predicate. The operand's absolute GM
@@ -228,7 +254,7 @@ struct GraphDefinition {
     uint32_t off_tensors;
     uint32_t off_tensor_sources;
     uint32_t off_scalars;
-    uint32_t off_scalar_sources;
+    uint32_t off_scalar_inheritance;
     uint32_t off_boundary_signatures;
     uint32_t off_predicates;
 };
@@ -237,8 +263,9 @@ static_assert(std::is_trivially_copyable_v<GraphTensorSourceRef>);
 static_assert(std::is_standard_layout_v<GraphTensorSourceRef>);
 static_assert(std::is_trivially_copyable_v<GraphTensor>);
 static_assert(std::is_standard_layout_v<GraphTensor>);
-static_assert(std::is_trivially_copyable_v<GraphScalarSourceRef>);
-static_assert(std::is_standard_layout_v<GraphScalarSourceRef>);
+static_assert(std::is_trivially_copyable_v<GraphScalarInheritance>);
+static_assert(std::is_standard_layout_v<GraphScalarInheritance>);
+static_assert(sizeof(GraphScalarInheritance) == 4, "the image's scalar section assumes this layout");
 static_assert(std::is_trivially_copyable_v<InGraphTaskDefinition>);
 static_assert(std::is_standard_layout_v<InGraphTaskDefinition>);
 // graph_fill_definition assigns this struct field by field, so its interior padding
@@ -262,7 +289,7 @@ static_assert(std::is_standard_layout_v<GraphDefinition>);
 static_assert(
     alignof(InGraphTaskDefinition) <= alignof(std::max_align_t) && alignof(GraphTensor) <= alignof(std::max_align_t) &&
         alignof(GraphTensorSourceRef) <= alignof(std::max_align_t) &&
-        alignof(GraphScalarSourceRef) <= alignof(std::max_align_t) &&
+        alignof(GraphScalarInheritance) <= alignof(std::max_align_t) &&
         alignof(GraphBoundarySignature) <= alignof(std::max_align_t) &&
         alignof(GraphPredicate) <= alignof(std::max_align_t),
     "a Definition section type must not be over-aligned: its storage is a byte vector"
diff --git a/src/common/host_build_graph/graph_recorder_pool.h b/src/common/host_build_graph/graph_recorder_pool.h
index 20033ae226..95adc67610 100644
--- a/src/common/host_build_graph/graph_recorder_pool.h
+++ b/src/common/host_build_graph/graph_recorder_pool.h
@@ -11,20 +11,18 @@
 #pragma once
 
 /**
- * The bounded pool of threads that record Graph bodies, and the copy of a boundary's
- * arguments each queued recording owns.
+ * The bounded pool of threads that record Graph bodies.
  *
  * Runtime-owned, and only in the host target: one pool per process, serving every
  * registered callable. It used to live in orchestration_api.h, which put a
- * function-local static — and therefore a pool, eight threads and their recording
- * storage — inside every orchestration .so. Those are dlopen'd one per callable with
- * RTLD_LOCAL and are not released as cases finish, so a process held one pool per
- * registered callable: measured at 3 concurrently mapped orchestration images over a
- * four-case pytest session and 5 over the a2a3 host_build_graph corpus, i.e. 24-40
- * recorder threads where 8 suffice.
+ * function-local static — and therefore a pool and eight threads — inside every
+ * orchestration .so. Those are dlopen'd one per callable with RTLD_LOCAL and are not
+ * released as cases finish, so a process held one pool per registered callable: measured
+ * at 3 concurrently mapped orchestration images over a four-case pytest session and 5
+ * over the a2a3 host_build_graph corpus, i.e. 24-40 recorder threads where 8 suffice.
  *
  * A worker runs jobs the orchestration .so builds (each captures that .so's generated
- * orchestration function), reached through the ops table. Two properties make that
+ * orchestration function), reached through the ops table. Three properties make that
  * sound and are relied on here:
  *
  *   - No job outlives the bind that queued it. rt_orchestration_done() ->
@@ -35,6 +33,13 @@
  *     (orchestration/common.cpp, deliberately not thread_local), so a worker shared
  *     across callables reads the right one: the job's own inlined code reads its own
  *     .so's global.
+ *   - **The boundary a job reads is owned by the in-flight entry, not by this pool.**
+ *     start() forwards that reference rather than copying the arguments, so the entry
+ *     must outlive the job. It does: graph_end only marks the entry, and graph_commit
+ *     frees it only after waiting for every recording to finish. Teardown does not
+ *     escape that either -- shutdown() drains the queue before it sets stopping_, so a
+ *     worker's stopping_ exit always finds an empty queue and no job runs after the
+ *     entry it would read.
  */
 
 #include <pthread.h>
@@ -54,62 +59,9 @@
 #include "graph_host_state.h"        // GRAPH_MAX_DEFINITIONS
 #include "host_build_graph/types.h"  // GraphTaskArgs, GRAPH_MAX_{TENSOR,SCALAR}_ARGS
 
-class GraphOwnedArgs {
-public:
-    GraphOwnedArgs() { std::memset(tensors_.data(), 0, sizeof(tensors_)); }
-
-    // The arrays below are sized to the Graph boundary's own capacity, so a
-    // source GraphTaskArgs cannot report more args than they hold and the copy
-    // loops need no runtime bound.
-    void assign(const GraphTaskArgs &source) {
-        args_.reset();
-        for (int32_t i = 0; i < source.tensor_count(); ++i) {
-            tensors_[static_cast<size_t>(i)].copy(source.tensor(i).ref());
-            switch (source.tag(i)) {
-            case TensorArgType::INPUT:
-                args_.add_input(tensors_[static_cast<size_t>(i)]);
-                break;
-            case TensorArgType::OUTPUT_EXISTING:
-                args_.add_output(tensors_[static_cast<size_t>(i)]);
-                break;
-            case TensorArgType::INOUT:
-                args_.add_inout(tensors_[static_cast<size_t>(i)]);
-                break;
-            case TensorArgType::NO_DEP:
-                args_.add_no_dep(tensors_[static_cast<size_t>(i)]);
-                break;
-            case TensorArgType::OUTPUT:
-                args_.set_error("Runtime-allocated output is not supported at a Graph boundary");
-                break;
-            }
-        }
-        for (int32_t i = 0; i < source.scalar_count(); ++i) {
-            scalars_[static_cast<size_t>(i)] = source.scalar(i);
-            args_.add_scalar(scalars_[static_cast<size_t>(i)]);
-        }
-        args_.launch_spec = source.launch_spec;
-        args_.set_allow_early_resolve(source.allow_early_resolve());
-        if (source.task_timing_slot() != TASK_TIMING_SLOT_NONE) {
-            args_.set_task_timing_slot(source.task_timing_slot());
-        }
-        args_.set_predicate(source.predicate());
-    }
-
-    GraphTaskArgs &args() { return args_; }
-
-private:
-    std::array<simpler::hbg::Tensor, GRAPH_MAX_TENSOR_ARGS> tensors_{};
-    std::array<uint64_t, GRAPH_MAX_SCALAR_ARGS> scalars_{};
-    GraphTaskArgs args_;
-};
-
 class GraphAsyncRecordingState {
 public:
-    GraphAsyncRecordingState() {
-        for (size_t i = 0; i < kJobCapacity; ++i) {
-            free_owned_args_[i] = kJobCapacity - i - 1;
-        }
-    }
+    GraphAsyncRecordingState() = default;
     ~GraphAsyncRecordingState() { shutdown(); }
 
     GraphAsyncRecordingState(const GraphAsyncRecordingState &) = delete;
@@ -132,31 +84,24 @@ class GraphAsyncRecordingState {
         return !stopping_ && !storage_failed_ && target == kPrewarmedWorkerCount;
     }
 
+    // `args` is the in-flight entry's own boundary and is only forwarded, never copied:
+    // the entry outlives every job that reads it, because graph_commit drains this pool
+    // before freeing one. Const because that boundary is shared -- the submitting thread
+    // compares later same-key submissions against it while a worker records.
     template <typename Job>
     bool start(const GraphTaskArgs &args, Job &&job) {
-        std::function<void(GraphTaskArgs &)> next;
+        std::function<void(const GraphTaskArgs &)> next;
         try {
             next = std::forward<Job>(job);
         } catch (...) {
             return false;
         }
 
-        size_t owned_args_index;
-        {
-            std::scoped_lock lock(mutex_);
-            if (stopping_ || free_owned_args_count_ == 0 || job_count_ == kJobCapacity) return false;
-            owned_args_index = free_owned_args_[--free_owned_args_count_];
-        }
-        owned_args_[owned_args_index].assign(args);
-
         std::unique_lock<std::mutex> lock(mutex_);
-        if (stopping_) {
-            free_owned_args_[free_owned_args_count_++] = owned_args_index;
-            return false;
-        }
+        if (stopping_ || job_count_ == kJobCapacity) return false;
         PendingJob &pending = jobs_[job_tail_];
         pending.function = std::move(next);
-        pending.owned_args_index = owned_args_index;
+        pending.args = &args;
         job_tail_ = (job_tail_ + 1) % kJobCapacity;
         job_count_++;
         const size_t desired_workers = std::min(kMaxWorkerCount, job_count_ + active_jobs_);
@@ -167,8 +112,8 @@ class GraphAsyncRecordingState {
             job_tail_ = (job_tail_ + kJobCapacity - 1) % kJobCapacity;
             PendingJob &rollback = jobs_[job_tail_];
             rollback.function = {};
+            rollback.args = nullptr;
             job_count_--;
-            free_owned_args_[free_owned_args_count_++] = rollback.owned_args_index;
             return false;
         }
         lock.unlock();
@@ -235,8 +180,9 @@ class GraphAsyncRecordingState {
     static constexpr size_t kJobCapacity = GRAPH_MAX_DEFINITIONS;
 
     struct PendingJob {
-        std::function<void(GraphTaskArgs &)> function;
-        size_t owned_args_index{0};
+        std::function<void(const GraphTaskArgs &)> function;
+        // The in-flight entry's boundary. Borrowed, not owned -- see start().
+        const GraphTaskArgs *args{nullptr};
     };
 
     bool is_worker_thread_locked(std::thread::id id) const {
@@ -288,16 +234,16 @@ class GraphAsyncRecordingState {
                 if (stopping_ && job_count_ == 0) return;
                 PendingJob &pending = jobs_[job_head_];
                 current.function = std::move(pending.function);
-                current.owned_args_index = pending.owned_args_index;
+                current.args = pending.args;
+                pending.args = nullptr;
                 job_head_ = (job_head_ + 1) % kJobCapacity;
                 job_count_--;
                 active_jobs_++;
             }
-            current.function(owned_args_[current.owned_args_index].args());
+            current.function(*current.args);
             current.function = {};
             {
                 std::scoped_lock lock(mutex_);
-                free_owned_args_[free_owned_args_count_++] = current.owned_args_index;
                 active_jobs_--;
             }
             cv_.notify_all();
@@ -326,10 +272,7 @@ class GraphAsyncRecordingState {
     std::vector<std::thread> workers_;
     std::mutex mutex_;
     std::condition_variable cv_;
-    std::array<GraphOwnedArgs, kJobCapacity> owned_args_;
-    std::array<size_t, kJobCapacity> free_owned_args_{};
     std::array<PendingJob, kJobCapacity> jobs_;
-    size_t free_owned_args_count_{kJobCapacity};
     size_t job_head_{0};
     size_t job_tail_{0};
     size_t job_count_{0};
diff --git a/src/common/host_build_graph/host/graph_recorder_pool.cpp b/src/common/host_build_graph/host/graph_recorder_pool.cpp
index a92c9168c4..3b386d7b8f 100644
--- a/src/common/host_build_graph/host/graph_recorder_pool.cpp
+++ b/src/common/host_build_graph/host/graph_recorder_pool.cpp
@@ -32,9 +32,13 @@ bool graph_recorder_prewarm() { return graph_recorder_pool().prewarm(); }
 // whether or not it queues it -- start() takes the callable before it checks capacity --
 // so the caller must treat it as spent on return. No ownership crosses the .so boundary
 // either way: the caller's std::function destructs on its own side.
+//
+// `args` is the in-flight entry's own boundary, which outlives every job that reads it
+// (graph_commit drains the pool before freeing an entry), so the pool forwards it by
+// reference rather than copying it.
 bool graph_record_start_impl(RuntimeContext *, const GraphTaskArgs &args, void *job) {
     if (job == nullptr) return false;
-    auto *record = static_cast<std::function<void(GraphTaskArgs &)> *>(job);
+    auto *record = static_cast<std::function<void(const GraphTaskArgs &)> *>(job);
     return graph_recorder_pool().start(args, std::move(*record));
 }
 
diff --git a/src/common/host_build_graph/host/orchestrator.cpp b/src/common/host_build_graph/host/orchestrator.cpp
index 0fed8469d4..c3e728e021 100644
--- a/src/common/host_build_graph/host/orchestrator.cpp
+++ b/src/common/host_build_graph/host/orchestrator.cpp
@@ -266,17 +266,6 @@ struct GraphRecordedTensorSourceRef {
     uint64_t packed_offset{0};
 };
 
-enum class GraphRecordedScalarSourceKind : uint8_t {
-    STATIC_VALUE,
-    BOUNDARY,
-    INVALIDATED_BOUNDARY,
-};
-
-struct GraphRecordedScalarSourceRef {
-    GraphRecordedScalarSourceKind source_kind{GraphRecordedScalarSourceKind::STATIC_VALUE};
-    size_t source_index{0};
-};
-
 // A recorded task's dispatch predicate, held as the operand tensor plus the element
 // index within it rather than the absolute address submit would resolve. The tensor is
 // copied because the caller only lends it for the duration of the submit call.
@@ -361,17 +350,34 @@ struct GraphBoundaryByAddress {
     uint32_t index;
 };
 
-// The Graph boundary as the submitting thread captured it, deep-copied because the
-// caller only lends its arguments for the duration of the submit call. It anchors
-// boundary scalar sources while the main thread submits outer shells from later
-// invocation arguments, and later same-key submissions compare against it under
-// recording_mutex — so it belongs to the in-flight entry, not to the recorder's
-// storage, which the submitting thread must never touch.
+// The Graph boundary as the submitting thread captured it, deep-copied because the caller
+// only lends its arguments for the duration of the submit call.
+//
+// This is the boundary: the recorded body reads it, later same-key submissions compare
+// against it, and a task slot that follows one of its parameters names a slot in `args`.
+// Everything is written once, by the submitting thread in graph_begin, and only read
+// afterwards -- the recorder never writes here.
+//
+// `tensors` is the storage `args` points into: a TensorRef holds a Tensor*, so holding a
+// GraphTaskArgs is not the same as owning its tensor data. The array is fixed-size so
+// those pointers cannot be invalidated by a reallocation.
 struct GraphBoundary {
-    const GraphTaskArgs *args{nullptr};
-    int32_t scalar_count{0};
-    std::vector<simpler::hbg::Tensor> tensors;
-    std::vector<TensorArgType> types;
+    // Deliberately user-provided rather than `= default`: a defaulted constructor here is
+    // trivial, so make_unique<GraphInflightRecording>()'s value-initialization would zero
+    // all 13.3 KB of `tensors` on the submitting thread. A user-provided one leaves the
+    // array default-initialized -- Tensor is trivially default constructible, so those
+    // pages cost nothing until a Graph writes the tensors it actually has.
+    // NOLINTNEXTLINE(modernize-use-equals-default)
+    GraphBoundary() {}
+
+    std::array<simpler::hbg::Tensor, GRAPH_MAX_TENSOR_ARGS> tensors;
+    // A dynamic parameter here names itself, a static one names nothing -- so this list is
+    // the basis recording resolves against: scalar(i) folds to &params.scalars_[i] either
+    // way, and graph_classify_scalars turns that into the index i. A parameter that named
+    // the caller's variable instead would hand out an address outside this array, the task
+    // slot following it would be recorded as static, and it would silently stop being
+    // refreshed on replay.
+    GraphTaskArgs params;
 };
 
 // Storage for one recorded body, owned by the recorder thread and reset per
@@ -409,7 +415,9 @@ struct GraphRecording {
     // growth either.
     std::vector<GraphRecordedTensorSourceRef> tensor_sources;
     std::vector<uint64_t> scalars;
-    std::vector<GraphRecordedScalarSourceRef> scalar_sources;
+    // The wire form directly: recording resolves an inherited slot into a boundary
+    // parameter index as it classifies, so there is no host-side kind left to translate.
+    std::vector<GraphScalarInheritance> scalar_inheritance;
     std::vector<int32_t> internal_fanins;
     // Indexed by RecordedInGraphTask::predicate_index; only predicated tasks
     // contribute an entry.
@@ -451,10 +459,14 @@ struct GraphRecording {
         return task_tensor_pool.get() + task.tensor_offset;
     }
 
-    const GraphTaskArgs *boundary_args() const { return boundary == nullptr ? nullptr : boundary->args; }
-    int32_t boundary_scalar_count() const { return boundary == nullptr ? 0 : boundary->scalar_count; }
-    const std::vector<simpler::hbg::Tensor> &boundary_tensors() const { return boundary->tensors; }
-    const std::vector<TensorArgType> &boundary_types() const { return boundary->types; }
+    // The entry's boundary this recording is bound to. graph_prepare binds it and
+    // graph_end/graph_abort clears it, so every read below sits between those two points:
+    // a body only runs once graph_prepare has succeeded, and graph_layout_definition runs
+    // before the unbind.
+    const GraphBoundary &bound_boundary() const {
+        debug_assert(boundary != nullptr && "a recording reads its boundary only while bound");
+        return *boundary;
+    }
 };
 
 struct GraphPendingUpload {
@@ -599,7 +611,7 @@ bool graph_tensor_from_boundary(
     // exact one may still lie ahead of it, so the earliest view is carried to the end of
     // the run and used only if none was found. Entries of one address are in boundary
     // order, so "earliest" is the same tensor a walk over the whole boundary would pick.
-    const size_t boundary_count = recording.boundary_tensors().size();
+    const size_t boundary_count = static_cast<size_t>(recording.bound_boundary().params.tensor_count());
     size_t view_index = boundary_count;
     uint64_t view_offset = 0;
     const auto &by_address = recording.boundary_by_address;
@@ -610,17 +622,18 @@ bool graph_tensor_from_boundary(
     );
     for (; entry != by_address.end() && entry->addr == tensor.buffer.addr; ++entry) {
         const size_t i = entry->index;
-        const simpler::hbg::Tensor &boundary = recording.boundary_tensors()[i];
-        if (tensor.buffer.size != boundary.buffer.size) continue;
-        if (graph_tensor_exact(tensor, boundary)) {
+        const simpler::hbg::Tensor &param_tensor =
+            recording.bound_boundary().params.tensor(static_cast<int32_t>(i)).ref();
+        if (tensor.buffer.size != param_tensor.buffer.size) continue;
+        if (graph_tensor_exact(tensor, param_tensor)) {
             source->source_kind = GraphRecordedTensorSourceKind::BOUNDARY_EXACT;
             source->source_index = i;
             source->packed_offset = 0;
             return true;
         }
-        if (view_index == boundary_count && tensor.start_offset >= boundary.start_offset) {
+        if (view_index == boundary_count && tensor.start_offset >= param_tensor.start_offset) {
             view_index = i;
-            view_offset = tensor.start_offset - boundary.start_offset;
+            view_offset = tensor.start_offset - param_tensor.start_offset;
         }
     }
     if (view_index == boundary_count) return false;
@@ -630,32 +643,39 @@ bool graph_tensor_from_boundary(
     return true;
 }
 
+// Turn one task's scalar slots into wire source refs. A slot is BOUNDARY exactly when it
+// inherits a parameter of THIS recording's boundary, which the subtraction below both
+// converts to an index and proves; everything else is static Definition data.
+//
+// Recording is where this must happen: an origin is only valid while the body runs, and
+// the Definition it feeds has to be position-independent.
 template <typename ArgT>
-GraphRecordedScalarSourceRef
-graph_classify_scalar(const GraphRecording &recording, const ArgT &args, int32_t scalar_index) {
-    if (recording.boundary_args() == nullptr) return {};
-    // Identity, not type: an in-graph task's Arg and the boundary Arg have
-    // different capacities, so compare the addresses through void.
-    if (static_cast<const void *>(&args) == static_cast<const void *>(recording.boundary_args()) &&
-        scalar_index < recording.boundary_args()->scalar_count()) {
-        return GraphRecordedScalarSourceRef{GraphRecordedScalarSourceKind::BOUNDARY, static_cast<size_t>(scalar_index)};
-    }
-
-    const void *source = args.scalar_source(scalar_index);
-    const void *invalidated_source = args.invalidated_scalar_source(scalar_index);
-    if (source == nullptr && invalidated_source == nullptr) return {};
-    for (int32_t i = 0; i < recording.boundary_args()->scalar_count(); ++i) {
-        const void *boundary_source = static_cast<const void *>(&recording.boundary_args()->scalar(i));
-        if (source == boundary_source) {
-            return GraphRecordedScalarSourceRef{GraphRecordedScalarSourceKind::BOUNDARY, static_cast<size_t>(i)};
+void graph_classify_scalars(GraphRecording &recording, const ArgT &args, int32_t scalar_offset) {
+    const GraphTaskArgs &params = recording.bound_boundary().params;
+    const uintptr_t base = reinterpret_cast<uintptr_t>(params.scalar_slot_base());
+    const uintptr_t span = static_cast<uintptr_t>(params.scalar_count()) * sizeof(uint64_t);
+    for (int32_t i = 0; i < args.scalar_count(); ++i) {
+        GraphScalarInheritance &ref = recording.scalar_inheritance[scalar_offset + i];
+        if (!args.scalar_dynamic(i)) {
+            ref = GraphScalarInheritance::self_value();
+            continue;
         }
-        if (invalidated_source == boundary_source) {
-            return GraphRecordedScalarSourceRef{
-                GraphRecordedScalarSourceKind::INVALIDATED_BOUNDARY, static_cast<size_t>(i)
-            };
+        // Integer arithmetic, not pointer comparison: relational operators on pointers are
+        // only defined within one array object, and an origin outside this boundary's slot
+        // array is admitted here. The origin is never dereferenced -- it may already point
+        // at a caller local that has gone out of scope.
+        const uintptr_t origin = reinterpret_cast<uintptr_t>(args.scalar_origin(i));
+        if (origin < base || origin - base >= span || (origin - base) % sizeof(uint64_t) != 0) {
+            // A parameter whose origin is not one of this boundary's own is static
+            // Definition data: the Definition's inheritance entries index this boundary
+            // alone, so no other slot can be refreshed on replay, and the value the handle
+            // already carried into this slot is what the image should hold.
+            // GRAPH_EXECUTION.md states this as part of the boundary-scalar contract.
+            ref = GraphScalarInheritance::self_value();
+            continue;
         }
+        ref = GraphScalarInheritance::from_boundary(static_cast<uint16_t>((origin - base) / sizeof(uint64_t)));
     }
-    return {};
 }
 
 // Entry capacity for one recorded body's hazard map. A Definition is capped at
@@ -766,7 +786,7 @@ bool graph_recording_reserve_storage(GraphRecording &recording) {
     recording.tasks.resize(kInGraphTaskCap);
     recording.tensor_sources.reserve(kInGraphTaskCap * static_cast<size_t>(CORE_MAX_TENSOR_ARGS));
     recording.scalars.reserve(kInGraphTaskCap * static_cast<size_t>(CORE_MAX_SCALAR_ARGS));
-    recording.scalar_sources.reserve(kInGraphTaskCap * static_cast<size_t>(CORE_MAX_SCALAR_ARGS));
+    recording.scalar_inheritance.reserve(kInGraphTaskCap * static_cast<size_t>(CORE_MAX_SCALAR_ARGS));
     recording.predicates.reserve(kInGraphTaskCap);
     return true;
 }
@@ -826,7 +846,7 @@ bool graph_recording_reset(GraphRecording &recording, const GraphInflightRecordi
     recording.task_tensor_cursor = 0;
     recording.tensor_sources.clear();
     recording.scalars.clear();
-    recording.scalar_sources.clear();
+    recording.scalar_inheritance.clear();
     recording.internal_fanins.clear();
     recording.predicates.clear();
     // Built once here rather than per lookup: a body classifies every tensor argument of
@@ -834,10 +854,10 @@ bool graph_recording_reset(GraphRecording &recording, const GraphInflightRecordi
     // equal addresses in boundary order, which is what makes the walk in
     // graph_tensor_from_boundary pick the same tensor the boundary order would.
     recording.boundary_by_address.clear();
-    recording.boundary_by_address.reserve(recording.boundary_tensors().size());
-    for (size_t i = 0; i < recording.boundary_tensors().size(); ++i) {
+    recording.boundary_by_address.reserve(static_cast<size_t>(recording.bound_boundary().params.tensor_count()));
+    for (int32_t i = 0; i < recording.bound_boundary().params.tensor_count(); ++i) {
         recording.boundary_by_address.push_back(
-            {recording.boundary_tensors()[i].buffer.addr, static_cast<uint32_t>(i)}
+            {recording.bound_boundary().params.tensor(i).ref().buffer.addr, static_cast<uint32_t>(i)}
         );
     }
     std::stable_sort(
@@ -952,19 +972,6 @@ std::optional<GraphTensorSourceRef> graph_pack_tensor_source(const GraphRecorded
     return packed;
 }
 
-std::optional<GraphScalarSourceRef> graph_pack_scalar_source(const GraphRecordedScalarSourceRef &source) {
-    if (source.source_kind == GraphRecordedScalarSourceKind::INVALIDATED_BOUNDARY || source.source_index > UINT16_MAX) {
-        return std::nullopt;
-    }
-
-    GraphScalarSourceRef packed{};
-    packed.source_kind = source.source_kind == GraphRecordedScalarSourceKind::BOUNDARY ?
-                             static_cast<uint8_t>(GraphScalarSourceKind::BOUNDARY) :
-                             static_cast<uint8_t>(GraphScalarSourceKind::STATIC_VALUE);
-    packed.source_index = static_cast<uint16_t>(source.source_index);
-    return packed;
-}
-
 template <typename T>
 bool graph_layout_section(size_t count, size_t *cursor, uint32_t *offset) {
     if (count == 0) {
@@ -991,9 +998,8 @@ T *graph_image_section(std::byte *image, uint32_t offset) {
 // tasks in order.
 std::optional<GraphDefinition> graph_layout_definition(const GraphRecording &recording) {
     if (recording.unsupported || recording.task_count == 0 || recording.task_count > MAX_IN_GRAPH_TASKS ||
-        recording.boundary_tensors().empty() || recording.boundary_tensors().size() > UINT16_MAX ||
-        recording.boundary_tensors().size() != recording.boundary_types().size() ||
-        recording.boundary_args() == nullptr) {
+        recording.bound_boundary().params.tensor_count() <= 0 ||
+        recording.bound_boundary().params.tensor_count() > UINT16_MAX) {
         return std::nullopt;
     }
 
@@ -1009,7 +1015,7 @@ std::optional<GraphDefinition> graph_layout_definition(const GraphRecording &rec
     // int32 and every range test below stays in the offsets' own domain.
     const int32_t recorded_tensor_sources = static_cast<int32_t>(recording.tensor_sources.size());
     const int32_t recorded_scalars = static_cast<int32_t>(recording.scalars.size());
-    const int32_t recorded_scalar_sources = static_cast<int32_t>(recording.scalar_sources.size());
+    const int32_t recorded_scalar_inheritance = static_cast<int32_t>(recording.scalar_inheritance.size());
     const int32_t recorded_fanins = static_cast<int32_t>(recording.internal_fanins.size());
     for (int32_t i = 0; i < recording.task_count; ++i) {
         const RecordedInGraphTask &source = recording.tasks[i];
@@ -1023,8 +1029,8 @@ std::optional<GraphDefinition> graph_layout_definition(const GraphRecording &rec
             source.fanin_count > INT32_MAX - total_fanins || source.tensor_source_offset > recorded_tensor_sources ||
             source.tensor_count > recorded_tensor_sources - source.tensor_source_offset ||
             source.scalar_offset > recorded_scalars || source.scalar_count > recorded_scalars - source.scalar_offset ||
-            source.scalar_offset > recorded_scalar_sources ||
-            source.scalar_count > recorded_scalar_sources - source.scalar_offset ||
+            source.scalar_offset > recorded_scalar_inheritance ||
+            source.scalar_count > recorded_scalar_inheritance - source.scalar_offset ||
             source.fanin_offset > recorded_fanins || source.fanin_count > recorded_fanins - source.fanin_offset) {
             return std::nullopt;
         }
@@ -1041,8 +1047,8 @@ std::optional<GraphDefinition> graph_layout_definition(const GraphRecording &rec
     definition.task_count = recording.task_count;
     definition.edge_count = total_fanins;
     definition.root_count = root_count;
-    definition.boundary_count = static_cast<int32_t>(recording.boundary_tensors().size());
-    definition.boundary_scalar_count = recording.boundary_scalar_count();
+    definition.boundary_count = recording.bound_boundary().params.tensor_count();
+    definition.boundary_scalar_count = recording.bound_boundary().params.scalar_count();
     definition.tensor_arg_count = total_tensors;
     definition.scalar_arg_count = total_scalars;
     definition.predicate_count = predicate_count;
@@ -1068,9 +1074,12 @@ std::optional<GraphDefinition> graph_layout_definition(const GraphRecording &rec
         !graph_layout_section<GraphTensor>(total_tensors, &image_bytes, &definition.off_tensors) ||
         !graph_layout_section<GraphTensorSourceRef>(total_tensors, &image_bytes, &definition.off_tensor_sources) ||
         !graph_layout_section<uint64_t>(total_scalars, &image_bytes, &definition.off_scalars) ||
-        !graph_layout_section<GraphScalarSourceRef>(total_scalars, &image_bytes, &definition.off_scalar_sources) ||
+        !graph_layout_section<GraphScalarInheritance>(
+            total_scalars, &image_bytes, &definition.off_scalar_inheritance
+        ) ||
         !graph_layout_section<GraphBoundarySignature>(
-            recording.boundary_tensors().size(), &image_bytes, &definition.off_boundary_signatures
+            static_cast<size_t>(recording.bound_boundary().params.tensor_count()), &image_bytes,
+            &definition.off_boundary_signatures
         ) ||
         !graph_layout_section<GraphPredicate>(predicate_count, &image_bytes, &definition.off_predicates)) {
         return std::nullopt;
@@ -1111,7 +1120,7 @@ bool graph_fill_definition(const GraphRecording &recording, GraphDefinition defi
     auto *tensors = graph_image_section<GraphTensor>(image, definition.off_tensors);
     auto *tensor_sources = graph_image_section<GraphTensorSourceRef>(image, definition.off_tensor_sources);
     auto *scalars = graph_image_section<uint64_t>(image, definition.off_scalars);
-    auto *scalar_sources = graph_image_section<GraphScalarSourceRef>(image, definition.off_scalar_sources);
+    auto *scalar_inheritance = graph_image_section<GraphScalarInheritance>(image, definition.off_scalar_inheritance);
     auto *signatures = graph_image_section<GraphBoundarySignature>(image, definition.off_boundary_signatures);
     auto *predicates = graph_image_section<GraphPredicate>(image, definition.off_predicates);
     uint64_t required_heap = 0;
@@ -1220,18 +1229,19 @@ bool graph_fill_definition(const GraphRecording &recording, GraphDefinition defi
             tensor_cursor++;
         }
         for (int32_t scalar_index = 0; scalar_index < source.scalar_count; ++scalar_index) {
-            std::optional<GraphScalarSourceRef> packed_source =
-                graph_pack_scalar_source(recording.scalar_sources[source.scalar_offset + scalar_index]);
-            if (!packed_source.has_value() ||
-                (packed_source->source_kind == static_cast<uint8_t>(GraphScalarSourceKind::BOUNDARY) &&
-                 packed_source->source_index >= recording.boundary_args()->scalar_count())) {
+            const GraphScalarInheritance &inheritance =
+                recording.scalar_inheritance[source.scalar_offset + scalar_index];
+            // The last guard before the device indexes this: classification accounting
+            // being right does not prove the image's own layout is.
+            if (inheritance.inherited() &&
+                inheritance.boundary_index() >= recording.bound_boundary().params.scalar_count()) {
                 return false;
             }
-            scalar_sources[scalar_cursor] = *packed_source;
+            scalar_inheritance[scalar_cursor] = inheritance;
+            // An inherited slot's Definition value is a placeholder: materialize overwrites
+            // it from the invocation's own boundary.
             scalars[scalar_cursor++] =
-                packed_source->source_kind == static_cast<uint8_t>(GraphScalarSourceKind::BOUNDARY) ?
-                    0 :
-                    recording.scalars[source.scalar_offset + scalar_index];
+                inheritance.inherited() ? 0 : recording.scalars[source.scalar_offset + scalar_index];
         }
     }
     if (tensor_cursor != total_tensors || scalar_cursor != total_scalars || fanin_cursor != total_fanins ||
@@ -1248,18 +1258,19 @@ bool graph_fill_definition(const GraphRecording &recording, GraphDefinition defi
             fanout_indices[cursors[producer]++] = static_cast<uint16_t>(consumer);
         }
     }
-    for (size_t i = 0; i < recording.boundary_tensors().size(); ++i) {
-        const simpler::hbg::Tensor &tensor = recording.boundary_tensors()[i];
+    const GraphTaskArgs &boundary_params = recording.bound_boundary().params;
+    for (int32_t i = 0; i < boundary_params.tensor_count(); ++i) {
+        const simpler::hbg::Tensor &tensor = boundary_params.tensor(i).ref();
         if (tensor.ndims > MAX_TENSOR_DIMS) return false;
         uint16_t alias_rep = static_cast<uint16_t>(i);
-        for (size_t j = 0; j < i; ++j) {
-            if (recording.boundary_tensors()[j].buffer.addr == tensor.buffer.addr &&
-                recording.boundary_tensors()[j].buffer.size == tensor.buffer.size) {
+        for (int32_t j = 0; j < i; ++j) {
+            if (boundary_params.tensor(j).ref().buffer.addr == tensor.buffer.addr &&
+                boundary_params.tensor(j).ref().buffer.size == tensor.buffer.size) {
                 alias_rep = static_cast<uint16_t>(j);
                 break;
             }
         }
-        signatures[i] = graph_boundary_signature(tensor, recording.boundary_types()[i], alias_rep);
+        signatures[i] = graph_boundary_signature(tensor, boundary_params.tag(i), alias_rep);
     }
     std::memcpy(image, &definition, sizeof(definition));
     return true;
@@ -1708,7 +1719,7 @@ resolve_dispatch_predicate(OrchestratorState *orch, const CoreTaskPredicate &pre
 }
 
 // Shared body for submit_task / submit_dummy_task. Caller has already validated
-// args.has_error, decided active_mask (empty for dummy), and resolved the per-slot
+// args.has_error(), decided active_mask (empty for dummy), and resolved the per-slot
 // kernel_ids (all INVALID_KERNEL_ID for dummy). Performs tensormap sync, fanin
 // computation (explicit_deps + auto), output registration, slot init, and
 // Orch-side wiring/ready publication.
@@ -1971,18 +1982,16 @@ bool graph_boundary_matches(const GraphDefinition &definition, const GraphTaskAr
 }
 
 bool graph_boundary_matches(const GraphBoundary &boundary, const GraphTaskArgs &args) {
-    if (args.scalar_count() != boundary.scalar_count || args.explicit_dep_count() != 0 ||
-        args.tensor_count() != static_cast<int32_t>(boundary.tensors.size()) ||
-        boundary.tensors.size() != boundary.types.size()) {
+    if (args.scalar_count() != boundary.params.scalar_count() || args.explicit_dep_count() != 0 ||
+        args.tensor_count() != boundary.params.tensor_count()) {
         return false;
     }
     for (int32_t i = 0; i < args.tensor_count(); ++i) {
-        const simpler::hbg::Tensor &expected = boundary.tensors[static_cast<size_t>(i)];
+        const simpler::hbg::Tensor &expected = boundary.params.tensor(i).ref();
         const simpler::hbg::Tensor &actual = args.tensor(i).ref();
         if (actual.ndims > MAX_TENSOR_DIMS || actual.buffer.size != expected.buffer.size ||
-            actual.ndims != expected.ndims || actual.dtype != expected.dtype ||
-            args.tag(i) != boundary.types[static_cast<size_t>(i)] || actual.manual_dep != expected.manual_dep ||
-            actual.is_contiguous != expected.is_contiguous ||
+            actual.ndims != expected.ndims || actual.dtype != expected.dtype || args.tag(i) != boundary.params.tag(i) ||
+            actual.manual_dep != expected.manual_dep || actual.is_contiguous != expected.is_contiguous ||
             !std::equal(
                 std::begin(actual.shapes), std::begin(actual.shapes) + actual.ndims, std::begin(expected.shapes)
             ) ||
@@ -1994,7 +2003,7 @@ bool graph_boundary_matches(const GraphBoundary &boundary, const GraphTaskArgs &
         uint16_t expected_alias = static_cast<uint16_t>(i);
         uint16_t actual_alias = static_cast<uint16_t>(i);
         for (int32_t j = 0; j < i; ++j) {
-            const simpler::hbg::Tensor &expected_other = boundary.tensors[static_cast<size_t>(j)];
+            const simpler::hbg::Tensor &expected_other = boundary.params.tensor(j).ref();
             if (expected_other.buffer.addr == expected.buffer.addr &&
                 expected_other.buffer.size == expected.buffer.size) {
                 expected_alias = static_cast<uint16_t>(j);
@@ -2113,10 +2122,10 @@ bool graph_submit_outer(
     for (int32_t i = 0; i < args.tensor_count(); ++i)
         new (&boundary_tensors[i]) GraphTensor{graph_tensor_pack(args.tensor(i).ref())};
     if (args.scalar_count() != 0) {
-        std::memcpy(
-            payload.scalar_data(), args.scalar_data(),
-            CHIP_ALIGN_UP(static_cast<size_t>(args.scalar_count()) * sizeof(uint64_t), ARG_POOL_ALIGN)
-        );
+        // Resolved, not copied: this is the boundary the device patches BOUNDARY-sourced
+        // in-graph task scalars from, so it has to hold values. Only scalar_count entries
+        // are written; the region's alignment padding keeps whatever it held.
+        args.pack_scalars(payload.scalar_data());
     }
 
     // graph_reset_outer_payload above zeroed the count; the region delta is resolved
@@ -2247,7 +2256,7 @@ TaskOutputTensors graph_record_submit_in_graph_task(
     const TaskId task_id = TaskId::make_in_graph(GRAPH_RECORD_NO_OWNING_GRAPH, task_index);
     result.set_task_id(task_id);
 
-    if (task_index >= MAX_IN_GRAPH_TASKS || args.has_error) {
+    if (task_index >= MAX_IN_GRAPH_TASKS || args.has_error()) {
         recording.unsupported = true;
     }
 
@@ -2335,22 +2344,21 @@ TaskOutputTensors graph_record_submit_in_graph_task(
     }
     task.scalar_offset = static_cast<int32_t>(recording.scalars.size());
     task.scalar_count = args.scalar_count();
-    recording.scalars.insert(recording.scalars.end(), args.scalars(), args.scalars() + args.scalar_count());
+    // Resolved values, not slots: an inherited slot's word is a host pointer, and this is
+    // the recording's working copy of what the task passed. What a BOUNDARY-sourced slot
+    // contributes to the Definition is overwritten with a placeholder at build time.
+    recording.scalars.resize(recording.scalars.size() + static_cast<size_t>(task.scalar_count));
+    args.pack_scalars(recording.scalars.data() + task.scalar_offset);
 #if SIMPLER_DFX
     task.dump_metadata.dump_arg_mask = args.dump_arg_mask();
     task.dump_metadata.dump_arg_flags = args.dump_arg_index_ambiguous_mask();
     memcpy(task.dump_metadata.scalar_dtypes, args.scalar_dtypes(), args.scalar_count() * sizeof(uint8_t));
 #endif
 
-    // Classify each scalar's source: a plain literal is static Definition data,
-    // while a value copied from a boundary scalar is refreshed on replay. A
-    // mutable tracked boundary scalar is not supported and falls back.
-    recording.scalar_sources.resize(task.scalar_offset + task.scalar_count);
-    for (int32_t i = 0; i < args.scalar_count(); ++i) {
-        GraphRecordedScalarSourceRef source = graph_classify_scalar(recording, args, i);
-        if (source.source_kind == GraphRecordedScalarSourceKind::INVALIDATED_BOUNDARY) recording.unsupported = true;
-        recording.scalar_sources[task.scalar_offset + i] = source;
-    }
+    // Classify each scalar's source: a slot holding its own value is static Definition
+    // data, while an inherited slot names a boundary parameter refreshed on every replay.
+    recording.scalar_inheritance.resize(task.scalar_offset + task.scalar_count);
+    graph_classify_scalars(recording, args, task.scalar_offset);
 
     // Classify each tensor's source, then derive internal fanins from the
     // INTERNAL classifications plus any explicit internal dependency.
@@ -2481,12 +2489,14 @@ TaskOutputTensors graph_record_submit_in_graph_task(
             // Only the outer shell can order the body behind a pre-Graph task, and it
             // does so through its boundary args -- so a dep no boundary tensor carries
             // has no edge in the Definition and the body cannot be recorded.
-            const bool represented_by_boundary = std::any_of(
-                recording.boundary_tensors().begin(), recording.boundary_tensors().end(),
-                [dep](const simpler::hbg::Tensor &tensor) {
-                    return tensor.owner_task_id == dep;
+            bool represented_by_boundary = false;
+            const GraphTaskArgs &boundary_params = recording.bound_boundary().params;
+            for (int32_t i = 0; i < boundary_params.tensor_count(); ++i) {
+                if (boundary_params.tensor(i).ref().owner_task_id == dep) {
+                    represented_by_boundary = true;
+                    break;
                 }
-            );
+            }
             if (!represented_by_boundary) recording.unsupported = true;
             continue;
         }
@@ -2602,13 +2612,47 @@ OrchestratorState::graph_begin_inner(uint64_t graph_key, const GraphTaskArgs &ar
     // up here would sit on the submitting thread, between two outer shells.
     auto entry = std::make_unique<GraphInflightRecording>();
     entry->full_key = full_key;
-    entry->boundary.scalar_count = args.scalar_count();
-    entry->boundary.tensors.reserve(static_cast<size_t>(args.tensor_count()));
-    entry->boundary.types.reserve(static_cast<size_t>(args.tensor_count()));
+    // The boundary is built once, here, and only read afterwards -- by the recorder that
+    // picks this entry up, and by later same-key submissions comparing against it.
+    //
+    // Tensors are filled before any TensorRef is made to point at them: `tensors` is a
+    // fixed-size array precisely so those pointers cannot move, but a slot must still hold
+    // its value before args names it.
+    GraphBoundary &boundary = entry->boundary;
+    for (int32_t i = 0; i < args.tensor_count(); ++i) {
+        boundary.tensors[static_cast<size_t>(i)] = args.tensor(i).ref();
+    }
     for (int32_t i = 0; i < args.tensor_count(); ++i) {
-        entry->boundary.tensors.push_back(args.tensor(i).ref());
-        entry->boundary.types.push_back(args.tag(i));
+        simpler::hbg::Tensor &owned = boundary.tensors[static_cast<size_t>(i)];
+        switch (args.tag(i)) {
+        case TensorArgType::INPUT:
+            boundary.params.add_input(owned);
+            break;
+        case TensorArgType::OUTPUT_EXISTING:
+            boundary.params.add_output(owned);
+            break;
+        case TensorArgType::INOUT:
+            boundary.params.add_inout(owned);
+            break;
+        case TensorArgType::NO_DEP:
+            boundary.params.add_no_dep(owned);
+            break;
+        case TensorArgType::OUTPUT:
+            // GraphTaskArgs::add_output rejects a TensorCreateInfo at compile time, so no
+            // boundary carries this tag. The case exists because the switch is exhaustive.
+            debug_assert(false && "a Graph boundary cannot hold a runtime-allocated output");
+            break;
+        }
     }
+    // Values resolved, declarations carried over. A dynamic parameter names itself, so
+    // `params` stays the basis recording resolves against.
+    boundary.params.gen_scalar_params_from_args(args);
+    boundary.params.launch_spec = args.launch_spec;
+    boundary.params.set_allow_early_resolve(args.allow_early_resolve());
+    if (args.task_timing_slot() != TASK_TIMING_SLOT_NONE) {
+        boundary.params.set_task_timing_slot(args.task_timing_slot());
+    }
+    boundary.params.set_predicate(args.predicate());
     GraphInflightRecording *entry_ptr = entry.get();
     state->inflight.emplace(full_key, std::move(entry));
     state->inflight_count.store(state->inflight.size(), std::memory_order_release);
@@ -2619,6 +2663,7 @@ OrchestratorState::graph_begin_inner(uint64_t graph_key, const GraphTaskArgs &ar
         result.execute_block = false;
         result.recording = true;
         result.recording_handle = entry_ptr;
+        result.params = &entry_ptr->boundary.params;
         result.task_id = submitted;
         ORCH_PHASE_END(HostPhaseKind::OrchGraphSubmit, submitted.raw);
 #if SIMPLER_DFX
@@ -2634,7 +2679,13 @@ OrchestratorState::graph_begin_inner(uint64_t graph_key, const GraphTaskArgs &ar
     return result;
 }
 
-bool OrchestratorState::graph_prepare(void *recording_handle, const GraphTaskArgs &args) {
+// The parameter list is the entry's own -- graph_begin published it as
+// GraphScopeResult::params, and both the queued job and the synchronous fallback forward
+// that same object here, which is why it arrives unnamed: there is no second boundary for
+// it to agree with. A later same-key submission does arrive with the caller's own args,
+// and graph_begin_inner compares those against this entry before publishing a shell
+// against it.
+bool OrchestratorState::graph_prepare(void *recording_handle, const GraphTaskArgs &) {
     GraphHostState *state = graph_state_from(this);
     if (state == nullptr || recording_handle == nullptr || g_active_graph_recording != nullptr) return false;
     auto *entry = static_cast<GraphInflightRecording *>(recording_handle);
@@ -2642,22 +2693,13 @@ bool OrchestratorState::graph_prepare(void *recording_handle, const GraphTaskArg
     // the entry's address is stable for as long as the recording lives, so the
     // recording thread reaches its own state without searching for it. Until this
     // thread calls graph_end/graph_abort, later graph_begin calls only read the
-    // boundary vectors under recording_mutex, and only this thread writes the
+    // boundary under recording_mutex, and only this thread writes the
     // fields it binds below. Taking that mutex here lets the main thread's
     // same-key submit burst starve prepare and collapse the intended overlap, so
     // the status read goes through the atomic instead.
     if (entry->status() != GraphRecordingStatus::RECORDING) {
         return false;
     }
-    // The entry was created from this very boundary at graph_begin, and the handle names
-    // that entry rather than being searched for, so a mismatch here is unreachable. The
-    // comparison walks up to 128 simpler::hbg::Tensor descriptors on the thread whose start-up
-    // latency this path exists to keep short, so it is an assertion: debug builds still
-    // catch a boundary that stopped matching, release builds compile it out.
-    debug_assert(
-        graph_boundary_matches(entry->boundary, args) &&
-        "the entry's boundary copy must match the boundary graph_begin recorded"
-    );
     // This thread's own storage, emptied rather than allocated -- see
     // recorder_recording(). Failure is reachable only on this thread's first recording,
     // where the hazard map is stood up; the caller then aborts the recording, and the
@@ -2667,8 +2709,6 @@ bool OrchestratorState::graph_prepare(void *recording_handle, const GraphTaskArg
         LOG_WARN("%s", "[GraphExecution] recording hazard map allocation failed; recording abandoned");
         return false;
     }
-    args.anchor_scalar_sources();
-    entry->boundary.args = &args;
     g_active_graph_entry = entry;
     g_active_graph_recording = &recording;
     g_active_graph_owner = state;
@@ -2686,8 +2726,9 @@ void OrchestratorState::graph_abort(void *recording_handle) {
     // The storage outlives the entry it was bound to, and graph_commit destroys the
     // entries, so leaving the pointer behind parks a stale one in thread_local state for
     // the rest of the process. The next graph_prepare rebinds before anything reads it,
-    // which is why this is hygiene rather than a fix -- but boundary_tensors() does not
-    // null-check, so a future reader outside a recording would follow it.
+    // which is why this is hygiene rather than a fix -- but clearing it turns a dangling
+    // boundary into a null one, and null is the only state bound_boundary()'s assertion
+    // can catch a read outside a recording by.
     unbind_recorder_boundary();
     g_active_graph_entry = nullptr;
     g_active_graph_recording = nullptr;
@@ -2828,11 +2869,11 @@ TaskOutputTensors OrchestratorState::submit_task(const MixedKernels &mixed_kerne
     }
 
     // Validate Arg construction (errors recorded by add_input/add_output/etc.)
-    if (args.has_error) {
+    if (args.has_error()) {
         LOG_ERROR("========================================");
         LOG_ERROR("FATAL: Invalid Arg Detected!");
         LOG_ERROR("========================================");
-        LOG_ERROR("Error: %s", args.error_msg ? args.error_msg : "(unknown)");
+        LOG_ERROR("Error: %s", args.error_msg() ? args.error_msg() : "(unknown)");
         LOG_ERROR("  tensor_count: %d, scalar_count: %d", args.tensor_count(), args.scalar_count());
         LOG_ERROR("This is a bug in the orchestration code.");
         LOG_ERROR("========================================");
@@ -2916,11 +2957,11 @@ TaskOutputTensors OrchestratorState::submit_dummy_task(const CoreTaskArgs &args)
         return TaskOutputTensors{};
     }
 
-    if (args.has_error) {
+    if (args.has_error()) {
         LOG_ERROR("========================================");
         LOG_ERROR("FATAL: Invalid Arg in submit_dummy_task!");
         LOG_ERROR("========================================");
-        LOG_ERROR("Error: %s", args.error_msg ? args.error_msg : "(unknown)");
+        LOG_ERROR("Error: %s", args.error_msg() ? args.error_msg() : "(unknown)");
         LOG_ERROR("  tensor_count: %d, scalar_count: %d", args.tensor_count(), args.scalar_count());
         LOG_ERROR("========================================");
         orch_mark_fatal(orch, SIMPLER_ERROR_INVALID_ARGS);
@@ -2974,10 +3015,10 @@ TaskOutputTensors OrchestratorState::alloc_tensors(const CoreTaskArgs &args) {
     ORCH_STEP_START();
     ORCH_PHASE_START();
 
-    if (args.has_error) {
+    if (args.has_error()) {
         report_fatal(
             SIMPLER_ERROR_INVALID_ARGS, __FUNCTION__, "%s",
-            args.error_msg ? args.error_msg : "alloc_tensors failed to construct output-only Arg"
+            args.error_msg() ? args.error_msg() : "alloc_tensors failed to construct output-only Arg"
         );
         return TaskOutputTensors{};
     }
diff --git a/src/common/host_build_graph/runtime_ops.h b/src/common/host_build_graph/runtime_ops.h
index 67255ada98..caa085f4a1 100644
--- a/src/common/host_build_graph/runtime_ops.h
+++ b/src/common/host_build_graph/runtime_ops.h
@@ -86,7 +86,7 @@ struct RuntimeOps {
     // layout does not move with SIMPLER_DFX; nullptr when DFX is off.
     void (*record_orch_phase)(uint32_t kind, uint64_t start_ns, uint64_t end_ns, uint64_t detail);
     // Queue one Graph body for asynchronous recording, and drain every queued one.
-    // `job` is a `std::function<void(GraphTaskArgs &)> *` the pool moves out of --
+    // `job` is a `std::function<void(const GraphTaskArgs &)> *` the pool moves out of --
     // whether or not it queues it, since start() takes the callable before it checks
     // capacity -- so the caller must not invoke it afterwards. Nothing is owned across
     // the boundary either way: the caller's std::function destructs normally, empty or
diff --git a/src/common/host_build_graph/runtime_types.h b/src/common/host_build_graph/runtime_types.h
index d90cabd0d3..e26f6544a8 100644
--- a/src/common/host_build_graph/runtime_types.h
+++ b/src/common/host_build_graph/runtime_types.h
@@ -491,11 +491,10 @@ struct TaskPayload {
                 result.materialize_output(dst[i]);
             }
         }
-        // Round up to cache line boundary. Every scalar region is a whole number of
-        // cache lines (ARG_POOL_ALIGN), so the rounded copy stays inside this
-        // task's own region. Eliminates branches; extra bytes within the same CL have
-        // zero additional cost.
-        memcpy(scalar_data(), args.scalars(), CHIP_ALIGN_UP(args.scalar_count() * sizeof(uint64_t), 64));
+        // A slot is a value, so this is one memcpy of scalar_count entries. The trailing
+        // padding of this task's cache-line-aligned region keeps whatever it held --
+        // nothing reads past the count.
+        args.pack_scalars(scalar_data());
 
         // The task table's payload storage is raw shared memory that no constructor
         // runs over, so an unset predicate reads back as whatever the slot last held —
diff --git a/src/common/host_build_graph/types.h b/src/common/host_build_graph/types.h
index a5c9cbf6ed..f88968d9d5 100644
--- a/src/common/host_build_graph/types.h
+++ b/src/common/host_build_graph/types.h
@@ -183,12 +183,76 @@ class TensorRef {
     bool refers_to(const TensorCreateInfo *ci) const { return create_info_ == ci; }
 };
 
+template <size_t MaxT, size_t MaxS>
+class Arg;
+class InheritableScalar;
+
+/**
+ * A scalar parameter handed out for forwarding: its value, and the address of the slot it
+ * ultimately comes from.
+ *
+ * Carrying both is what lets a destination store the value without ever dereferencing the
+ * origin. So the origin is allowed to dangle -- a caller's local goes out of scope once
+ * submit returns -- and the value is still right. The address is only ever tested against
+ * null (dynamic vs static parameter) or subtracted from a boundary's slot array base to
+ * recover a parameter index (recording).
+ *
+ * Arg::scalar(i) folds: an already-inherited slot yields its own origin rather than
+ * itself, so a chain is exactly one hop and recording resolves it without a walk.
+ */
+class InheritableScalar {
+public:
+    constexpr InheritableScalar(uint64_t bits, const void *origin) :
+        bits_(bits),
+        origin_(origin) {}
+
+    constexpr const void *origin() const { return origin_; }
+
+    // Read the parameter as a value, with to_u64's actual inverse applied. static_cast on
+    // the pattern is not that inverse: a float slot holds a bit pattern, so
+    // static_cast<float> of 1.0f's pattern yields 1065353216.0. It is also the only
+    // spelling that reaches an enum, since a conversion to an enumeration does not accept
+    // a user-defined one on the way.
+    template <typename T>
+    T to() const {
+        return from_u64<T>(bits_);
+    }
+
+    // Implicit, so an existing value read still compiles; deprecated, so it says so.
+    // Forwarding stores origin() and never reaches here, which is what keeps a
+    // pass-through silent.
+    [[deprecated(
+        "scalar slot read as a value, which breaks inheritance: the destination "
+        "slot holds this invocation's number instead of following the source, so "
+        "a Graph Definition freezes it. Forward it (add_scalar(args.scalar(i))) "
+        "to keep it per-invocation; use args.scalar(i).to<T>() if a static "
+        "value is intended; pass it as a construction parameter if it selects "
+        "the Graph's structure."
+    )]]
+    operator uint64_t() const {
+        return bits_;
+    }
+
+private:
+    uint64_t bits_;
+    const void *origin_;
+};
+static_assert(sizeof(InheritableScalar) == 16, "InheritableScalar is passed by value and holds a value and an origin");
+
+// Defined here rather than beside is_supported_scalar_arg_v in data_type.h: that trait
+// is shared with the tensormap_and_ringbuffer runtime, which has no Graph and whose
+// dtype_of/mark_dump_arg would silently accept a type they cannot describe.
+template <typename T>
+inline constexpr bool is_inheritable_scalar_v = std::is_same_v<std::decay_t<T>, InheritableScalar>;
+
 /**
  * Aggregated argument container for rt_submit_task
  *
  * Inherits storage from TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS, TensorArgType>.
  * Each tensor slot stores a TensorRef union (simpler::hbg::Tensor* or TensorCreateInfo)
- * discriminated by the corresponding tag().
+ * discriminated by the corresponding tag(). Each scalar slot stores a value; the parallel
+ * scalar_inherited_ array names where that value came from, or is null when the parameter
+ * is static.
  * ChipTensors are dispatched first in kernel args, followed by scalars.
  *
  * Output arguments follow two distinct ownership models:
@@ -228,15 +292,21 @@ struct CoreTaskPredicate {
 };
 
 template <size_t MaxT, size_t MaxS>
-struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
+class Arg : private TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
     using Base = TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType>;
-    // Make dependent-base members visible for unqualified use (two-phase lookup
-    // does not search a dependent base in a class template).
-    using Base::scalar_count_;
-    using Base::scalars_;
-    using Base::tags_;
-    using Base::tensor_count_;
-    using Base::tensors_;
+
+public:
+    // The base's own API, re-exported one name at a time. Private inheritance is what
+    // makes that a choice rather than a default: a public base is reachable by an
+    // implicit derived-to-base conversion, through which its members are public again
+    // however this class hides their names, so `static_cast<const Base &>(args).tags_`
+    // would read the tag array the accessors exist to mediate.
+    using Base::scalar_count;
+    using Base::tag;
+    using Base::tag_data;
+    using Base::tensor;
+    using Base::tensor_count;
+    using Base::tensor_data;
 
     // Minimal-permission: an Arg is built in place and consumed by reference;
     // it is never copied/moved (it is a large object, and its TensorRef slots
@@ -247,34 +317,17 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
     Arg &operator=(const Arg &) = delete;
     Arg &operator=(Arg &&) = delete;
 
-    bool has_error{false};
-    const char *error_msg{nullptr};
-    LaunchSpec launch_spec;  // SPMD launch parameters (block_num, etc.)
+    bool has_error() const { return has_error_; }
+    const char *error_msg() const { return error_msg_; }
 
-    // Speculative early-dispatch hint (codegen-author set, off by default). When
-    // true, the scheduler may stage this task on an idle core before its producer
-    // finishes, gating execution on the DATA_MAIN_BASE doorbell — only safe when
-    // the author knows the task's data dependencies allow it. Read in-process by
-    // the runtime; never crosses the wire format.
-    bool allow_early_resolve_{false};
     void set_allow_early_resolve(bool v = true) { allow_early_resolve_ = v; }
     bool allow_early_resolve() const { return allow_early_resolve_; }
 
-    // Dispatch predicate (codegen-author set; default op == NONE = always
-    // dispatch). A FALSE result at the dispatch point retires the task inline
-    // through the dep-only path — never dispatched to an AICore — while still
-    // resolving fanin/fanout so consumers unlock. The predicate tensor's producer
-    // MUST be a dependency of this task so the value is current when the task
-    // becomes ready. Read in-process; never crosses the wire.
-    CoreTaskPredicate predicate_;
     void set_predicate(const CoreTaskPredicate &pred) { predicate_ = pred; }
     const CoreTaskPredicate &predicate() const { return predicate_; }
 
-    // Selective task-timing slot: tag this task to have the scheduler record its
-    // AICPU dispatch/finish cycles into fixed slot `slot` (0..15). Untagged by
-    // default. An out-of-range id fails through the standard invalid-arg path so
-    // the scheduler never stamps out of bounds.
-    int32_t task_timing_slot_{TASK_TIMING_SLOT_NONE};
+    // An out-of-range id fails through the standard invalid-arg path so the scheduler
+    // never stamps out of bounds.
     void set_task_timing_slot(int32_t slot) {
         if (slot < 0 || slot >= NUM_TASK_TIMING_SLOTS) {
             set_error("task_timing_slot out of range (valid: 0..15)");
@@ -286,8 +339,9 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
 
     void clear() {
         Base::clear();
-        scalar_sources_.fill(nullptr);
-        scalar_sources_invalidated_.fill(false);
+        // All-null is a correct empty state: every parameter reads as static, so nothing
+        // below scalar_count_ can be mistaken for one that follows a source.
+        scalar_inherited_.fill(nullptr);
 #if SIMPLER_DFX
         dump_arg_selection_.clear();
 #endif
@@ -300,14 +354,14 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
 
     void reset() {
         clear();
-        has_error = false;
-        error_msg = nullptr;
+        has_error_ = false;
+        error_msg_ = nullptr;
     }
 
     void set_error(const char *msg) {
-        if (!has_error) {
-            has_error = true;
-            error_msg = msg;
+        if (!has_error_) {
+            has_error_ = true;
+            error_msg_ = msg;
         }
     }
 
@@ -429,141 +483,136 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
     const TaskId *explicit_deps_data() const { return explicit_deps_; }
 
     /**
-     * Add scalar values. Types are deduced per argument; each value is
-     * bit-cast to uint64_t for storage. Mixed types are allowed:
+     * Add scalar values, declaring each one a **dynamic** parameter. Types are deduced per
+     * argument; each value is bit-cast to uint64_t for storage. Mixed types are allowed:
+     *
+     *   args.add_scalar(token_pos);                   // single
+     *   args.add_scalar(3.14f, int32_t(42), 7u);      // mixed batch
+     *
+     * "Dynamic" means the value may differ per invocation, so a Graph cache lookup must
+     * not compare it. What declares it is the argument's value category: an lvalue (the
+     * caller holds it somewhere) or an InheritableScalar (it already names a parameter)
+     * is dynamic; a literal or any other rvalue is static. Say it explicitly with
+     * add_static_scalar when an lvalue holds a value that does not change.
      *
-     *   args.add_scalar(uint64_val);                  // single
-     *   args.add_scalar(3.14f, int32_t(42), 7u);     // mixed batch
+     * A GraphTaskArgs::scalar(i) may be passed alongside plain values; that slot then
+     * follows the boundary parameter instead of freezing its current value.
      */
     template <typename... Args>
     void add_scalar(Args &&...args) {
         static_assert(sizeof...(Args) >= 1, "add_scalar: at least one argument required");
-        static_assert((is_supported_scalar_arg_v<Args> && ...), "add_scalar: all types must be arithmetic or enum");
+        static_assert(
+            ((is_supported_scalar_arg_v<Args> || is_inheritable_scalar_v<Args>) && ...),
+            "add_scalar: all types must be arithmetic, enum, or a Graph boundary scalar"
+        );
         if (scalar_count_ + sizeof...(Args) > MaxS) {
             set_error(scalar_cap_msg());
             return;
         }
-        (add_scalar_one(std::forward<Args>(args)), ...);
-    }
-
-    void add_scalars(const uint64_t *values, int count) {
-        if (count < 0 || scalar_count_ + count > MaxS) {
-            set_error(scalar_cap_msg());
-            return;
-        }
-        memcpy(&scalars_[scalar_count_], values, count * sizeof(uint64_t));
-        std::fill_n(scalar_sources_.begin() + scalar_count_, count, nullptr);
-        std::fill_n(scalar_sources_invalidated_.begin() + scalar_count_, count, false);
-#if SIMPLER_DFX
-        dump_arg_selection_.clear_scalar_metadata(scalar_count_, count);
-#endif
-        scalar_count_ += count;
+        (add_scalar_one<true>(std::forward<Args>(args)), ...);
     }
 
     /**
-     * Zero-extend int32 bit patterns into uint64 scalar slots.
-     * Negative values are treated as their unsigned 32-bit representation
-     * (e.g., -1 → 0x00000000FFFFFFFF, not 0xFFFFFFFFFFFFFFFF).
-     * Uses NEON to process 4 elements per iteration on aarch64.
+     * Add scalar values, declaring each one a **static** parameter regardless of value
+     * category.
+     *
+     * On an in-graph task's arguments the declaration takes effect at once: recording
+     * reads it back through scalar_dynamic() and records the slot as static Definition
+     * data instead of following a parameter.
+     *
+     * On a Graph's own parameter list it is inert for now. gen_scalar_params_from_args
+     * carries it across, but graph_full_key is callable_hash and graph_key, so no lookup
+     * compares a scalar value; callers accordingly declare every parameter dynamic. This
+     * is what they will say otherwise with, once a scalar value is part of the condition
+     * a Definition is reused under -- which must not precede their migration (#2170),
+     * since a parameter a body freezes while declared dynamic would then match on a
+     * Definition holding a stale number.
+     *
+     * An InheritableScalar passed here is resolved to its value rather than followed --
+     * this is how an enclosing Graph's parameter is deliberately frozen into an inner
+     * boundary.
      */
-    void add_scalars_i32(const int32_t *values, int count) {
-        if (count < 0 || scalar_count_ + count > MaxS) {
+    template <typename... Args>
+    void add_static_scalar(Args &&...args) {
+        static_assert(sizeof...(Args) >= 1, "add_static_scalar: at least one argument required");
+        static_assert(
+            ((is_supported_scalar_arg_v<Args> || is_inheritable_scalar_v<Args>) && ...),
+            "add_static_scalar: all types must be arithmetic, enum, or a Graph boundary scalar"
+        );
+        if (scalar_count_ + sizeof...(Args) > MaxS) {
             set_error(scalar_cap_msg());
             return;
         }
-        uint64_t *dst = &scalars_[scalar_count_];
-#if defined(__aarch64__)
-        int i = 0;
-        for (; i + 4 <= count; i += 4) {
-            uint32x4_t v = vld1q_u32(reinterpret_cast<const uint32_t *>(values + i));
-            uint64x2_t lo = vmovl_u32(vget_low_u32(v));
-            uint64x2_t hi = vmovl_u32(vget_high_u32(v));
-            vst1q_u64(dst + i, lo);
-            vst1q_u64(dst + i + 2, hi);
-        }
-        for (; i < count; i++) {
-            dst[i] = static_cast<uint64_t>(static_cast<uint32_t>(values[i]));
-        }
-#else
-        for (int i = 0; i < count; i++) {
-            dst[i] = static_cast<uint64_t>(static_cast<uint32_t>(values[i]));
-        }
-#endif
-        std::fill_n(scalar_sources_.begin() + scalar_count_, count, nullptr);
-        std::fill_n(scalar_sources_invalidated_.begin() + scalar_count_, count, false);
-#if SIMPLER_DFX
-        dump_arg_selection_.clear_scalar_metadata(scalar_count_, count);
-#endif
-        scalar_count_ += count;
+        (add_scalar_one<false>(std::forward<Args>(args)), ...);
     }
 
-    /**
-     * Copy scalars from another Arg's scalar array.
-     * Useful when multiple tasks share the same scalar data (e.g., block indices).
-     */
-    void copy_scalars_from(const Arg &src, int src_offset, int count) {
-        if (src_offset < 0 || count < 0 || src_offset + count > src.scalar_count_) {
-            set_error("Source scalar range out of bounds in copy_scalars_from");
-            return;
-        }
-        if (scalar_count_ + count > MaxS) {
-            set_error(scalar_cap_msg());
-            return;
-        }
-        memcpy(&scalars_[scalar_count_], &src.scalars_[src_offset], count * sizeof(uint64_t));
-        for (int i = 0; i < count; ++i) {
-            const int src_index = src_offset + i;
-            scalar_sources_[scalar_count_ + i] = src.scalar_sources_[src_index] != nullptr ?
-                                                     src.scalar_sources_[src_index] :
-                                                     static_cast<const void *>(&src.scalars_[src_index]);
-            scalar_sources_invalidated_[scalar_count_ + i] = src.scalar_sources_invalidated_[src_index];
-        }
-#if SIMPLER_DFX
-        dump_arg_selection_.copy_scalar_dtypes_from(src.dump_arg_selection_, scalar_count_, src_offset, count);
-#endif
-        scalar_count_ += count;
+    // Bulk form of add_scalar: an array element is an lvalue, so every parameter added
+    // here is dynamic. Use add_static_scalars for a run of values that does not change.
+    template <typename T>
+    void add_scalars(const T *values, int count) {
+        add_scalars_impl(values, count, true);
     }
 
-    const uint64_t &scalar(int32_t i) const { return scalars_[i]; }
-    uint64_t &scalar(int32_t i) {
-        // A mutable reference may escape and be written later. Conservatively
-        // invalidate inherited Graph-boundary provenance as soon as it is requested.
-        scalar_sources_invalidated_[i] = scalar_sources_[i] != nullptr;
-        return scalars_[i];
-    }
-    const void *scalar_source(int32_t i) const { return scalar_sources_invalidated_[i] ? nullptr : scalar_sources_[i]; }
-    const void *invalidated_scalar_source(int32_t i) const {
-        return scalar_sources_invalidated_[i] ? scalar_sources_[i] : nullptr;
-    }
-    // Graph recording starts before its function runs. Anchor the host-only
-    // provenance at that point so every boundary slot has a unique identity,
-    // even when multiple slots were initialized from the same lvalue.
-    void anchor_scalar_sources() const {
-        for (int32_t i = 0; i < scalar_count_; ++i) {
-            scalar_sources_[i] = static_cast<const void *>(&scalars_[i]);
-            scalar_sources_invalidated_[i] = false;
-        }
+    template <typename T>
+    void add_static_scalars(const T *values, int count) {
+        add_scalars_impl(values, count, false);
     }
 
+    // Hand out parameter i for forwarding: its value, plus the slot that value comes from.
+    // Passing the result to another Arg's add_scalar makes that slot follow this one;
+    // reading it as a value goes through InheritableScalar's deprecated conversion, which
+    // is what makes breaking the chain visible at the call site.
+    //
+    // An already-inherited slot yields its own origin rather than itself, so the chain a
+    // destination records is always one hop: "C inherits B, B inherits A" records C -> A.
+    InheritableScalar scalar(int32_t i) const {
+        return {scalars_[i], scalar_inherited_[i] != nullptr ? scalar_inherited_[i] : &scalars_[i]};
+    }
+
+    // Copy every value into a compact uint64 array -- what submit and payload
+    // materialization need. A slot is always a value, so this is one memcpy: no
+    // discriminator to consult, nothing to resolve.
+    void pack_scalars(uint64_t *out) const {
+        memcpy(out, scalars_, static_cast<size_t>(scalar_count_) * sizeof(uint64_t));
+    }
+
+    // A parameter is dynamic exactly when it names where its value came from. The address
+    // is never dereferenced -- see scalar_origin.
+    bool scalar_dynamic(int32_t i) const { return scalar_inherited_[i] != nullptr; }
+
+    // Where parameter i's value came from, or null when it is static.
+    //
+    // This address MUST NOT be dereferenced: it may already dangle, since a caller's local
+    // goes out of scope once submit returns. It is only ever tested against null and
+    // subtracted from a slot array base, and that is safe precisely because the value was
+    // copied at add_scalar time.
+    const void *scalar_origin(int32_t i) const { return scalar_inherited_[i]; }
+
+    // Base of the slot array. Its only use is turning an origin pointer into a boundary
+    // parameter index; named apart from a value accessor because what it hands out is an
+    // identity baseline rather than a run of values.
+    const void *scalar_slot_base() const { return scalars_; }
+
 #if SIMPLER_DFX
     const uint8_t *scalar_dtypes() const { return dump_arg_selection_.scalar_dtypes(); }
 #else
     const uint8_t *scalar_dtypes() const { return nullptr; }
 #endif
 
+protected:
+    // Capacity-overflow message — spells the actual limit (MaxS, whatever the
+    // instantiation is) into the text via std::to_string. Built once into a
+    // function-local static so set_error() can hold the const char* safely.
+    static const char *scalar_cap_msg() {
+        static const std::string msg = "Too many scalar args (max " + std::to_string(MaxS) + ")";
+        return msg.c_str();
+    }
+
 private:
-    // In-process recording metadata only; it is never copied into a task payload
-    // or any host-device wire image.
-    mutable std::array<const void *, MaxS> scalar_sources_{};
-    // Kept separately so invalidated boundary ancestry can reject Graph caching
-    // instead of being mistaken for an unrelated static scalar.
-    mutable std::array<bool, MaxS> scalar_sources_invalidated_{};
-    // Caller-owned dependency array; lifetime must extend through submit.
-#if SIMPLER_DFX
-    DumpArgSelection dump_arg_selection_;
-#endif
-    const TaskId *explicit_deps_{nullptr};
-    uint32_t explicit_dep_count_{0};
+    // Held fully private: these two hand out the slot array with nothing to qualify it.
+    using Base::scalar_data;
+    using Base::scalars;
+
 #if SIMPLER_DFX
     template <typename T>
     static constexpr bool is_supported_dump_arg_v =
@@ -571,36 +620,85 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
         is_supported_scalar_arg_v<T>;
 #endif
 
-    // Capacity-overflow messages — spell the actual limit (MaxS/MaxT, whatever
-    // the instantiation is) into the text via std::to_string. Built once into a
-    // function-local static so set_error() can hold the const char* safely.
-    static const char *scalar_cap_msg() {
-        static const std::string msg = "Too many scalar args (max " + std::to_string(MaxS) + ")";
-        return msg.c_str();
-    }
     static const char *tensor_cap_msg() {
         static const std::string msg = "Too many tensor args (max " + std::to_string(MaxT) + ")";
         return msg.c_str();
     }
 
-    template <typename T>
+    // Add one value plus its declaration. Dynamic == true is add_scalar's promise that the
+    // caller may change this parameter between invocations; the origin recorded for it is
+    // what a Graph recording turns into a boundary parameter index.
+    //
+    // The value is converted here, where T is still known -- which is why a parameter of
+    // any width can be dynamic. Deferring the conversion to a read of the origin would
+    // require the origin's width, and the slot has nowhere to keep it.
+    template <bool Dynamic, typename T>
     void add_scalar_one(T &&value) {
-        scalars_[scalar_count_] = to_u64(value);
-        scalar_sources_[scalar_count_] = nullptr;
-        scalar_sources_invalidated_[scalar_count_] = false;
-        if constexpr (std::is_lvalue_reference_v<T>) {
-            scalar_sources_[scalar_count_] = static_cast<const void *>(&value);
-        }
+        if constexpr (is_inheritable_scalar_v<T>) {
+            // The value travels with the handle, so following the parameter costs no
+            // dereference of the origin -- and lets the origin dangle harmlessly. Read
+            // through to<uint64_t>() rather than a bits() getter: a public bits() would be
+            // a second silent value-read path, exactly what the deprecated conversion
+            // exists to surface.
+            scalars_[scalar_count_] = value.template to<uint64_t>();
+            scalar_inherited_[scalar_count_] = Dynamic ? value.origin() : nullptr;
 #if SIMPLER_DFX
-        uintptr_t scalar_source_ptr = 0;
-        scalar_source_ptr = reinterpret_cast<uintptr_t>(scalar_sources_[scalar_count_]);
-        dump_arg_selection_.record_scalar_source(
-            scalar_count_, scalar_source_ptr, dtype_of<std::remove_cv_t<std::remove_reference_t<T>>>()
-        );
+            // No host address to identify this slot by: it names a boundary parameter,
+            // not a caller variable, so dump() cannot match it by pointer. The dtype is
+            // u64 for the same reason -- an InheritableScalar carries bits and origin and
+            // nothing else -- so a forwarded float or int32 slot prints as its bit
+            // pattern rather than as its source type.
+            dump_arg_selection_.record_scalar_source(scalar_count_, 0, dtype_of<uint64_t>());
 #endif
+        } else {
+            scalars_[scalar_count_] = to_u64(value);
+            // An lvalue is a parameter the caller holds and may change; an rvalue cannot
+            // be changed by anyone, so it is static however add_scalar was called.
+            if constexpr (Dynamic && std::is_lvalue_reference_v<T>) {
+                scalar_inherited_[scalar_count_] = &value;
+            } else {
+                scalar_inherited_[scalar_count_] = nullptr;
+            }
+#if SIMPLER_DFX
+            uintptr_t scalar_source_ptr = 0;
+            if constexpr (std::is_lvalue_reference_v<T>) {
+                scalar_source_ptr = reinterpret_cast<uintptr_t>(&value);
+            }
+            dump_arg_selection_.record_scalar_source(
+                scalar_count_, scalar_source_ptr, dtype_of<std::remove_cv_t<std::remove_reference_t<T>>>()
+            );
+#endif
+        }
         scalar_count_++;
     }
 
+    // Shared body of add_scalars / add_static_scalars. An array element is an lvalue, so
+    // value category cannot tell a dynamic parameter from a static one the way it does in
+    // add_scalar. The caller states the declaration instead, which is why it arrives as a
+    // run-time argument rather than as add_scalar_one's template parameter.
+    template <typename T>
+    void add_scalars_impl(const T *values, int count, bool dynamic) {
+        static_assert(is_supported_scalar_arg_v<T>, "add_scalars: element type must be arithmetic or enum");
+        if (count < 0 || scalar_count_ + count > MaxS) {
+            set_error(scalar_cap_msg());
+            return;
+        }
+        if constexpr (std::is_same_v<std::remove_cv_t<T>, uint64_t>) {
+            memcpy(&scalars_[scalar_count_], values, static_cast<size_t>(count) * sizeof(uint64_t));
+        } else {
+            for (int i = 0; i < count; ++i) {
+                scalars_[scalar_count_ + i] = to_u64(values[i]);
+            }
+        }
+        for (int i = 0; i < count; ++i) {
+            scalar_inherited_[scalar_count_ + i] = dynamic ? static_cast<const void *>(&values[i]) : nullptr;
+        }
+#if SIMPLER_DFX
+        dump_arg_selection_.clear_scalar_metadata(scalar_count_, count);
+#endif
+        scalar_count_ += count;
+    }
+
 #if SIMPLER_DFX
     // No-arg dump(): mark every arg already added to this Arg.
     void mark_all_dump_args() {
@@ -681,6 +779,61 @@ struct Arg : TaskArgsTpl<TensorRef, uint64_t, MaxT, MaxS, TensorArgType> {
         }
         return true;
     }
+
+public:
+    LaunchSpec launch_spec;  // SPMD launch parameters (block_num, etc.)
+
+protected:
+    // The dependent base's storage. These declarations are what makes it reachable by
+    // unqualified lookup inside this template, which does not search a dependent base.
+    //
+    // Protected, not public: a derived Arg writes these slot by slot
+    // (GraphTaskArgs::gen_scalar_params_from_args), while every reader outside goes
+    // through an accessor that pairs each array with what qualifies it -- tensor() and
+    // tag() for the tensors, pack_scalars() and scalar(i) for the values and origins.
+    using Base::scalar_count_;
+    using Base::scalars_;
+    using Base::tags_;
+    using Base::tensor_count_;
+    using Base::tensors_;
+
+    // Where this parameter's value came from, or null when it is static.
+    //
+    // These addresses are compared and subtracted, never dereferenced -- one may point at
+    // a caller local that has already gone out of scope. That is safe precisely because
+    // the value was copied at add_scalar time.
+    std::array<const void *, MaxS> scalar_inherited_{};
+
+private:
+    bool has_error_{false};
+    const char *error_msg_{nullptr};
+
+    // Speculative early-dispatch hint (codegen-author set, off by default). When true,
+    // the scheduler may stage this task on an idle core before its producer finishes,
+    // gating execution on the DATA_MAIN_BASE doorbell — only safe when the author knows
+    // the task's data dependencies allow it. Read in-process by the runtime; never
+    // crosses the wire format.
+    bool allow_early_resolve_{false};
+
+    // Dispatch predicate (codegen-author set; default op == NONE = always dispatch). A
+    // FALSE result at the dispatch point retires the task inline through the dep-only
+    // path — never dispatched to an AICore — while still resolving fanin/fanout so
+    // consumers unlock. The predicate tensor's producer MUST be a dependency of this task
+    // so the value is current when the task becomes ready. Read in-process; never crosses
+    // the wire.
+    CoreTaskPredicate predicate_;
+
+    // Scheduler records this task's AICPU dispatch/finish cycles into fixed slot 0..15.
+    // TASK_TIMING_SLOT_NONE leaves it untagged.
+    int32_t task_timing_slot_{TASK_TIMING_SLOT_NONE};
+
+#if SIMPLER_DFX
+    DumpArgSelection dump_arg_selection_;
+#endif
+
+    // Caller-owned dependency array; lifetime must extend through submit.
+    const TaskId *explicit_deps_{nullptr};
+    uint32_t explicit_dep_count_{0};
 };
 
 // =============================================================================
@@ -695,12 +848,72 @@ using CoreTaskArgs = Arg<MAX_TENSOR_ARGS, MAX_SCALAR_ARGS>;
 inline constexpr int32_t GRAPH_MAX_TENSOR_ARGS = 128;
 inline constexpr int32_t GRAPH_MAX_SCALAR_ARGS = 64;
 
-// Boundary arguments of a Graph. Sized independently of CoreTaskArgs because the
-// outer GRAPH payload carries the whole boundary, while materialize stages only
-// one in-graph task's arguments at a time. The compact boundary values live in that
-// payload's argument-pool regions, so widening these caps costs pool bytes only
-// for Graphs that use them; TaskPayload itself stays fixed-size.
-using GraphTaskArgs = Arg<GRAPH_MAX_TENSOR_ARGS, GRAPH_MAX_SCALAR_ARGS>;
+/**
+ * A Graph's argument list.
+ *
+ * One type serves both sides of the call. A caller fills one with the arguments it is
+ * passing; the in-flight entry holds one that is the Graph's formal parameters. The two
+ * are separated by gen_scalar_params_from_args below -- before it, a slot's origin names
+ * wherever the caller's value came from; after it, a dynamic parameter names itself, and
+ * that is what a recorded body resolves against.
+ *
+ * Sized independently of CoreTaskArgs because the outer GRAPH payload carries the whole
+ * boundary, while materialize stages only one in-graph task's arguments at a time. The
+ * compact boundary values live in that payload's argument-pool regions, so widening these
+ * caps costs pool bytes only for Graphs that use them; TaskPayload itself stays
+ * fixed-size.
+ *
+ * A type of its own rather than an alias of Arg, because generating a parameter list is
+ * something only a Graph does: the one method below writes the slot arrays directly,
+ * which is why they are protected rather than private.
+ */
+struct GraphTaskArgs : Arg<GRAPH_MAX_TENSOR_ARGS, GRAPH_MAX_SCALAR_ARGS> {
+    /**
+     * Existing tensors only: a Graph boundary cannot take a runtime-allocated output.
+     *
+     * A Definition records the device addresses its body resolved against, so every
+     * boundary tensor must already own its buffer when the body is recorded and again on
+     * every replay. A TensorCreateInfo names a buffer the runtime would allocate at
+     * submit, which is a different address each time and none at record time.
+     */
+    template <typename... Args>
+    void add_output(Args &&...args) {
+        static_assert(
+            !(std::is_same_v<std::decay_t<Args>, TensorCreateInfo> || ...),
+            "a Graph boundary cannot take a runtime-allocated output (TensorCreateInfo); "
+            "allocate the tensor before the Graph and pass it as an existing Tensor"
+        );
+        Arg<GRAPH_MAX_TENSOR_ARGS, GRAPH_MAX_SCALAR_ARGS>::add_output(std::forward<Args>(args)...);
+    }
+
+    /**
+     * Build this Graph's scalar parameters from the arguments a caller passed.
+     *
+     * Values are resolved straight into the slots, because an argument's origin does not
+     * outlive the submit call that lent it. The declaration does carry over, and a dynamic
+     * parameter ends up naming **itself** -- which is what keeps this list the basis
+     * recording resolves against: scalar(i) folds to &scalars_[i] whether the parameter is
+     * dynamic or static, so a task slot that follows parameter i reports this array's i-th
+     * slot and graph_classify_scalars turns that into the index i.
+     *
+     * Naming the caller's variable instead would hand out an address outside this array,
+     * the task slot would be recorded as static, and the parameter would silently stop
+     * being refreshed on replay.
+     */
+    void gen_scalar_params_from_args(const GraphTaskArgs &args) {
+        // The whole list is generated at once, into an Arg that has none yet -- graph_begin
+        // calls this on an object it has just built. Appending to an existing list has no
+        // meaning here: these are the Graph's parameters, not values accumulated by a
+        // caller. Both sides are GraphTaskArgs, so the source can never overflow this.
+        debug_assert(scalar_count_ == 0 && "a parameter list is generated whole, not appended to");
+        const int32_t count = args.scalar_count();
+        args.pack_scalars(scalars_);
+        for (int32_t i = 0; i < count; ++i) {
+            scalar_inherited_[i] = args.scalar_dynamic(i) ? &scalars_[i] : nullptr;
+        }
+        scalar_count_ = count;
+    }
+};
 
 // ChipTaskArgs — chip-level entry-arg holding the orchestration entry's
 // already-allocated inputs (capacity matches simpler::hbg::EntryArgsStorage).
diff --git a/tests/st/a2a3/host_build_graph/paged_attention_unroll/kernels/orchestration/paged_attention_orch.cpp b/tests/st/a2a3/host_build_graph/paged_attention_unroll/kernels/orchestration/paged_attention_orch.cpp
index 4d75b920c3..36e6556d29 100644
--- a/tests/st/a2a3/host_build_graph/paged_attention_unroll/kernels/orchestration/paged_attention_orch.cpp
+++ b/tests/st/a2a3/host_build_graph/paged_attention_unroll/kernels/orchestration/paged_attention_orch.cpp
@@ -111,7 +111,7 @@ static void process_qtile_scope(const CoreTaskArgs &ctx) {
     uint64_t scale_value = ctx.scalar(7);
     uint64_t bn_this_batch = ctx.scalar(8);
     uint64_t cur_seq = ctx.scalar(9);
-    DataType data_type = static_cast<DataType>(ctx.scalar(10));
+    DataType data_type = ctx.scalar(10).to<DataType>();
 
     CYCLE_COUNT_START();
 
diff --git a/tests/ut/cpp/common/test_hbg_graph_async_submit.cpp b/tests/ut/cpp/common/test_hbg_graph_async_submit.cpp
index 8370068336..18e6c694ef 100644
--- a/tests/ut/cpp/common/test_hbg_graph_async_submit.cpp
+++ b/tests/ut/cpp/common/test_hbg_graph_async_submit.cpp
@@ -36,6 +36,14 @@ RuntimeContext *g_bound_runtime = nullptr;
 extern "C" RuntimeContext *framework_current_runtime(void) { return g_bound_runtime; }
 extern "C" void framework_bind_runtime(RuntimeContext *rt) { g_bound_runtime = rt; }
 
+// One recording's formal parameters and the tensor storage its TensorRefs point at.
+// The two travel together because gen_scalar_params_from_args resolves against storage
+// this object owns, the way an entry's GraphBoundary does.
+struct FakeBoundary {
+    std::array<simpler::hbg::Tensor, GRAPH_MAX_TENSOR_ARGS> tensors{};
+    GraphTaskArgs params;
+};
+
 struct FakeRuntime {
     const RuntimeOps *ops;
     ScopeMode pending_scope_mode{ScopeMode::AUTO};
@@ -68,8 +76,8 @@ struct FakeRuntime {
     std::thread::id prepare_thread;
     std::thread::id submit_thread;
 
-    // What graph_prepare actually received, so the deep copy GraphOwnedArgs makes
-    // can be compared against the boundary the caller passed.
+    // What graph_prepare actually received, so the parameter list the body resolves
+    // against can be compared against the boundary the caller passed.
     bool prepare_saw_args{false};
     const void *prepare_handle{nullptr};
     int32_t recorded_tensor_count{-1};
@@ -81,6 +89,15 @@ struct FakeRuntime {
     TensorArgType recorded_tag{};
     const void *recorded_args_object{nullptr};
     const void *recorded_tensor_storage{nullptr};
+
+    // Stands in for the in-flight entry's own copy of the formal parameters: the real
+    // graph_begin builds one per entry and hands it out through GraphScopeResult, and the
+    // recorded body reads that rather than the caller's arguments. One per recording, so
+    // a queued body still reads what its own begin built after a later begin has run.
+    // The storage is owned by the test, not by this struct, which is what keeps
+    // FakeRuntime standard-layout so the offsetof guards below stay well-defined.
+    FakeBoundary *boundaries{nullptr};
+    int32_t boundary_capacity{0};
 };
 
 static_assert(offsetof(FakeRuntime, ops) == 0);
@@ -90,15 +107,40 @@ FakeRuntime *as_fake(RuntimeContext *rt) { return reinterpret_cast<FakeRuntime *
 
 bool fake_is_fatal(RuntimeContext *rt) { return as_fake(rt)->fatal.load(std::memory_order_acquire); }
 
-GraphScopeResult fake_graph_begin(RuntimeContext *rt, uint64_t, const GraphTaskArgs &) {
+GraphScopeResult fake_graph_begin(RuntimeContext *rt, uint64_t, const GraphTaskArgs &args) {
     FakeRuntime &fake = *as_fake(rt);
     std::lock_guard<std::mutex> lock(fake.mutex);
     fake.begin_calls++;
     GraphScopeResult result;
     result.execute_block = false;
     result.recording = fake.record_every_begin || fake.begin_calls == 1;
-    // The handle the recording thread must hand back to graph_prepare.
-    if (result.recording) result.recording_handle = &fake;
+    if (result.recording) {
+        // Each recording resolves against its own list, so a body queued by an earlier
+        // begin is unaffected by a later one -- as with an entry's own GraphBoundary.
+        // A non-void return rules out ASSERT_, and EXPECT_ would record the failure and
+        // then index past the array anyway, so a test owning too few boundaries reports
+        // no recording instead.
+        if (fake.begin_calls > fake.boundary_capacity) {
+            ADD_FAILURE() << "test must own one boundary per recording";
+            result.recording = false;
+            return result;
+        }
+        FakeBoundary &boundary = fake.boundaries[fake.begin_calls - 1];
+        // Deep-copy the parameters the way graph_begin does: tensors first, so the
+        // TensorRefs that follow point at storage this object owns, then resolved values.
+        boundary.params.reset();
+        for (int32_t i = 0; i < args.tensor_count(); ++i) {
+            boundary.tensors[static_cast<size_t>(i)] = args.tensor(i).ref();
+        }
+        for (int32_t i = 0; i < args.tensor_count(); ++i) {
+            boundary.params.add_input(boundary.tensors[static_cast<size_t>(i)]);
+        }
+        boundary.params.gen_scalar_params_from_args(args);
+        // The handle the recording thread must hand back to graph_prepare, and the
+        // parameters the recorded body reads.
+        result.recording_handle = &fake;
+        result.params = &boundary.params;
+    }
     if (!result.recording) {
         fake.later_submit_entered = true;
         fake.cv.notify_all();
@@ -125,7 +167,7 @@ bool fake_graph_prepare(RuntimeContext *rt, void *recording_handle, const GraphT
         fake.recorded_tensor_storage = &tensor;
     }
     if (args.scalar_count() > 0) {
-        fake.recorded_scalar = args.scalar(0);
+        fake.recorded_scalar = args.scalar(0).to<uint64_t>();
     }
     if (fake.gate_four_prepares) {
         fake.cv.notify_all();
@@ -182,7 +224,7 @@ GraphAsyncRecordingState &test_pool() {
 }
 
 bool fake_graph_record_start(RuntimeContext *, const GraphTaskArgs &args, void *job) {
-    auto *record = static_cast<std::function<void(GraphTaskArgs &)> *>(job);
+    auto *record = static_cast<std::function<void(const GraphTaskArgs &)> *>(job);
     return test_pool().start(args, std::move(*record));
 }
 
@@ -217,7 +259,7 @@ TEST(HbgGraphAsyncSubmit, PrewarmedRecorderPoolGrowsPastThePrewarmedCount) {
     GraphTaskArgs empty_args;
 
     for (int i = 0; i < kGraphCount; ++i) {
-        ASSERT_TRUE(pool.start(empty_args, [&](GraphTaskArgs &) {
+        ASSERT_TRUE(pool.start(empty_args, [&](const GraphTaskArgs &) {
             std::unique_lock<std::mutex> lock(gate_mutex);
             worker_ids.insert(std::this_thread::get_id());
             entered++;
@@ -245,8 +287,11 @@ TEST(HbgGraphAsyncSubmit, PrewarmedRecorderPoolGrowsPastThePrewarmedCount) {
 }
 
 TEST(HbgGraphAsyncSubmit, FourDistinctGraphMissesDoNotInsertAnIntermediateCommit) {
+    std::array<FakeBoundary, 4> boundaries;
     FakeRuntime fake{};
     fake.ops = &kFakeOps;
+    fake.boundaries = boundaries.data();
+    fake.boundary_capacity = static_cast<int32_t>(boundaries.size());
     fake.record_every_begin = true;
     fake.gate_four_prepares = true;
     framework_bind_runtime(reinterpret_cast<RuntimeContext *>(&fake));
@@ -278,8 +323,11 @@ TEST(HbgGraphAsyncSubmit, FourDistinctGraphMissesDoNotInsertAnIntermediateCommit
 }
 
 TEST(HbgGraphAsyncSubmit, WorkerRecordsWhileMainSubmitsLaterGraphs) {
+    std::array<FakeBoundary, 4> boundaries;
     FakeRuntime fake{};
     fake.ops = &kFakeOps;
+    fake.boundaries = boundaries.data();
+    fake.boundary_capacity = static_cast<int32_t>(boundaries.size());
     framework_bind_runtime(reinterpret_cast<RuntimeContext *>(&fake));
 
     uint32_t storage[4]{};
@@ -370,8 +418,11 @@ TEST(HbgGraphAsyncSubmit, WorkerRecordsWhileMainSubmitsLaterGraphs) {
 // pins both halves of its contract — the values survive, and the storage they
 // live in is not the caller's.
 TEST(HbgGraphAsyncSubmit, RecordingReadsAnOwnedCopyOfTheBoundary) {
+    std::array<FakeBoundary, 4> boundaries;
     FakeRuntime fake{};
     fake.ops = &kFakeOps;
+    fake.boundaries = boundaries.data();
+    fake.boundary_capacity = static_cast<int32_t>(boundaries.size());
     framework_bind_runtime(reinterpret_cast<RuntimeContext *>(&fake));
 
     uint32_t storage[8]{};
@@ -427,8 +478,11 @@ TEST(HbgGraphAsyncSubmit, RecordingReadsAnOwnedCopyOfTheBoundary) {
 // directly: both live in the wrapper, which is why this case is here and not with the
 // orchestrator's own graph tests.
 TEST(HbgGraphAsyncSubmit, AFatalInsideARecordedBodyReachesGraphEndAndAbortsNothing) {
+    std::array<FakeBoundary, 4> boundaries;
     FakeRuntime fake{};
     fake.ops = &kFakeOps;
+    fake.boundaries = boundaries.data();
+    fake.boundary_capacity = static_cast<int32_t>(boundaries.size());
     framework_bind_runtime(reinterpret_cast<RuntimeContext *>(&fake));
 
     uint32_t storage[4]{};
diff --git a/tests/ut/cpp/common/test_hbg_graph_cache.cpp b/tests/ut/cpp/common/test_hbg_graph_cache.cpp
index 1f015f561b..a17356b8af 100644
--- a/tests/ut/cpp/common/test_hbg_graph_cache.cpp
+++ b/tests/ut/cpp/common/test_hbg_graph_cache.cpp
@@ -88,10 +88,10 @@ make_test_definition(uint64_t graph_key, uint64_t boundary_address, uint32_t bou
     tensor_sources[1].source_kind = static_cast<uint8_t>(GraphTensorSourceKind::INTERNAL);
     tensor_sources[1].packed_offset = 16;
     std::vector<uint64_t> scalars{0, 18};
-    std::vector<GraphScalarSourceRef> scalar_sources(2);
-    scalar_sources[0].source_kind = static_cast<uint8_t>(GraphScalarSourceKind::BOUNDARY);
-    scalar_sources[0].source_index = boundary_scalar_count - 1;
-    scalar_sources[1].source_kind = static_cast<uint8_t>(GraphScalarSourceKind::STATIC_VALUE);
+    std::vector<GraphScalarInheritance> scalar_inheritance{
+        GraphScalarInheritance::from_boundary(boundary_scalar_count - 1),
+        GraphScalarInheritance::self_value(),
+    };
 
     GraphDefinition definition{};
     definition.full_key = graph_key;
@@ -113,7 +113,7 @@ make_test_definition(uint64_t graph_key, uint64_t boundary_address, uint32_t bou
     definition.off_tensors = append_section(image, tensors);
     definition.off_tensor_sources = append_section(image, tensor_sources);
     definition.off_scalars = append_section(image, scalars);
-    definition.off_scalar_sources = append_section(image, scalar_sources);
+    definition.off_scalar_inheritance = append_section(image, scalar_inheritance);
     size_t execution_storage_bytes = 0;
     graph_execution_storage_bytes(
         definition.task_count, definition.tensor_arg_count, definition.scalar_arg_count, &execution_storage_bytes
@@ -279,34 +279,179 @@ TEST(GraphCache, ConfigValuesSelectDifferentDefinitions) {
     EXPECT_EQ(rt_graph_make_key(GRAPH_ID, 0), rt_graph_make_key(GRAPH_ID, 0));
 }
 
+// Arg's storage stays unreachable only while the base is private. Under a public base an
+// implicit derived-to-base conversion reaches the same subobject, whose members are
+// public there however Arg hides their names -- so tags_, tensors_ and scalars_ would be
+// readable raw, each without the array that qualifies it.
+static_assert(
+    !std::is_convertible_v<
+        const CoreTaskArgs *,
+        const TaskArgsTpl<TensorRef, uint64_t, MAX_TENSOR_ARGS, MAX_SCALAR_ARGS, TensorArgType> *>,
+    "Arg must not be convertible to its storage base"
+);
+
+// The slot array is handed out as void*, so index it the way recording does.
+const void *slot_addr(const void *base, int32_t i) { return static_cast<const uint64_t *>(base) + i; }
+
+// A boundary parameter must name no origin of its own. That is what makes the boundary's
+// slot array the basis recording resolves against: scalar(i) then folds to
+// &scalars_[i], and subtracting the base yields i. Were a boundary parameter to name an
+// origin -- a caller local, or a scratch buffer used while deep-copying -- scalar(i)
+// would hand out that address instead, it would fall outside the slot array, and
+// graph_classify_scalars would record the parameter as static. Every dynamic parameter
+// would silently stop being refreshed on replay.
+TEST(GraphScalarProvenance, AStaticParameterIsItsOwnOrigin) {
+    GraphTaskArgs boundary_args;
+    uint64_t scratch = 18;
+    boundary_args.add_static_scalar(scratch);
+    CoreTaskArgs task_args;
+
+    task_args.add_scalar(boundary_args.scalar(0));
+
+    EXPECT_EQ(task_args.scalar_origin(0), boundary_args.scalar_slot_base())
+        << "a task slot must name the boundary's own slot, not whatever the value came from";
+    EXPECT_NE(task_args.scalar_origin(0), &scratch);
+}
+
+TEST(GraphScalarProvenance, AnLvalueDeclaresADynamicParameter) {
+    GraphTaskArgs args;
+    uint32_t token_pos = 17;
+
+    // An lvalue is a parameter the caller holds and may change between invocations; a
+    // literal cannot be changed by anyone. Width does not enter into it -- the value is
+    // converted here, where its type is still known.
+    args.add_scalar(token_pos);
+    args.add_scalar(uint32_t{18});
+    args.add_static_scalar(token_pos);
+
+    EXPECT_TRUE(args.scalar_dynamic(0));
+    EXPECT_EQ(args.scalar_origin(0), &token_pos);
+    EXPECT_FALSE(args.scalar_dynamic(1));
+    EXPECT_EQ(args.scalar_origin(1), nullptr);
+    EXPECT_FALSE(args.scalar_dynamic(2)) << "add_static_scalar overrides value category";
+    EXPECT_EQ(args.scalar(2).to<uint32_t>(), 17u);
+}
+
 TEST(GraphScalarProvenance, ForwardedScalarRetainsBoundarySource) {
-    uint32_t value = 17;
-    CoreTaskArgs boundary_args;
-    boundary_args.add_scalar(value, value);
-    boundary_args.anchor_scalar_sources();
+    GraphTaskArgs boundary_args;
+    boundary_args.add_scalar(uint32_t{17}, uint32_t{18});
+    CoreTaskArgs task_args;
+
+    task_args.add_scalar(boundary_args.scalar(1));
+
+    EXPECT_TRUE(task_args.scalar_dynamic(0));
+    EXPECT_EQ(task_args.scalar_origin(0), slot_addr(boundary_args.scalar_slot_base(), 1));
+    EXPECT_EQ(task_args.scalar(0).to<uint64_t>(), uint64_t{18});
+}
+
+TEST(GraphScalarProvenance, ForwardedScalarNamesOriginThroughAnIntermediary) {
+    GraphTaskArgs boundary_args;
+    boundary_args.add_scalar(uint32_t{17}, uint32_t{18});
     CoreTaskArgs forwarded_args;
-    forwarded_args.copy_scalars_from(boundary_args, 1, 1);
+    forwarded_args.add_scalar(boundary_args.scalar(1));
     CoreTaskArgs task_args;
 
-    task_args.copy_scalars_from(forwarded_args, 0, 1);
+    task_args.add_scalar(forwarded_args.scalar(0));
 
-    EXPECT_EQ(task_args.scalar_source(0), static_cast<const void *>(&std::as_const(boundary_args).scalar(1)));
+    // A -> B -> C still records A: the handle names the origin, not the Arg it came through.
+    EXPECT_TRUE(task_args.scalar_dynamic(0));
+    EXPECT_EQ(task_args.scalar_origin(0), slot_addr(boundary_args.scalar_slot_base(), 1));
 }
 
-TEST(GraphScalarProvenance, MutableAccessInvalidatesForwardedSource) {
-    CoreTaskArgs boundary_args;
-    boundary_args.add_scalar(uint32_t{17});
-    boundary_args.anchor_scalar_sources();
+TEST(GraphScalarProvenance, FreezingAParameterDropsItsOrigin) {
+    GraphTaskArgs boundary_args;
+    boundary_args.add_scalar(uint32_t{17}, uint32_t{18});
     CoreTaskArgs task_args;
-    task_args.copy_scalars_from(boundary_args, 0, 1);
-    ASSERT_NE(task_args.scalar_source(0), nullptr);
 
-    task_args.scalar(0) = 18;
+    // add_static_scalar resolves the handle to its value: this is how an enclosing Graph's
+    // parameter is deliberately frozen rather than followed.
+    task_args.add_static_scalar(boundary_args.scalar(1));
 
-    EXPECT_EQ(task_args.scalar_source(0), nullptr);
-    EXPECT_EQ(
-        task_args.invalidated_scalar_source(0), static_cast<const void *>(&std::as_const(boundary_args).scalar(0))
-    );
+    EXPECT_FALSE(task_args.scalar_dynamic(0));
+    EXPECT_EQ(task_args.scalar_origin(0), nullptr);
+    EXPECT_EQ(task_args.scalar(0).to<uint64_t>(), uint64_t{18});
+}
+
+TEST(GraphScalarProvenance, ValueScalarHoldsItsOwnValue) {
+    CoreTaskArgs task_args;
+    task_args.add_scalar(uint32_t{17}, 2.5F);
+
+    EXPECT_FALSE(task_args.scalar_dynamic(0));
+    EXPECT_FALSE(task_args.scalar_dynamic(1));
+    EXPECT_EQ(task_args.scalar(0).to<uint64_t>(), uint64_t{17});
+    EXPECT_EQ(task_args.scalar(1).to<float>(), 2.5F);
+}
+
+TEST(GraphScalarProvenance, ZeroInitialisedSlotsReadAsStaticZero) {
+    CoreTaskArgs task_args;
+    uint64_t packed[4] = {1, 2, 3, 4};
+
+    task_args.add_scalar(uint64_t{0});
+    task_args.pack_scalars(packed);
+
+    EXPECT_FALSE(task_args.scalar_dynamic(0));
+    EXPECT_EQ(packed[0], uint64_t{0});
+}
+
+TEST(GraphScalarProvenance, PackCopiesEveryValue) {
+    GraphTaskArgs boundary_args;
+    boundary_args.add_scalar(uint64_t{100}, uint64_t{200});
+    CoreTaskArgs task_args;
+    // A slot is always a value, whether or not it names an origin, so a mixed run copies
+    // out whole.
+    task_args.add_scalar(uint64_t{7});
+    task_args.add_scalar(boundary_args.scalar(1));
+    task_args.add_scalar(uint64_t{9});
+    uint64_t packed[3] = {0, 0, 0};
+
+    task_args.pack_scalars(packed);
+
+    EXPECT_EQ(packed[0], uint64_t{7});
+    EXPECT_EQ(packed[1], uint64_t{200});
+    EXPECT_EQ(packed[2], uint64_t{9});
+}
+
+TEST(GraphScalarProvenance, AValueOutlivesItsOrigin) {
+    CoreTaskArgs task_args;
+    {
+        uint64_t transient = 42;
+        task_args.add_scalar(transient);
+    }
+
+    // The origin now dangles, and that is by design: the value was copied at add_scalar
+    // time, and the address is only ever compared, never read through.
+    EXPECT_TRUE(task_args.scalar_dynamic(0));
+    EXPECT_EQ(task_args.scalar(0).to<uint64_t>(), uint64_t{42});
+}
+
+TEST(GraphScalarProvenance, TaskSlotsInheritFromEachOther) {
+    GraphTaskArgs boundary_args;
+    boundary_args.add_scalar(uint32_t{17}, uint32_t{18});
+    CoreTaskArgs a;
+    a.add_scalar(boundary_args.scalar(1));
+    CoreTaskArgs b;
+
+    // Any slot is an inheritance source, not just a boundary parameter.
+    b.add_scalar(a.scalar(0));
+
+    // scalar(i) folds: a[0] already names an origin, so b records that origin rather than
+    // a[0] itself. A chain is therefore one hop and recording resolves it without a walk.
+    EXPECT_TRUE(b.scalar_dynamic(0));
+    EXPECT_EQ(b.scalar_origin(0), slot_addr(boundary_args.scalar_slot_base(), 1));
+    EXPECT_EQ(b.scalar(0).to<uint64_t>(), uint64_t{18});
+}
+
+TEST(GraphScalarProvenance, InheritingAValueSlotNamesThatSlot) {
+    CoreTaskArgs a;
+    a.add_scalar(uint64_t{7});
+    CoreTaskArgs b;
+
+    b.add_scalar(a.scalar(0));
+
+    // a[0] names no origin, so it is itself the origin.
+    EXPECT_TRUE(b.scalar_dynamic(0));
+    EXPECT_EQ(b.scalar_origin(0), a.scalar_slot_base());
+    EXPECT_EQ(b.scalar(0).to<uint64_t>(), uint64_t{7});
 }
 
 TEST(GraphExecutionStorage, ComputesAlignedExactSize) {