| from dataclasses import dataclass |
| from typing import List, Tuple, Dict, Callable, Optional |
|
|
| |
| |
| |
| @dataclass(frozen=True) |
| class FragmentElement: |
| operand: str |
| lane: int |
| vgpr: int |
| packed_half: Optional[int] |
| logical_row: int |
| logical_col: int |
|
|
| @dataclass(frozen=True) |
| class ReadOp: |
| lane: int |
| operand: str |
| address: int |
| width_bytes: int = 64 |
|
|
| |
| |
| |
| def mfma_16x16x16_f16_a_coords(lane: int) -> List[Tuple[int, int, int, int]]: |
| if not 0 <= lane < 64: |
| raise ValueError("lane must be in 0..63") |
| m = lane >> 2 |
| k0 = (lane & 0x3) << 2 |
| return [ |
| (m, k0 + 0, 0, 0), |
| (m, k0 + 1, 0, 1), |
| (m, k0 + 2, 1, 0), |
| (m, k0 + 3, 1, 1), |
| ] |
|
|
| def mfma_16x16x16_f16_b_coords(lane: int) -> List[Tuple[int, int, int, int]]: |
| if not 0 <= lane < 64: |
| raise ValueError("lane must be in 0..63") |
| k0 = (lane >> 4) << 2 |
| n = lane & 0xF |
| return [ |
| (k0 + 0, n, 0, 0), |
| (k0 + 1, n, 0, 1), |
| (k0 + 2, n, 1, 0), |
| (k0 + 3, n, 1, 1), |
| ] |
|
|
| def mfma_16x16x16_f16_cd_coords(lane: int) -> List[Tuple[int, int, int]]: |
| if not 0 <= lane < 64: |
| raise ValueError("lane must be in 0..63") |
| n = lane & 0xF |
| m0 = lane >> 4 |
| return [ |
| (m0 + 0, n, 0), |
| (m0 + 4, n, 1), |
| (m0 + 8, n, 2), |
| (m0 + 12, n, 3), |
| ] |
|
|
| def generate_v_mfma_f32_16x16x16f16_fragments() -> Dict[str, List[FragmentElement]]: |
| fragments: Dict[str, List[FragmentElement]] = {"A": [], "B": [], "C": [], "D": []} |
| for lane in range(64): |
| for m, k, reg, half in mfma_16x16x16_f16_a_coords(lane): |
| fragments["A"].append(FragmentElement( |
| operand="A", lane=lane, vgpr=reg, packed_half=half, |
| logical_row=m, logical_col=k |
| )) |
| for k, n, reg, half in mfma_16x16x16_f16_b_coords(lane): |
| fragments["B"].append(FragmentElement( |
| operand="B", lane=lane, vgpr=reg, packed_half=half, |
| logical_row=k, logical_col=n |
| )) |
| for m, n, reg in mfma_16x16x16_f16_cd_coords(lane): |
| fragments["C"].append(FragmentElement( |
| operand="C", lane=lane, vgpr=reg, packed_half=None, |
| logical_row=m, logical_col=n |
| )) |
| fragments["D"].append(FragmentElement( |
| operand="D", lane=lane, vgpr=reg, packed_half=None, |
| logical_row=m, logical_col=n |
| )) |
| return fragments |
|
|
| |
| |
| |
| def validate_fragment_map( |
| fragments: Dict[str, List[FragmentElement]], |
| m: int = 16, |
| n: int = 16, |
| k: int = 16, |
| ) -> None: |
| expected = { |
| "A": m * k, |
| "B": k * n, |
| "C": m * n, |
| "D": m * n, |
| } |
| for operand, count in expected.items(): |
| actual = len(fragments[operand]) |
| if actual != count: |
| raise ValueError( |
| f"{operand}: expected {count} logical elements, got {actual}" |
| ) |
| coords = { |
| (x.logical_row, x.logical_col) |
| for x in fragments[operand] |
| } |
| if len(coords) != count: |
| raise ValueError( |
| f"{operand}: logical-coordinate map is not bijective; " |
| f"{len(coords)} unique coordinates for {count} elements" |
| ) |
|
|
| |
| |
| |
| def build_read_plan_b64( |
| elements: List[FragmentElement], |
| operand: str, |
| opcode: str = "v_mfma_f32_16x16x16f16" |
| ) -> List[ReadOp]: |
| """ |
| Assumes each lane's four FP16 elements are to be loaded with one ds_read_b64. |
| The four elements must be stored in LDS as two consecutive 32-bit words: |
| word0: [elem0, elem1] at address A |
| word1: [elem2, elem3] at address A+4 |
| and the address A must be 4-byte aligned. |
| We compute the address per lane from the logical coordinates and a layout function |
| that will be provided later (here we just return a placeholder; the address will be |
| filled in by the layout function). |
| """ |
| |
| lane_to_elements: Dict[int, List[FragmentElement]] = {} |
| for elem in elements: |
| lane_to_elements.setdefault(elem.lane, []).append(elem) |
| |
| reads: List[ReadOp] = [] |
| for lane in range(64): |
| elems = lane_to_elements[lane] |
| if len(elems) != 4: |
| raise ValueError(f"Lane {lane} has {len(elems)} elements, expected 4") |
| |
| elems.sort(key=lambda e: (e.logical_row, e.logical_col)) |
| |
| reads.append(ReadOp( |
| lane=lane, |
| operand=operand, |
| address=0, |
| width_bytes=64 |
| )) |
| return reads |
|
|
| |
| |
| |
| def address_A( |
| lane: int, |
| row_stride_fp16: int, |
| ) -> int: |
| """ |
| Compute LDS byte address for the b64 read of operand A for a given lane. |
| Assumes row-major storage with row stride = row_stride_fp16 (FP16 elements). |
| Address = 2 * [ m * row_stride_fp16 + k_start ] |
| where m = lane >> 2, k_start = (lane & 0x3) << 2 |
| """ |
| m = lane >> 2 |
| k_start = (lane & 0x3) << 2 |
| index = m * row_stride_fp16 + k_start |
| return 2 * index |
|
|
| def address_B( |
| lane: int, |
| col_stride_fp16: int, |
| ) -> int: |
| """ |
| Compute LDS byte address for the b64 read of operand B for a given lane. |
| Assumes column-major storage with column stride = col_stride_fp16 (FP16 elements). |
| Address = 2 * [ n * col_stride_fp16 + k_start ] |
| where k_start = (lane >> 4) << 2, n = lane & 0xF |
| """ |
| k_start = (lane >> 4) << 2 |
| n = lane & 0xF |
| index = n * col_stride_fp16 + k_start |
| return 2 * index |
|
|
| |
| |
| |
| DS_READ_B128_GROUPS = [ |
| list(range(0, 4)) + list(range(20, 24)), |
| list(range(4, 8)) + list(range(16, 20)), |
| list(range(8, 12)) + list(range(28, 32)), |
| list(range(12, 16)) + list(range(24, 28)), |
| list(range(32, 36)) + list(range(52, 56)), |
| list(range(36, 40)) + list(range(48, 52)), |
| list(range(40, 44)) + list(range(60, 64)), |
| list(range(44, 48)) + list(range(56, 60)), |
| ] |
|
|
| def conflict_report_b64( |
| read_ops: List[ReadOp], |
| address_of: Callable[[int], int] |
| ) -> List[dict]: |
| conflicts = [] |
| for gid, group in enumerate(DS_READ_B128_GROUPS): |
| for q in range(2): |
| bank_to_entries: Dict[int, List[Tuple[int, int]]] = {} |
| for lane in group: |
| addr = address_of(lane) |
| if addr % 4 != 0: |
| conflicts.append({ |
| "kind": "misalignment", |
| "group": gid, |
| "q": q, |
| "lane": lane, |
| "base_addr": addr, |
| }) |
| continue |
| word_addr = (addr // 4) + q |
| bank = word_addr % 32 |
| bank_to_entries.setdefault(bank, []).append((lane, word_addr)) |
| for bank, entries in bank_to_entries.items(): |
| distinct = {wd for _, wd in entries} |
| if len(distinct) > 1: |
| conflicts.append({ |
| "kind": "bank-conflict", |
| "group": gid, |
| "q": q, |
| "bank": bank, |
| "accesses": entries, |
| "way": len(distinct), |
| }) |
| return conflicts |
|
|
| def has_conflict_b64(read_ops: List[ReadOp], address_of: Callable[[int], int]) -> bool: |
| return bool(conflict_report_b64(read_ops, address_of)) |
|
|
| |
| |
| |
| def find_layout_padding( |
| address_func: Callable[[int, int], int], |
| max_padding: int = 32 |
| ) -> Optional[Dict]: |
| """ |
| Tries padding (making the stride even) to eliminate b64 bank conflicts. |
| Returns the first layout (dict) that yields zero conflicts and 4-byte alignment. |
| """ |
| for P in range(max_padding + 1): |
| stride = 16 + P |
| if stride % 2 != 0: |
| continue |
| |
| def addr_fn(lane_id: int) -> int: |
| return address_func(lane_id, stride) |
| |
| |
| reads = [ReadOp(lane=i, operand="dummy", address=0, width_bytes=64) for i in range(64)] |
| |
| reads_with_addr = [ |
| ReadOp( |
| lane=read.lane, |
| operand=read.operand, |
| address=addr_fn(read.lane), |
| width_bytes=read.width_bytes |
| ) |
| for read in reads |
| ] |
| if not has_conflict_b64(reads_with_addr, addr_fn): |
| return { |
| "kind": "padded", |
| "pad_words": P, |
| "stride_fp16": stride, |
| "conflicts": [] |
| } |
| return None |
|
|
| |
| |
| |
| if __name__ == "__main__": |
| |
| frags = generate_v_mfma_f32_16x16x16f16_fragments() |
| validate_fragment_map(frags) |
| print("Fragment map validation passed.") |
| |
| |
| plan_a = build_read_plan_b64(frags["A"], operand="A") |
| plan_b = build_read_plan_b64(frags["B"], operand="B") |
| |
| print("\n=== Operand A (row-major) ===") |
| layout_a = find_layout_padding(address_A, max_padding=32) |
| if layout_a: |
| print(f"Layout: {layout_a['kind']}") |
| print(f" Padding: {layout_a['pad_words']} FP16 elements") |
| print(f" Row stride: {layout_a['stride_fp16']} FP16 elements") |
| print(f" = {layout_a['stride_fp16'] * 2} bytes") |
| else: |
| print("No conflict-free padding found for A") |
|
|
| print("\n=== Operand B (column-major) ===") |
| layout_b = find_layout_padding(address_B, max_padding=32) |
| if layout_b: |
| print(f"Layout: {layout_b['kind']}") |
| print(f" Padding: {layout_b['pad_words']} FP16 elements") |
| print(f" Column stride: {layout_b['stride_fp16']} FP16 elements") |
| print(f" = {layout_b['stride_fp16'] * 2} bytes") |
| else: |
| print("No conflict-free padding found for B") |
|
|
| |
| if layout_a and layout_b: |
| cert = { |
| "target": "gfx942", |
| "opcode": "v_mfma_f32_16x16x16f16", |
| "wavefront_size": 64, |
| "mfma_tile": { "M": 16, "N": 16, "K": 16 }, |
| "operand_A": { |
| "fragment_map_sha256": "TODO", |
| "lds_layout": { |
| "kind": layout_a["kind"], |
| "row_stride_fp16": layout_a["stride_fp16"], |
| "pad_words": layout_a["pad_words"], |
| }, |
| "load": "ds_read_b64", |
| "conflicts": layout_a["conflicts"] |
| }, |
| "operand_B": { |
| "fragment_map_sha256": "TODO", |
| "lds_layout": { |
| "kind": layout_b["kind"], |
| "col_stride_fp16": layout_b["stride_fp16"], |
| "pad_words": layout_b["pad_words"], |
| }, |
| "load": "ds_read_b64", |
| "conflicts": layout_b["conflicts"] |
| } |
| } |
| import json |
| print("\n=== Layout Certificate ===") |
| print(json.dumps(cert, indent=2)) |