diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh
index 22547932f..e1fc406a1 100644
--- a/crates/ml/src/cuda_pipeline/trade_physics.cuh
+++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh
@@ -46,6 +46,40 @@ __device__ __forceinline__ float compute_target_position(
     return fraction * max_position;
 }
 
+/* ── 4-branch hierarchical position mapping (Phase 2) ────────────────
+ * direction(3) × magnitude(3) replaces flat exposure(9).
+ * Direction: 0=Short(-1), 1=Flat(0), 2=Long(+1)
+ * Magnitude: 0=Quarter(0.25), 1=Half(0.50), 2=Full(1.00)
+ * Position = direction × magnitude × max_position
+ *
+ * Grid: S25  S50  S100
+ *       0    0    0     (Flat × any = 0)
+ *       L25  L50  L100
+ */
+__device__ __forceinline__ float compute_target_position_4branch(
+    int dir_idx, int mag_idx, float max_position
+) {
+    float direction = (dir_idx == 0) ? -1.0f : (dir_idx == 2) ? 1.0f : 0.0f;
+    float magnitude = (mag_idx == 0) ? 0.25f : (mag_idx == 1) ? 0.5f : 1.0f;
+    return direction * magnitude * max_position;
+}
+
+/* ── 4-branch action decoding ────────────────────────────────────────
+ * Action = dir * (b1*b2*b3) + mag * (b2*b3) + order * b3 + urgency
+ * With all sizes=3: action = dir*27 + mag*9 + order*3 + urgency */
+__device__ __forceinline__ int decode_direction_4b(int action, int b1, int b2, int b3) {
+    return action / (b1 * b2 * b3);
+}
+__device__ __forceinline__ int decode_magnitude_4b(int action, int b1, int b2, int b3) {
+    return (action / (b2 * b3)) % b1;
+}
+__device__ __forceinline__ int decode_order_4b(int action, int b2, int b3) {
+    return (action / b3) % b2;
+}
+__device__ __forceinline__ int decode_urgency_4b(int action, int b3) {
+    return action % b3;
+}
+
 /* ── Margin-aware position cap with drawdown scaling ──────────────────── */
 /* Real brokers enforce margin requirements AND risk managers cut exposure
  * as drawdown increases. Without drawdown scaling, the margin cap is static