kernbench2/docs/adr-ko/ADR-0026-par-dppolicy-intra-device.md

# ADR-0026: DPPolicy = Intra-Device Only — sip/num_sips 필드 제거

## Status

Accepted (Revision 6 — cube_start 추가 2026-06-04; Revision 5 landed 2026-04-14)

## Context

### 목표

`DPPolicy`를 **한 device(SIP) 내부의 cube × PE 분산**만 표현하는 순수한
intra-device 추상화로 명확화한다. SIP 간 분산(TP)은 별도 레이어로 분리
(ADR-0024의 `torch.ahbm.set_device(rank)` 또는 ADR-0027의 Megatron parallel
layers가 담당).

## Decision

### D1. `DPPolicy`에서 `sip` + `num_sips` 필드 제거

```python
@dataclass(frozen=True)
class DPPolicy:
    """Intra-device (cube × PE) data-parallel policy.

    SIP-level placement is controlled by ``torch.ahbm.set_device(rank)``
    (ADR-0024 D3) and, for model-level TP, by Megatron-style parallel
    layers (ADR-0027). DPPolicy does not cross SIP boundaries.
    """
    cube: Literal["replicate", "column_wise", "row_wise"] = "replicate"
    pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
    num_pes: int | None = None
    num_cubes: int | None = None
    cube_start: int = 0   # Revision 6 — SIP 내 첫 cube 인덱스; D8 참조
```

제거되는 필드: `sip`, `num_sips`.
추가되는 필드 (Revision 6): `cube_start` — D8 참조.

### D2. `ShardSpec` — structural (sip, cube, pe) 좌표, `pe_index` 완전 제거

현재 `ShardSpec.pe_index`는 **global flat index** (`sip × cubes × pes + cube ×
pes + pe`). 이는 ADR-0024 D4이 "abstraction leakage"로 지적한 형태.

본 ADR에서 ShardSpec을 **structural 좌표로 재정의**하고, `pe_index`는
property로도 **남기지 않는다**:

```python
# src/kernbench/policy/placement/dp.py (after)
@dataclass(frozen=True)
class ShardSpec:
    """Structural shard placement — intra-SIP (cube × PE) coord.

    Global-flat `pe_index` was removed in ADR-0026. Callers must use
    structural coords (sip, cube, pe) directly. If a flat integer key is
    needed (e.g. dict lookup), compute it explicitly at the call site.
    """
    sip: int              # structural — which SIP this shard lives on
    cube: int             # local within SIP
    pe: int               # local within cube
    offset_bytes: int
    nbytes: int
```

**핵심 원칙**:
- ShardSpec의 정체성은 `(sip, cube, pe)` 3튜플.
- **`pe_index` property도 없음** — silent semantics drift 차단.
- Global flat을 기대한 기존 호출자는 `.pe_index` 접근 시 **즉시
  `AttributeError`** → 반드시 구조적 좌표로 migration.
- Flat integer key가 필요한 국소 문맥 (예: 내부 dict lookup)은 호출자가
  명시적으로 `spec.sip * N_CUBES * N_PE + spec.cube * N_PE + spec.pe`를 계산.

**Property 제거 정당화**: KernBench는 사내 프로젝트로 call site가 한정되어
있음. Silent drift 위험 (의미만 바뀌고 타입은 같은 int) 대비 explicit breakage
(AttributeError)가 훨씬 안전.

### D3. `resolve_dp_policy`가 `target_sip`을 받아 structural 좌표 생성

ADR-0024 D4의 계약 구현. Post-hoc shifting 없음.

```python
# src/kernbench/policy/placement/dp.py (after)

@dataclass(frozen=True)
class _LocalPeShard:
    """Internal — PE resolver의 반환. Cube 내 local PE 식별자 + payload."""
    local_pe: int                  # cube-local PE index (0..num_pe-1)
    offset_bytes: int
    nbytes: int


def resolve_dp_policy(
    policy: DPPolicy,
    *,
    shape: tuple[int, int],
    itemsize: int,
    num_pe: int,
    num_cubes: int = 1,
    target_sip: int,       # NEW — 어느 SIP에 배치할지 명시
) -> list[ShardSpec]:
    """2-level resolution (cube × PE) on a specified SIP.

    Returns ShardSpecs with structural coords (sip=target_sip, cube, pe).
    No SIP-level split — DPPolicy is intra-device only.
    """
    resolver = _PE_RESOLVERS[policy.pe]
    all_shards: list[ShardSpec] = []

    # Level 1: cube within SIP
    cube_splits = _split_shape(policy.cube, shape, num_cubes, itemsize)

    for cube_id, (cube_shape, cube_offset) in enumerate(cube_splits):
        # Level 2: PE within cube — resolver returns _LocalPeShard (local_pe)
        local_shards = resolver(shape=cube_shape, itemsize=itemsize,
                                 num_pe=num_pe)

        for ls in local_shards:
            all_shards.append(ShardSpec(
                sip=target_sip,                              # from caller (current_device)
                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)
                offset_bytes=cube_offset + ls.offset_bytes,
                nbytes=ls.nbytes,
            ))

    return all_shards
```

**내부 resolver** (`column_wise`, `row_wise`, `replicate`)는 `_LocalPeShard`
리스트 반환 — `local_pe` 필드명으로 **"cube-local PE identifier"임이 명시적**.
과거 `ShardSpec.pe_index`와 이름이 혼동되던 문제 해소.

**이름 규약 정리** (전체 ADR):
- `ShardSpec.pe`: 최종 외부 API — cube-local PE (structural coord)
- `_LocalPeShard.local_pe`: 내부 resolver 단계의 동일 의미
- `pe_index`: **제거**. 외부/내부 어디에도 남기지 않는다 (silent drift 차단의
  부가 효과: 이름 재등장 없음).

### D4. `_create_tensor` — 구조적 좌표로 직접 placement

ADR-0024 D4 연속선. Post-hoc shifting 제거, 구조적 좌표를 `resolve_dp_policy`
호출 시점에 직접 지정.

```python
# context.py _create_tensor (after)
current_sip = self.ahbm.current_device()
if current_sip is None:
    # Single-driver fallback (ADR-0024 D2와 일관).
    # Launcher 기반 코드가 set_device()를 빼먹으면 조용히 SIP 0에 박히는
    # 문제가 있음 → debug mode에서 경고.
    if os.environ.get("KERNBENCH_DEBUG"):
        import warnings
        warnings.warn(
            "torch.ahbm.current_device() is None; defaulting to SIP 0. "
            "If this is a multi-rank launcher context, you likely forgot "
            "torch.ahbm.set_device(rank) inside the worker.",
            stacklevel=2,
        )
    current_sip = 0

placement = resolve_dp_policy(
    dp,
    shape=shape_2d,
    itemsize=itemsize,
    num_pe=eff_num_pe,
    num_cubes=eff_num_cubes,
    target_sip=current_sip,          # ← 구조적 좌표 일차 지정
)

# placement의 각 ShardSpec은 이미 (sip=current_sip, cube=local, pe=local) 포함.
# 과거의 post-hoc shifting 블록은 완전히 제거.
```

**모든** 텐서가 current device SIP에 배치됨. Multi-SIP 텐서를 만들고 싶으면
ADR-0027의 TP primitive 사용.

**Single-driver fallback의 trade-off**: set_device 없는 호출에서 SIP 0으로
default는 기존 single-driver 테스트 호환을 위해 유지. `KERNBENCH_DEBUG=1`
환경에서는 launcher 컨텍스트의 실수로 set_device 누락 시 조용히 잘못된 SIP에
배치되는 것을 감지할 수 있도록 warning.

### D5. Downstream — allocator lookup은 구조적 tuple key로

기존 `deploy_tensor` (`src/kernbench/runtime_api/tensor.py`):

```python
for spec in placement:
    alloc = allocators[spec.pe_index]       # ← AttributeError (property 제거됨)
```

`pe_index`가 없어졌으므로 구조적 좌표로 **강제** migration:

```python
for spec in placement:
    alloc = allocators[(spec.sip, spec.cube, spec.pe)]
```

`_ensure_allocators`의 dict population도 tuple key로:

```python
# context.py _ensure_allocators (after)
for sip_id in sip_range:
    for cube_id in range(cubes_per_sip):
        for pe_id in range(pes_per_cube):
            self._allocators[(sip_id, cube_id, pe_id)] = PEMemAllocator(
                rack_id=0, sip_id=sip_id, cube_id=cube_id, pe_id=pe_id, cfg=cfg,
            )
```

`_free_tensor`도 동일: 기존 `flat_idx = sip * ... + cube * ... + pe` 계산
블록 제거, `(shard.sip, shard.cube, shard.pe)` 직접 사용.

**Tuple vs dataclass `PEIdentity`**: Tuple이 단순하고 hashable로 바로 써서
권고. `PEIdentity` 값객체는 명시적 타입 장점은 있지만 boilerplate가 크고 현재
allocator dict의 유일한 key라 오버엔지니어링. Tuple 유지.

### D7. 하위 호환 — 불가 (cleanup ADR)

이 ADR은 **breaking change**.

1. `DPPolicy(sip=...)` 또는 `DPPolicy(num_sips=...)` 호출 → `TypeError`
2. `ShardSpec.pe_index` 접근 → `AttributeError`

모두 **즉시 명시적 breakage**. Deprecation warning / fallback 경로 없음.
KernBench는 사내 프로젝트로 call site가 한정되어 있어 한 번에 migration.

**Silent drift 차단**이 property 완전 제거의 주된 이점: global flat을 기대한
코드가 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

- **ADR-0024** (launcher): `set_device(rank)` 및 current-device scoping이
  SIP 배치 메커니즘 제공. 본 ADR은 그 위에 서서 DPPolicy를 순수 intra-device로
  좁힘.
- **ADR-0027** (Megatron TP): 다중 SIP에 걸친 텐서가 필요한 경우의 대안 경로.
  이 ADR 적용 후 multi-SIP use case는 ADR-0027로 이관.

---

## Non-goals

- **`DPPolicy.cube` / `pe` 재설계**: 기존 replicate/column_wise/row_wise 의미
  유지.
- **Tiling 정책 통합**: `tiled_column_major` / `tiled_row_major`는 그대로.
- **Multi-device 텐서 추상화 신규**: DTensor-like는 ADR-0028.

---

## Open questions

- **`_create_tensor`의 current_sip 기본값**: set_device 없는 호출에서 rank=0
  (SIP 0)로 fallback할지, 아니면 error 낼지. 권고는 fallback (기존 single-driver
  테스트와의 호환).
- **`test_sip_parallel.py` 재작성 범위**: 기존 단위 테스트의 의도를 유지하며
  launcher 기반으로 옮기려면 추가 fixture 필요. 별도 작업으로 scope.
- **`DPPolicy`의 `num_sips=None` 의미**: 필드가 없어지면 `num_sips` 개념 자체가
  사라짐. Multi-SIP을 표현하고 싶으면 ADR-0027의 TP primitive를 쓰라는 것이
  명시적 답.

**Resolved (이전 rev에서 open이었던 것들)**:
- ~~`ShardSpec.pe_index` property 존치 여부~~ → **완전 제거** (D2)
- ~~`_ensure_allocators` dict key 형식~~ → **tuple `(sip, cube, pe)`** (D5)

---

## Consequences

### Positive

- **개념 분리 명확**: DPPolicy = intra-device, TP = inter-device.
- **API 단순화**: DPPolicy 생성자 필드 ~33% 축소.
- **Structural 좌표 일관성**: ShardSpec이 `(sip, cube, pe)` 튜플로 표현 →
  abstraction leakage 해소 (ADR-0024 D4 계약 충족).
- **`pe_index` 의미 명확**: SIP-local이 단일 해석. Global flat이 필요하면 명시.
- **Launcher 모델 일관성**: ADR-0024의 "1 worker per SIP" 모델이 유일한 SIP
  경계 제어 메커니즘.

### Negative

- **Breaking change (explicit)**: `DPPolicy(sip=...)` → `TypeError`,
  `spec.pe_index` → `AttributeError`. 모든 호출자 한 번에 수정 필요.
- **ShardSpec schema 변경**: `pe_index` 단일 필드 → `sip`/`cube`/`pe` 세 필드.
  Downstream (`deploy_tensor`, `_free_tensor`, `_ensure_allocators`,
  `allocators` dict key 등) 연쇄 수정.
- **Silent drift 없음**: property 완전 제거로 runtime에서 즉시 실패 →
  migration leakage 원천 차단. (Negative가 아니라 explicit tradeoff)
- `test_sip_parallel.py` 재작성 비용.

### Neutral

- 기존 `cube` / `pe` 필드 의미 불변.