policy(adr): ADR-0026 Revision 6 — document DPPolicy.cube_start
Amends ADR-0026 to reflect the cube_start field added in e2fe331.
The production code shipped without an ADR update; this fills that
gap. Documentation-only change (no production code, no test code).
Changes (mirrored in both EN and KO):
- Status: Revision 5 → Revision 6
- D1: add ``cube_start: int = 0`` to the canonical DPPolicy dataclass
- D3: ``cube = policy.cube_start + cube_id`` in resolve_dp_policy
- D8: new section explaining purpose (disjoint cube sub-meshes for
GQA Llama-70B 8-KV-group headline), semantics, default-0 backward
compatibility, intra-device constraint, design rationale
(scalar vs 2D origin vs cube_ids list), and the kernel-side
cube_start subtraction needed to compensate for ADR-0022's
physical-cube-id ``program_id(axis=1)`` semantics.
tools/verify_adr_lang_pairs.py passes (EN/KO Status keyword and
title in sync).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
@@ -2,7 +2,7 @@
|
|||||||
|
|
||||||
## Status
|
## Status
|
||||||
|
|
||||||
Accepted (Revision 5 — Phase 2 landed 2026-04-14, 523 passed + 1 strict xfail)
|
Accepted (Revision 6 — cube_start 추가 2026-06-04; Revision 5 landed 2026-04-14)
|
||||||
|
|
||||||
## Context
|
## Context
|
||||||
|
|
||||||
@@ -30,9 +30,11 @@ class DPPolicy:
|
|||||||
pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
|
pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
|
||||||
num_pes: int | None = None
|
num_pes: int | None = None
|
||||||
num_cubes: int | None = None
|
num_cubes: int | None = None
|
||||||
|
cube_start: int = 0 # Revision 6 — SIP 내 첫 cube 인덱스; D8 참조
|
||||||
```
|
```
|
||||||
|
|
||||||
제거되는 필드: `sip`, `num_sips`.
|
제거되는 필드: `sip`, `num_sips`.
|
||||||
|
추가되는 필드 (Revision 6): `cube_start` — D8 참조.
|
||||||
|
|
||||||
### D2. `ShardSpec` — structural (sip, cube, pe) 좌표, `pe_index` 완전 제거
|
### D2. `ShardSpec` — structural (sip, cube, pe) 좌표, `pe_index` 완전 제거
|
||||||
|
|
||||||
@@ -114,7 +116,7 @@ def resolve_dp_policy(
|
|||||||
for ls in local_shards:
|
for ls in local_shards:
|
||||||
all_shards.append(ShardSpec(
|
all_shards.append(ShardSpec(
|
||||||
sip=target_sip, # from caller (current_device)
|
sip=target_sip, # from caller (current_device)
|
||||||
cube=cube_id, # local within SIP
|
cube=policy.cube_start + cube_id, # Rev 6: D8 cube_start만큼 shift; 기본 0 = local within SIP
|
||||||
pe=ls.local_pe, # local within cube (explicit name)
|
pe=ls.local_pe, # local within cube (explicit name)
|
||||||
offset_bytes=cube_offset + ls.offset_bytes,
|
offset_bytes=cube_offset + ls.offset_bytes,
|
||||||
nbytes=ls.nbytes,
|
nbytes=ls.nbytes,
|
||||||
@@ -224,6 +226,47 @@ KernBench는 사내 프로젝트로 call site가 한정되어 있어 한 번에
|
|||||||
**Silent drift 차단**이 property 완전 제거의 주된 이점: global flat을 기대한
|
**Silent drift 차단**이 property 완전 제거의 주된 이점: global flat을 기대한
|
||||||
코드가 SIP-local 결과를 받아 조용히 잘못된 인덱싱을 할 가능성 제거.
|
코드가 SIP-local 결과를 받아 조용히 잘못된 인덱싱을 할 가능성 제거.
|
||||||
|
|
||||||
|
### D8. SIP 내 분리된 cube sub-mesh를 위한 `cube_start: int = 0` 추가
|
||||||
|
|
||||||
|
**Revision 6 추가 사항 (2026-06-04)**.
|
||||||
|
|
||||||
|
**목적**: 한 SIP 내에서 분리된 cube sub-mesh 지정 (예: 0..7과 8..15을 함께
|
||||||
|
사용). GQA Llama-70B 8-KV-group 헤드라인 타겟 — SIP당 2×4 KV-group 2개 ×
|
||||||
|
4 SIP = 64 cubes — 의 두 번째 KV-group을 기본값 0..7이 아닌 cubes 8..15에
|
||||||
|
배치하기 위해 필요.
|
||||||
|
|
||||||
|
**Semantics**: `resolve_dp_policy`가
|
||||||
|
`cube = policy.cube_start + cube_id`로 ShardSpec 생성 (cube_id는 launch 내에서
|
||||||
|
`0..num_cubes-1` 순회). 선택된 cube는 target SIP 내의
|
||||||
|
`[cube_start, cube_start + num_cubes)` 범위.
|
||||||
|
|
||||||
|
**기본값**: `cube_start = 0`은 기존 모든 호출 사이트 동작을 bit-for-bit
|
||||||
|
보존 (이전처럼 `ShardSpec.cube ∈ [0, num_cubes)`). CCL milestone bench의
|
||||||
|
full-SIP `DPPolicy(num_cubes=16)`은 그대로 cubes 0..15 생성.
|
||||||
|
|
||||||
|
**제약**: intra-device 불변성 보존.
|
||||||
|
`cube_start ∈ [0, cubes_per_sip)`이며
|
||||||
|
`cube_start + num_cubes ≤ cubes_per_sip`. SIP 경계 교차는 여전히
|
||||||
|
`ahbm.set_device(rank)` (ADR-0024)의 책임.
|
||||||
|
|
||||||
|
**왜 scalar인가** (2D `cube_mesh_origin`이나 임의 `cube_ids` list가 아닌):
|
||||||
|
consumer 커널 (예: `_attention_mesh_mlo_2d`)이 row-major contiguous cube를
|
||||||
|
가정; scalar `cube_start`는 `num_cubes`와 자연스럽게 쌍을 이룸 (범위 =
|
||||||
|
`[start, start + count)`); 기본값 0에서 마이그레이션 부담 없음. 비연속
|
||||||
|
케이스가 필요해지면 향후 더 일반적인 설계를 위에 얹을 수 있음.
|
||||||
|
|
||||||
|
**하위 호환성**: 기본값이 있는 가산적 변경. 기존 호출 사이트의 변경을
|
||||||
|
강제하지 않음. D7의 "breaking change" 입장은 원래 sip/num_sips 제거에만
|
||||||
|
해당하며, `cube_start`는 기존 호출자를 깨지 않음.
|
||||||
|
|
||||||
|
**커널 측 참고**: kernbench의 `tl.program_id(axis=1)`은 launch-local rank가
|
||||||
|
아닌 물리 cube id를 반환 (ADR-0022). `program_id(axis=1)`에서 ring 위치를
|
||||||
|
도출하는 커널은 `cube_start > 0`일 때 launch-local rank를 복원하기 위해
|
||||||
|
`cube_start`를 빼야 함 — 그렇지 않으면 out-of-bounds sub-mesh 위치를
|
||||||
|
계산하여 deadlock 발생. 참고 패턴은 `_attention_mesh_mlo_2d.py` 참조.
|
||||||
|
ADR-0022의 향후 개정에서 `program_id(axis=1)`을 launch-local rank semantics로
|
||||||
|
이동하면 커널 측 명시적 차감이 제거될 수 있음.
|
||||||
|
|
||||||
## Dependencies
|
## Dependencies
|
||||||
|
|
||||||
- **ADR-0024** (launcher): `set_device(rank)` 및 current-device scoping이
|
- **ADR-0024** (launcher): `set_device(rank)` 및 current-device scoping이
|
||||||
|
|||||||
@@ -2,7 +2,7 @@
|
|||||||
|
|
||||||
## Status
|
## Status
|
||||||
|
|
||||||
Accepted (Revision 5 — Phase 2 landed 2026-04-14, 523 passed + 1 strict xfail)
|
Accepted (Revision 6 — cube_start added 2026-06-04; Revision 5 landed 2026-04-14)
|
||||||
|
|
||||||
## Context
|
## Context
|
||||||
|
|
||||||
@@ -31,9 +31,11 @@ class DPPolicy:
|
|||||||
pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
|
pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
|
||||||
num_pes: int | None = None
|
num_pes: int | None = None
|
||||||
num_cubes: int | None = None
|
num_cubes: int | None = None
|
||||||
|
cube_start: int = 0 # Revision 6 — first cube in SIP; see D8
|
||||||
```
|
```
|
||||||
|
|
||||||
Removed fields: `sip`, `num_sips`.
|
Removed fields: `sip`, `num_sips`.
|
||||||
|
Added fields (Revision 6): `cube_start` — see D8.
|
||||||
|
|
||||||
### D2. `ShardSpec` — structural (sip, cube, pe) coordinates, `pe_index` fully removed
|
### D2. `ShardSpec` — structural (sip, cube, pe) coordinates, `pe_index` fully removed
|
||||||
|
|
||||||
@@ -120,7 +122,7 @@ def resolve_dp_policy(
|
|||||||
for ls in local_shards:
|
for ls in local_shards:
|
||||||
all_shards.append(ShardSpec(
|
all_shards.append(ShardSpec(
|
||||||
sip=target_sip, # from caller (current_device)
|
sip=target_sip, # from caller (current_device)
|
||||||
cube=cube_id, # local within SIP
|
cube=policy.cube_start + cube_id, # Rev 6: shifted by D8 cube_start; default 0 = local within SIP
|
||||||
pe=ls.local_pe, # local within cube (explicit name)
|
pe=ls.local_pe, # local within cube (explicit name)
|
||||||
offset_bytes=cube_offset + ls.offset_bytes,
|
offset_bytes=cube_offset + ls.offset_bytes,
|
||||||
nbytes=ls.nbytes,
|
nbytes=ls.nbytes,
|
||||||
@@ -239,6 +241,53 @@ property: code that expected a global flat could otherwise silently
|
|||||||
receive a SIP-local result and index incorrectly — that possibility is
|
receive a SIP-local result and index incorrectly — that possibility is
|
||||||
eliminated.
|
eliminated.
|
||||||
|
|
||||||
|
### D8. Add `cube_start: int = 0` for disjoint cube sub-meshes within a SIP
|
||||||
|
|
||||||
|
**Revision 6 addition (2026-06-04)**.
|
||||||
|
|
||||||
|
**Purpose**: address a disjoint cube sub-mesh within one SIP (e.g.
|
||||||
|
cubes 8..15 alongside cubes 0..7). Required by the GQA Llama-70B
|
||||||
|
8-KV-group headline target — two 2×4 KV-groups per SIP × 4 SIPs = 64
|
||||||
|
cubes — where the second KV-group on each SIP must land on cubes 8..15
|
||||||
|
instead of the default 0..7.
|
||||||
|
|
||||||
|
**Semantics**: `resolve_dp_policy` returns `ShardSpec` with
|
||||||
|
`cube = policy.cube_start + cube_id`, where `cube_id` iterates
|
||||||
|
`0..num_cubes-1` within the launch. Selected cubes lie in
|
||||||
|
`[cube_start, cube_start + num_cubes)` within the target SIP.
|
||||||
|
|
||||||
|
**Default**: `cube_start = 0` preserves every existing call site
|
||||||
|
bit-for-bit (`ShardSpec.cube ∈ [0, num_cubes)` as before). The CCL
|
||||||
|
milestone bench's full-SIP `DPPolicy(num_cubes=16)` continues to
|
||||||
|
produce cubes 0..15.
|
||||||
|
|
||||||
|
**Constraint**: intra-device invariant preserved.
|
||||||
|
`cube_start ∈ [0, cubes_per_sip)` and
|
||||||
|
`cube_start + num_cubes ≤ cubes_per_sip`. SIP boundary crossing remains
|
||||||
|
the job of `ahbm.set_device(rank)` (ADR-0024).
|
||||||
|
|
||||||
|
**Why scalar** (not 2D `cube_mesh_origin` or arbitrary `cube_ids` list):
|
||||||
|
consumer kernels (e.g. `_attention_mesh_mlo_2d`) assume row-major
|
||||||
|
contiguous cubes; scalar `cube_start` pairs naturally with `num_cubes`
|
||||||
|
(range = `[start, start + count)`); zero migration churn at default 0.
|
||||||
|
More general designs can be added on top later if a non-contiguous use
|
||||||
|
case appears.
|
||||||
|
|
||||||
|
**Backward compatibility**: additive change with default value. No
|
||||||
|
existing call site is forced to change. The "breaking change" stance
|
||||||
|
in D7 applies only to the original sip/num_sips removal — `cube_start`
|
||||||
|
does NOT break existing callers.
|
||||||
|
|
||||||
|
**Kernel-side note**: kernbench's `tl.program_id(axis=1)` returns the
|
||||||
|
physical cube id (ADR-0022), not a launch-local rank. Kernels that
|
||||||
|
derive ring positions from `program_id(axis=1)` must subtract
|
||||||
|
`cube_start` to recover launch-local rank when `cube_start > 0` —
|
||||||
|
otherwise they compute out-of-bounds sub-mesh positions and deadlock.
|
||||||
|
See `_attention_mesh_mlo_2d.py` for the reference pattern. A future
|
||||||
|
revision of ADR-0022 could move `program_id(axis=1)` to launch-local
|
||||||
|
rank semantics; that change would let kernels drop the explicit
|
||||||
|
subtraction.
|
||||||
|
|
||||||
## Dependencies
|
## Dependencies
|
||||||
|
|
||||||
- **ADR-0024** (launcher): `set_device(rank)` and current-device scoping
|
- **ADR-0024** (launcher): `set_device(rank)` and current-device scoping
|
||||||
|
|||||||
Reference in New Issue
Block a user