From c0fa4e1ad2ab8f3ac9fcf08736a09b58974633be Mon Sep 17 00:00:00 2001 From: poursoul Date: Wed, 9 Sep 2026 01:17:19 -0700 Subject: [PATCH] Refactor: make a Graph boundary scalar a formal parameter A body that read a boundary scalar and one that read a constant were the same thing in the type system: both a uint64_t in the same slot array. Recording recovered the difference by comparing the host address the caller happened to pass against the boundary's slot range, so a body that loaded a parameter into a local and forwarded the local recorded a constant instead, silently, and every later replay of that Definition reused the stale value. A scalar now carries two things in parallel arrays: its value, and the slot that value came from. Null origin means static. The value is converted by to_u64 at add_scalar time, where the argument's type is still known, so a parameter of any width can be dynamic -- a representation that packed both into one eight-byte union could not, because following a parameter would then mean reading through the origin in pack_scalars, where the type is long gone. The origin is only ever compared against null and subtracted from a slot array base, never dereferenced, so it may dangle once the caller's local goes out of scope; the value was copied long before. What declares a parameter dynamic is the argument's value category: add_scalar of an lvalue (the caller holds it and may change it) or of an InheritableScalar (it already names a parameter) is dynamic, a literal is static, and add_static_scalar says so explicitly whatever was passed. InheritableScalar has no public getter for its bits: the only place that reads one is add_scalar_one, forwarding a parameter into a destination slot, and it does that through to() rather than a bits() accessor. A public bits() would have been a second silent value-read path next to the deprecated conversion -- exactly what that conversion exists to close off. InheritableScalar carries value and origin together, so a destination stores the value without following the origin. Forwarding one keeps the parameter; reading its number goes through a deprecated conversion, so the point where a body freezes a parameter is a compiler diagnostic instead of a silent change of behaviour. to() is the deliberate read: it applies to_u64's actual inverse, and it is the only spelling that reaches an enum, since a conversion to an enumeration does not accept a user-defined one on the way. One limit on that diagnostic: 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 is the case found here. The warning is an inventory of the value reads written in this repo, which is what the migration needs, not a proof that none exists. A Graph's parameters are built by GraphTaskArgs::gen_scalar_params_from_args, and a dynamic parameter there names *itself*. That is the invariant the scheme rests on: scalar(i) folds to &scalars_[i] whether a parameter is dynamic or static, so a task slot following parameter i reports that array's i-th slot and recording turns it into the index i. A parameter naming the caller's variable instead would hand out an address outside the array, the task slot would be recorded as static, and it would stop being refreshed on replay with no diagnostic at all. GraphTaskArgs also rejects a runtime-allocated output at compile time. A Definition records the device addresses its body resolved against, so a boundary tensor must own its buffer when the body is recorded and again on every replay, while a TensorCreateInfo names a buffer the runtime would allocate at submit -- a different address each time, and none at record time. The submit-time check that stood in for that constraint goes with it: rt_graph_args_cacheable no longer walks the tags looking for one, and the boundary builder's branch for the tag is unreachable rather than an error path that left a boundary short a tensor. GraphScalarInheritance is the wire form: recording knows whether a slot inherits and which parameter it inherits, so a flag and an index say both. Both are uint16_t, which leaves the type no padding -- it is memcpy'd into the Definition image, and a byte no writer sets is a byte the image carries without meaning. Its fields are private and set only together, through self_value() and from_boundary(), so an entry claiming to inherit while naming no parameter cannot be spelled. The index is still a claim about a boundary the entry cannot see, so the packer and materialize each bound it against the boundary they do have. GraphBoundary holds the parameter list rather than a pointer to one. That pointer named a slot in the recorder pool, which is reused, so it went stale the moment a job finished and nothing cleared it. Nothing read it stale: classify, layout and fill all run inside the job, and a later same-key submission compares tensors, types and scalar_count instead of following the pointer. The old arrangement was therefore correct by that timing rather than by ownership, and holding the list removes the dependency. GraphOwnedArgs and the pool's deep copy go with it -- the pool forwards a reference to the entry's own list, which graph_commit keeps alive until every recording has finished. The submitting thread makes one deep copy instead of two, and a boundary is written once in graph_begin and only read after that, by either thread. GraphScopeResult carries the list out, because the entry type is private to orchestrator.cpp; the body reads it even on the synchronous fallback path, since reading the caller's arguments there would classify every parameter as static. Arg's surface closes down to its API. Its storage base is private, because a public one is reachable by an implicit derived-to-base conversion through which the members are public again however the derived class hides their names; a static_assert holds that shut. has_error and error_msg become accessors, the slot arrays are protected so a Graph can build a parameter list, and Arg is a class whose functions and data sit in separate runs. Removed along the way: the non-const Arg::scalar(), which had no legal caller -- a body that wrote through it marked the whole recording unsupported, and the only in-repo user was the test written for that rejection -- along with scalar_sources_invalidated_, invalidated_scalar_source() and the INVALIDATED_BOUNDARY source kind; copy_scalars_from, since add_scalar(args.scalar(i)) covers it and an InheritableScalar crosses two Arg capacities without either naming the other; and add_scalars_i32, since add_scalars is a template now. pack_scalars is one memcpy: a slot is a value, so there is nothing to resolve and no discriminator to scan. graph_prepare's boundary-consistency assertion goes as well. Both callers now forward the entry's own parameter list, so it compared an object with itself and held whatever the boundary had become. The check it stood for still runs where a caller's arguments do arrive against an existing entry -- a same-key submission, compared in graph_begin_inner on a path release builds keep. Three orchestration entry points therefore change shape or disappear: copy_scalars_from and add_scalars_i32 are gone, and add_scalars becomes a template whose default declaration flips from static to dynamic. None of the three has a call site in pypto (a18c4cf9), pypto-lib (56c01e1), or this repo, so no caller moves with them. That flip is toward the unsafe side of the two declarations, which is worth naming even though nothing observes it yet. A forward-only scalar wrongly declared static only makes matching stricter; a scalar the body freezes while declared dynamic would match on a Definition holding a stale number. Both are inert today because graph_full_key is callable_hash and graph_key, so no lookup compares a scalar value at all -- #2170 tracks migrating the callers before one does. _run_subprocess grows an opt-in way to surface a successful compile's stderr: warnings.warn, which pytest reports with no flag given, where the DEBUG log line it had displays nothing -- pytest hides logger output below ERROR unless --log-cli-level is passed, and the resource scheduler's child processes do not inherit that option. No call site turns it on. The orchestration sources still hold the value reads #2170 migrates, so surfacing them would repeat that known inventory on every scene test and bury the reads written after it; the kernel toolchains carry pre-existing warnings of their own for the same reason. Emitting a warning can raise, under an error-level filter, from a compile that already produced its output, so _compile_to_bytes deletes that output from a finally rather than from the success path. GRAPH_EXECUTION.md drops the paragraphs describing the removed invalidation rule, and its examples name GraphTaskArgs for the boundary and CoreTaskArgs for the in-graph tasks -- no rt_submit_graph overload accepts the former as the latter, so neither example compiled as written. Its add_static_scalar paragraph had described an entry point no code defined; that entry point exists now, so the paragraph says what it does and how it differs from to(), which reads a value rather than forwarding one with its origin dropped. The GCC blind spot above is recorded there as well, since the doc otherwise reads as though the warning were a complete census. --- ...-08-host-orch-phase-tail-is-page-faults.md | 2 +- simpler_setup/kernel_compiler.py | 45 +- .../orchestration/arg_with_deps.h | 6 +- .../orchestration/orchestration_api.h | 24 +- .../orchestration/arg_with_deps.h | 6 +- .../orchestration/orchestration_api.h | 24 +- .../device/graph_execution.cpp | 29 +- .../host_build_graph/docs/GRAPH_EXECUTION.md | 108 ++-- src/common/host_build_graph/graph_cache.h | 21 +- src/common/host_build_graph/graph_execution.h | 53 +- .../host_build_graph/graph_recorder_pool.h | 115 +--- .../host/graph_recorder_pool.cpp | 6 +- .../host_build_graph/host/orchestrator.cpp | 347 ++++++----- src/common/host_build_graph/runtime_ops.h | 2 +- src/common/host_build_graph/runtime_types.h | 9 +- src/common/host_build_graph/types.h | 549 ++++++++++++------ .../orchestration/paged_attention_orch.cpp | 2 +- .../common/test_hbg_graph_async_submit.cpp | 70 ++- tests/ut/cpp/common/test_hbg_graph_cache.cpp | 191 +++++- 19 files changed, 1058 insertions(+), 551 deletions(-) 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 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 ¶ms = *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 job = [invoke, handle](GraphTaskArgs &record_args) mutable { + std::function 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 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 ¶ms = *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 job = [invoke, handle](GraphTaskArgs &record_args) mutable { + std::function 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 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(definition, definition.off_scalars, definition.scalar_arg_count); - const GraphScalarSourceRef *scalar_sources = - definition.scalar_arg_count == 0 ? nullptr : - graph_definition_array( - definition, definition.off_scalar_sources, definition.scalar_arg_count - ); + const GraphScalarInheritance *scalar_inheritance = + definition.scalar_arg_count == 0 ? + nullptr : + graph_definition_array( + definition, definition.off_scalar_inheritance, definition.scalar_arg_count + ); const GraphPredicate *predicates = definition.predicate_count == 0 ? nullptr : graph_definition_array(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(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(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(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()`. It +applies `to_u64`'s actual inverse, which `static_cast` is not — a float slot +holds a bit pattern, so `static_cast` of `1.0f`'s pattern yields +`1065353216.0`. It is also the only spelling that reaches an enum: +`static_cast(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()` 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 &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); static_assert(std::is_standard_layout_v); static_assert(std::is_trivially_copyable_v); static_assert(std::is_standard_layout_v); -static_assert(std::is_trivially_copyable_v); -static_assert(std::is_standard_layout_v); +static_assert(std::is_trivially_copyable_v); +static_assert(std::is_standard_layout_v); +static_assert(sizeof(GraphScalarInheritance) == 4, "the image's scalar section assumes this layout"); static_assert(std::is_trivially_copyable_v); static_assert(std::is_standard_layout_v); // graph_fill_definition assigns this struct field by field, so its interior padding @@ -262,7 +289,7 @@ static_assert(std::is_standard_layout_v); 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 @@ -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(i)].copy(source.tensor(i).ref()); - switch (source.tag(i)) { - case TensorArgType::INPUT: - args_.add_input(tensors_[static_cast(i)]); - break; - case TensorArgType::OUTPUT_EXISTING: - args_.add_output(tensors_[static_cast(i)]); - break; - case TensorArgType::INOUT: - args_.add_inout(tensors_[static_cast(i)]); - break; - case TensorArgType::NO_DEP: - args_.add_no_dep(tensors_[static_cast(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(i)] = source.scalar(i); - args_.add_scalar(scalars_[static_cast(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 tensors_{}; - std::array 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 bool start(const GraphTaskArgs &args, Job &&job) { - std::function next; + std::function next; try { next = std::forward(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 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 function; - size_t owned_args_index{0}; + std::function 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 workers_; std::mutex mutex_; std::condition_variable cv_; - std::array owned_args_; - std::array free_owned_args_{}; std::array 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 *>(job); + auto *record = static_cast *>(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 tensors; - std::vector types; + // Deliberately user-provided rather than `= default`: a defaulted constructor here is + // trivial, so make_unique()'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 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 ¶ms.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 tensor_sources; std::vector scalars; - std::vector 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 scalar_inheritance; std::vector 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 &boundary_tensors() const { return boundary->tensors; } - const std::vector &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(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 ¶m_tensor = + recording.bound_boundary().params.tensor(static_cast(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 -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(&args) == static_cast(recording.boundary_args()) && - scalar_index < recording.boundary_args()->scalar_count()) { - return GraphRecordedScalarSourceRef{GraphRecordedScalarSourceKind::BOUNDARY, static_cast(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(&recording.boundary_args()->scalar(i)); - if (source == boundary_source) { - return GraphRecordedScalarSourceRef{GraphRecordedScalarSourceKind::BOUNDARY, static_cast(i)}; +void graph_classify_scalars(GraphRecording &recording, const ArgT &args, int32_t scalar_offset) { + const GraphTaskArgs ¶ms = recording.bound_boundary().params; + const uintptr_t base = reinterpret_cast(params.scalar_slot_base()); + const uintptr_t span = static_cast(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(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(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((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(CORE_MAX_TENSOR_ARGS)); recording.scalars.reserve(kInGraphTaskCap * static_cast(CORE_MAX_SCALAR_ARGS)); - recording.scalar_sources.reserve(kInGraphTaskCap * static_cast(CORE_MAX_SCALAR_ARGS)); + recording.scalar_inheritance.reserve(kInGraphTaskCap * static_cast(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(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(i)} + {recording.bound_boundary().params.tensor(i).ref().buffer.addr, static_cast(i)} ); } std::stable_sort( @@ -952,19 +972,6 @@ std::optional graph_pack_tensor_source(const GraphRecorded return packed; } -std::optional 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(GraphScalarSourceKind::BOUNDARY) : - static_cast(GraphScalarSourceKind::STATIC_VALUE); - packed.source_index = static_cast(source.source_index); - return packed; -} - template 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 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 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(recording.tensor_sources.size()); const int32_t recorded_scalars = static_cast(recording.scalars.size()); - const int32_t recorded_scalar_sources = static_cast(recording.scalar_sources.size()); + const int32_t recorded_scalar_inheritance = static_cast(recording.scalar_inheritance.size()); const int32_t recorded_fanins = static_cast(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 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 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(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 graph_layout_definition(const GraphRecording &rec !graph_layout_section(total_tensors, &image_bytes, &definition.off_tensors) || !graph_layout_section(total_tensors, &image_bytes, &definition.off_tensor_sources) || !graph_layout_section(total_scalars, &image_bytes, &definition.off_scalars) || - !graph_layout_section(total_scalars, &image_bytes, &definition.off_scalar_sources) || + !graph_layout_section( + total_scalars, &image_bytes, &definition.off_scalar_inheritance + ) || !graph_layout_section( - recording.boundary_tensors().size(), &image_bytes, &definition.off_boundary_signatures + static_cast(recording.bound_boundary().params.tensor_count()), &image_bytes, + &definition.off_boundary_signatures ) || !graph_layout_section(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(image, definition.off_tensors); auto *tensor_sources = graph_image_section(image, definition.off_tensor_sources); auto *scalars = graph_image_section(image, definition.off_scalars); - auto *scalar_sources = graph_image_section(image, definition.off_scalar_sources); + auto *scalar_inheritance = graph_image_section(image, definition.off_scalar_inheritance); auto *signatures = graph_image_section(image, definition.off_boundary_signatures); auto *predicates = graph_image_section(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 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(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(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(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(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(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(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(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(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(i); uint16_t actual_alias = static_cast(i); for (int32_t j = 0; j < i; ++j) { - const simpler::hbg::Tensor &expected_other = boundary.tensors[static_cast(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(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(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(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(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(); entry->full_key = full_key; - entry->boundary.scalar_count = args.scalar_count(); - entry->boundary.tensors.reserve(static_cast(args.tensor_count())); - entry->boundary.types.reserve(static_cast(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(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(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(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 *` the pool moves out of -- + // `job` is a `std::function *` 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 +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 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 + T to() const { + return from_u64(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() 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 +inline constexpr bool is_inheritable_scalar_v = std::is_same_v, InheritableScalar>; + /** * Aggregated argument container for rt_submit_task * * Inherits storage from TaskArgsTpl. * 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 -struct Arg : TaskArgsTpl { +class Arg : private TaskArgsTpl { using Base = TaskArgsTpl; - // 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(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 { 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 { 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 { 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 { 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 void add_scalar(Args &&...args) { static_assert(sizeof...(Args) >= 1, "add_scalar: at least one argument required"); - static_assert((is_supported_scalar_arg_v && ...), "add_scalar: all types must be arithmetic or enum"); + static_assert( + ((is_supported_scalar_arg_v || is_inheritable_scalar_v) && ...), + "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)), ...); - } - - 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(std::forward(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 + 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 || is_inheritable_scalar_v) && ...), + "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(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(static_cast(values[i])); - } -#else - for (int i = 0; i < count; i++) { - dst[i] = static_cast(static_cast(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(std::forward(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(&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 + 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(&scalars_[i]); - scalar_sources_invalidated_[i] = false; - } + template + 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(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 scalar_sources_{}; - // Kept separately so invalidated boundary ancestry can reject Graph caching - // instead of being mistaken for an unrelated static scalar. - mutable std::array 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 static constexpr bool is_supported_dump_arg_v = @@ -571,36 +620,85 @@ struct Arg : TaskArgsTpl { is_supported_scalar_arg_v; #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 + // 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 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) { - scalar_sources_[scalar_count_] = static_cast(&value); - } + if constexpr (is_inheritable_scalar_v) { + // 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() 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(); + scalar_inherited_[scalar_count_] = Dynamic ? value.origin() : nullptr; #if SIMPLER_DFX - uintptr_t scalar_source_ptr = 0; - scalar_source_ptr = reinterpret_cast(scalar_sources_[scalar_count_]); - dump_arg_selection_.record_scalar_source( - scalar_count_, scalar_source_ptr, dtype_of>>() - ); + // 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()); #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) { + 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) { + scalar_source_ptr = reinterpret_cast(&value); + } + dump_arg_selection_.record_scalar_source( + scalar_count_, scalar_source_ptr, dtype_of>>() + ); +#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 + void add_scalars_impl(const T *values, int count, bool dynamic) { + static_assert(is_supported_scalar_arg_v, "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, uint64_t>) { + memcpy(&scalars_[scalar_count_], values, static_cast(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(&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 { } 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 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; 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; +/** + * 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 { + /** + * 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 + void add_output(Args &&...args) { + static_assert( + !(std::is_same_v, 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::add_output(std::forward(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(ctx.scalar(10)); + DataType data_type = ctx.scalar(10).to(); 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 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_castfatal.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 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(i)] = args.tensor(i).ref(); + } + for (int32_t i = 0; i < args.tensor_count(); ++i) { + boundary.params.add_input(boundary.tensors[static_cast(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(); } 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 *>(job); + auto *record = static_cast *>(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 lock(gate_mutex); worker_ids.insert(std::this_thread::get_id()); entered++; @@ -245,8 +287,11 @@ TEST(HbgGraphAsyncSubmit, PrewarmedRecorderPoolGrowsPastThePrewarmedCount) { } TEST(HbgGraphAsyncSubmit, FourDistinctGraphMissesDoNotInsertAnIntermediateCommit) { + std::array boundaries; FakeRuntime fake{}; fake.ops = &kFakeOps; + fake.boundaries = boundaries.data(); + fake.boundary_capacity = static_cast(boundaries.size()); fake.record_every_begin = true; fake.gate_four_prepares = true; framework_bind_runtime(reinterpret_cast(&fake)); @@ -278,8 +323,11 @@ TEST(HbgGraphAsyncSubmit, FourDistinctGraphMissesDoNotInsertAnIntermediateCommit } TEST(HbgGraphAsyncSubmit, WorkerRecordsWhileMainSubmitsLaterGraphs) { + std::array boundaries; FakeRuntime fake{}; fake.ops = &kFakeOps; + fake.boundaries = boundaries.data(); + fake.boundary_capacity = static_cast(boundaries.size()); framework_bind_runtime(reinterpret_cast(&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 boundaries; FakeRuntime fake{}; fake.ops = &kFakeOps; + fake.boundaries = boundaries.data(); + fake.boundary_capacity = static_cast(boundaries.size()); framework_bind_runtime(reinterpret_cast(&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 boundaries; FakeRuntime fake{}; fake.ops = &kFakeOps; + fake.boundaries = boundaries.data(); + fake.boundary_capacity = static_cast(boundaries.size()); framework_bind_runtime(reinterpret_cast(&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(GraphTensorSourceKind::INTERNAL); tensor_sources[1].packed_offset = 16; std::vector scalars{0, 18}; - std::vector scalar_sources(2); - scalar_sources[0].source_kind = static_cast(GraphScalarSourceKind::BOUNDARY); - scalar_sources[0].source_index = boundary_scalar_count - 1; - scalar_sources[1].source_kind = static_cast(GraphScalarSourceKind::STATIC_VALUE); + std::vector 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 *>, + "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(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(), 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{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(&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(&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{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{17}); + EXPECT_EQ(task_args.scalar(1).to(), 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{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{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{7}); } TEST(GraphExecutionStorage, ComputesAlignedExactSize) {