From fc2d65c1a376deb4deb1a579fa81d292732e5762 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 24 Apr 2026 11:24:03 +0200 Subject: [PATCH] =?UTF-8?q?refactor(dqn-v2):=20Invariant=208=20=E2=80=94?= =?UTF-8?q?=20named=20constants=20for=20branch=20indices?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Add BRANCH_DIR/BRANCH_MAG/BRANCH_ORD/BRANCH_URG/NUM_BRANCHES to state_layout.cuh. Migrate the 4 clear literal branch-index accesses in branch_grad_balance_kernel.cu (branch_norms_dev[0..3] in grad_balance_isv_update). Other branch-keyed arrays (branch_starts, branch_lens, branch_norms, branch_scales) use the runtime `branch = blockIdx.y` variable or loop counter — no raw literal index, so no migration needed. Rust call sites use struct fields (branch_0_size etc.) or loop vars — not literals. No behavioural change; pure refactor. Plan 1 Task 4D. Spec §3 Invariant 8. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/branch_grad_balance_kernel.cu | 10 ++++++---- crates/ml/src/cuda_pipeline/state_layout.cuh | 11 +++++++++++ docs/dqn-named-dims.md | 3 ++- 3 files changed, 19 insertions(+), 5 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/branch_grad_balance_kernel.cu b/crates/ml/src/cuda_pipeline/branch_grad_balance_kernel.cu index 3d42424da..ee9b0456f 100644 --- a/crates/ml/src/cuda_pipeline/branch_grad_balance_kernel.cu +++ b/crates/ml/src/cuda_pipeline/branch_grad_balance_kernel.cu @@ -63,6 +63,8 @@ * between the post_aux phase and `compute_grad_norm_for_adam`. */ +#include "state_layout.cuh" + extern "C" __global__ void branch_grad_norm_reduce( const float* __restrict__ grad_buf, const int* __restrict__ branch_starts, /* [4] element offsets into grad_buf */ @@ -147,10 +149,10 @@ extern "C" __global__ void grad_balance_isv_update( ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; - float n0 = branch_norms_dev[0]; - float n1 = branch_norms_dev[1]; - float n2 = branch_norms_dev[2]; - float n3 = branch_norms_dev[3]; + float n0 = branch_norms_dev[BRANCH_DIR]; + float n1 = branch_norms_dev[BRANCH_MAG]; + float n2 = branch_norms_dev[BRANCH_ORD]; + float n3 = branch_norms_dev[BRANCH_URG]; /* All-zero branch norms → quiescent training step; skip the update * to avoid driving targets toward zero and the limit to 1.0 (which diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index ebeb3fb3f..da1fb8726 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -94,6 +94,17 @@ #define PLAN_PARAM_ASYMMETRY 5 // profit/stop ratio #define PLAN_PARAM_DIM 6 +// ──────────────────────────────────────────────────────────────────────────── +// Branch selector indices (4-branch factored action). +// Invariant 8: every branch has a named constant. +// These are used to index per-branch arrays (norms, targets, sizes, etc.). +// ──────────────────────────────────────────────────────────────────────────── +#define BRANCH_DIR 0 // direction branch (4 actions: Short/Hold/Long/Flat) +#define BRANCH_MAG 1 // magnitude branch (3 actions: Quarter/Half/Full) +#define BRANCH_ORD 2 // order-type branch (3 actions) +#define BRANCH_URG 3 // urgency branch (3 actions) +#define NUM_BRANCHES 4 + // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); diff --git a/docs/dqn-named-dims.md b/docs/dqn-named-dims.md index 808837d94..b7f5d86a3 100644 --- a/docs/dqn-named-dims.md +++ b/docs/dqn-named-dims.md @@ -120,4 +120,5 @@ From the trade_plan MLP output. - Task 4A (ps[0..PS_STRIDE) constants): commit 144c85b85 - Task 4B (plan_isv[0..6) constants): commit 741cb48d5 -- Task 4C (plan_params[0..6) constants): commit +- Task 4C (plan_params[0..6) constants): commit 0ac83479e +- Task 4D (BRANCH_* constants): commit