diff --git a/crates/ml-alpha/src/lib.rs b/crates/ml-alpha/src/lib.rs
index 03adf6afe..8a22b070a 100644
--- a/crates/ml-alpha/src/lib.rs
+++ b/crates/ml-alpha/src/lib.rs
@@ -35,6 +35,7 @@ pub mod pinned;
 pub mod pinned_mem;
 #[cfg(feature = "kernel-step-trace")]
 pub mod gpu_log;
+pub mod rl;
 pub mod trainer;
 
 // Preserved from prior crate state — used by Phase A.
diff --git a/crates/ml-alpha/src/rl/common.rs b/crates/ml-alpha/src/rl/common.rs
new file mode 100644
index 000000000..4f56155a9
--- /dev/null
+++ b/crates/ml-alpha/src/rl/common.rs
@@ -0,0 +1,62 @@
+//! Shared types for the integrated RL trainer.
+
+use cudarc::driver::CudaSlice;
+
+/// Discrete action grid cardinality. Matches the legacy DQN's 9-action
+/// layout (`crates/ml/src/cuda_pipeline/sp5_isv_slots.rs`) so the two
+/// systems stay comparable. See `action.rs` (Phase C/D) for the
+/// enum-style mapping {Short-large, Short-small, Hold, Flat-from-long,
+/// Flat-from-short, Long-small, Long-large, Trail-tighten, Trail-loosen}.
+pub const N_ACTIONS: usize = 9;
+
+/// C51 distributional Q-head atom count. Matches existing
+/// `pearl_per_branch_c51_atom_span` infrastructure (21 atoms over the
+/// learned support `[v_min, v_max]`).
+pub const Q_N_ATOMS: usize = 21;
+
+/// PPO update epochs per rollout. Algorithmic constant (no signal exists
+/// to drive it adaptively); standard PPO uses 3-10, we default to 4.
+pub const PPO_N_EPOCHS: usize = 4;
+
+/// Single `(state, action, reward, next_state)` transition.
+///
+/// Stored simultaneously in:
+/// - The off-policy replay buffer (DQN reads it for Bellman TD via PER).
+/// - The on-policy rollout buffer (PPO reads it for clipped surrogate).
+///
+/// `h_t` is the encoder representation, stored directly so the loss
+/// computation doesn't re-run the encoder. `q_value_old` and `log_pi_old`
+/// are recorded at sample time to bootstrap importance ratios (PPO) and
+/// priority updates (PER).
+///
+/// Phase A defines the struct shape only. Phase B wires it into the
+/// encoder forward path; Phase C/D into the loss heads; Phase E into the
+/// training loop and lobsim integration.
+#[derive(Debug)]
+pub struct Transition {
+    /// Encoder representation at step t. Shape `[B, HIDDEN_DIM=128]`,
+    /// device buffer (no host copy).
+    pub h_t: CudaSlice<f32>,
+
+    /// Action taken at step t. Range `0..N_ACTIONS`.
+    pub action: u32,
+
+    /// Per-trade realized PnL on trade close (USD). Zero between closes
+    /// (reward is sparse, fires only at position-flat events).
+    pub reward: f32,
+
+    /// Encoder representation at step t+1.
+    pub next_h_t: CudaSlice<f32>,
+
+    /// True at trade close (position returned to 0) or end-of-stream.
+    pub done: bool,
+
+    /// log π_old(action | state) recorded at sample time. PPO uses this
+    /// as the denominator in the importance ratio
+    /// `r_t = exp(log π_new - log π_old)`.
+    pub log_pi_old: f32,
+
+    /// Q values at sample time, all actions. Used to seed PER priority
+    /// (initial priority = |Q(a) - bootstrap_estimate|).
+    pub q_value_old: [f32; N_ACTIONS],
+}
diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs
new file mode 100644
index 000000000..a5e9ecf29
--- /dev/null
+++ b/crates/ml-alpha/src/rl/isv_slots.rs
@@ -0,0 +1,70 @@
+//! ISV slot allocations for the integrated RL trainer.
+//!
+//! Each constant names a slot index in the shared GPU ISV bus
+//! (allocated 400-411 to live above the existing SP5/SP7/SP11
+//! blocks which end at 383). Per `pearl_controller_anchors_isv_driven`
+//! and `feedback_isv_for_adaptive_bounds`, every adaptive hyperparameter
+//! is sourced from one of these slots, with the value emitted by a
+//! corresponding controller kernel.
+//!
+//! Phase A defines the constants. Phase C onward wires consumers (loss
+//! kernels read these via `__ldg(isv + INDEX)`). The controller kernels
+//! themselves (producers) land in their respective phases (C/D/E/F).
+//!
+//! Bootstrap discipline: every slot follows
+//! `pearl_first_observation_bootstrap`. The sentinel value `0.0` at the
+//! slot's address means "uninitialised — emit the bootstrap value on the
+//! next controller fire." First-observation values for each slot are
+//! documented in the plan §"ISV-driven adaptive control" table.
+
+/// Discount factor γ. Controller input: mean trade duration / anchor.
+/// Bootstrap 0.99.
+pub const RL_GAMMA_INDEX: usize = 400;
+
+/// Target-net soft-update fraction τ. Controller input: Q-divergence EMA.
+/// Bootstrap 0.005.
+pub const RL_TARGET_TAU_INDEX: usize = 401;
+
+/// PPO clipped surrogate ε. Controller input: KL(π_new‖π_old) EMA.
+/// Bootstrap 0.2.
+pub const RL_PPO_CLIP_INDEX: usize = 402;
+
+/// Entropy bonus weight. Controller input: action-entropy deficit.
+/// Bootstrap 0.01.
+pub const RL_ENTROPY_COEF_INDEX: usize = 403;
+
+/// Rollout buffer flush size. Controller input: advantage-variance ratio.
+/// Bootstrap 2048.
+pub const RL_N_ROLLOUT_STEPS_INDEX: usize = 404;
+
+/// PER priority exponent α. Controller input: TD-error kurtosis EMA.
+/// Bootstrap 0.6.
+pub const RL_PER_ALPHA_INDEX: usize = 405;
+
+/// Per-trade reward standardization scale. Controller input: mean
+/// |realized_pnl_usd| over last 100 closes. Bootstrap 1.0.
+pub const RL_REWARD_SCALE_INDEX: usize = 406;
+
+/// Diagnostic: agreement rate between argmax(Q) and mode(π). Not
+/// consumed by any loss; alarm fires if value < 0.5. Bootstrap 0.0
+/// (EMA-tracked from first observation, no init magic).
+pub const RL_Q_ARG_VS_PI_AGREE_INDEX: usize = 407;
+
+/// Loss-balance λ for the BCE direction head. Produced by the existing
+/// `pearl_loss_balance_controller` Wiener-α controller (extended from 3
+/// horizons to 5 heads). Bootstrap 1.0 (equal-weight 1/5 init applied
+/// by the controller's first emit).
+pub const RL_LOSS_LAMBDA_BCE_INDEX: usize = 408;
+
+/// Loss-balance λ for the C51 distributional Q-head.
+pub const RL_LOSS_LAMBDA_Q_INDEX: usize = 409;
+
+/// Loss-balance λ for the PPO policy head.
+pub const RL_LOSS_LAMBDA_PI_INDEX: usize = 410;
+
+/// Loss-balance λ for the scalar value (V) head.
+pub const RL_LOSS_LAMBDA_V_INDEX: usize = 411;
+
+/// Last RL-allocated slot index (exclusive). The integrated trainer
+/// extends `ISV_TOTAL_DIM` to at least this value at trainer init time.
+pub const RL_SLOTS_END: usize = 412;
diff --git a/crates/ml-alpha/src/rl/mod.rs b/crates/ml-alpha/src/rl/mod.rs
new file mode 100644
index 000000000..300d17130
--- /dev/null
+++ b/crates/ml-alpha/src/rl/mod.rs
@@ -0,0 +1,19 @@
+//! Integrated RL trainer for the alpha-perception model.
+//!
+//! Per the plan in `docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md`:
+//! single trainer combining BCE direction + C51 distributional Q + categorical
+//! π + scalar V + aux prof+size losses on a shared Mamba2+CfC encoder. Joint
+//! training with adaptive loss-balance weights; per-trade realized PnL reward
+//! from `ml_backtesting::sim::LobSimCuda`; discrete 9-action grid.
+//!
+//! All hyperparameters (γ, target-net τ, PPO clip ε, entropy coef, rollout
+//! length, PER α, reward scale, loss-balance λs) are sourced from ISV slots
+//! per `pearl_controller_anchors_isv_driven` and
+//! `feedback_isv_for_adaptive_bounds`. ISV slot definitions live in `isv_slots`.
+//!
+//! Phase A (this commit) lands the module structure + `Transition` type +
+//! ISV slot constants only. Subsequent phases (B-H) wire the heads, controllers,
+//! training loop, Argo plumbing, and backtest gate.
+
+pub mod common;
+pub mod isv_slots;