nf3_replan.py 7.4 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165
  1. """
  2. NF3 "2+1 split-plane" repack (load-time, VRAM) + expert-space permutation helpers.
  3. Disk checkpoint stays in the native "837" pack (8 codes -> 24-bit LE word -> 3 bytes,
  4. code k = (word >> 3*(k%8)) & 7, uint8 [E, N, K//8*3]). At load we repack to the
  5. split-plane layout the b12x NF3 kernel format consumes:
  6. per 32-code scale group -> 12 bytes:
  7. LO plane: 2 u32 = 32 x 2-bit fields (code j's bits 1:0 at bits 2j+1:2j of the 64-bit lo pair)
  8. HI plane: 1 u32 = 32 x 1-bit fields (code j's bit 2 at bit j)
  9. 128-code superblock = 48 B = 3 x 16 B cp.async chunks (25% fewer B bytes than 4-bit).
  10. Rows stay 16B-multiples at K=6144 (2304 B) and K=512 (192 B).
  11. In-kernel selector build for one prmt (4 codes), documented here, implemented in fp4.py:
  12. lo8 = (lo_u32 >> (8*g)) & 0xFF # 4 codes' 2-bit fields, group g of 4
  13. hi4 = (hi_u32 >> (4*g)) & 0xF # 4 codes' high bits
  14. # spread2: 4 x 2-bit fields (stride 2) -> nibble slots (stride 4); shift-ladder, carry-safe
  15. t = (lo8 | (lo8 << 4)) & 0x0F0F
  16. s2 = (t | (t << 2)) & 0x3333
  17. # spread1: 4 x 1-bit fields -> nibble slots
  18. u = (hi4 | (hi4 << 6)) & 0x0303
  19. s1 = (u | (u << 3)) & 0x1111
  20. sel = s2 | (s1 << 2) # 4 prmt selector nibbles, each = the 3-bit code
  21. -> 2 pool-prmt (8 lo / 8 hi codebook bytes in 4 regs) -> 2 interleave-prmt
  22. (0x5140, 0x7362) -> 2 x bf16x2; then mul by the widened e4m3_k32 scale (bf16x2).
  23. Kernel-facing tensor (per proj): uint32 [E, N, K//32, 3] contiguous, u32[0..1]=LO, u32[2]=HI.
  24. Scales stay checkpoint-native: float8_e4m3fn [E, N, K//32].
  25. NOTE: the intra-group code order (which K element lands at field j) is IDENTITY for now;
  26. if the b12x fragment map consumes K in an interleaved order (k16 two-lane style), set
  27. CODE_ORDER from the K1 read-map so field j = the j-th code the kernel consumes.
  28. """
  29. import torch
  30. NF3_CODEBOOK = (-1.0, -0.6047, -0.3563, -0.1275, 0.1275, 0.3563, 0.6047, 1.0)
  31. CODE_ORDER = None # optional [32] permutation: field j holds code group_base + CODE_ORDER[j]
  32. def unpack_837_codes(packed: torch.Tensor, K: int) -> torch.Tensor:
  33. """uint8 [E, N, K//8*3] (837 layout) -> int32 codes [E, N, K]."""
  34. E, N, _ = packed.shape
  35. b = packed.reshape(E, N, K // 8, 3).to(torch.int32)
  36. word = b[..., 0] | (b[..., 1] << 8) | (b[..., 2] << 16) # [E,N,K//8]
  37. shifts = torch.arange(8, device=packed.device, dtype=torch.int32) * 3
  38. codes = (word.unsqueeze(-1) >> shifts) & 7 # [E,N,K//8,8]
  39. return codes.reshape(E, N, K)
  40. def pack_2p1(codes: torch.Tensor, code_order=None) -> torch.Tensor:
  41. """int codes [E, N, K] (0..7) -> uint32 [E, N, K//32, 3] split-plane (LO,LO,HI)."""
  42. E, N, K = codes.shape
  43. assert K % 32 == 0, K
  44. c = codes.to(torch.int64).reshape(E, N, K // 32, 32) # [E,N,G,32]
  45. if code_order is not None:
  46. c = c[..., code_order]
  47. j = torch.arange(32, device=codes.device, dtype=torch.int64)
  48. lo = c & 3 # bits 1:0
  49. hi = (c >> 2) & 1 # bit 2
  50. lo64 = (lo << (2 * j)).sum(-1, dtype=torch.int64) # [E,N,G] 64-bit
  51. lo_w0 = lo64 & 0xFFFFFFFF
  52. lo_w1 = (lo64 >> 32) & 0xFFFFFFFF
  53. hi_w = (hi << j).sum(-1, dtype=torch.int64) & 0xFFFFFFFF
  54. return torch.stack([lo_w0, lo_w1, hi_w], dim=-1).to(torch.uint32)
  55. def unpack_2p1(planes: torch.Tensor, code_order=None) -> torch.Tensor:
  56. """uint32 [E, N, G, 3] -> int32 codes [E, N, G*32] (inverse of pack_2p1)."""
  57. E, N, G, _ = planes.shape
  58. p = planes.to(torch.int64)
  59. lo64 = p[..., 0] | (p[..., 1] << 32)
  60. j = torch.arange(32, device=planes.device, dtype=torch.int64)
  61. lo = (lo64.unsqueeze(-1) >> (2 * j)) & 3 # [E,N,G,32]
  62. hi = (p[..., 2].unsqueeze(-1) >> j) & 1
  63. c = (lo | (hi << 2))
  64. if code_order is not None:
  65. inv = torch.empty_like(torch.as_tensor(code_order))
  66. inv[torch.as_tensor(code_order)] = torch.arange(32)
  67. c = c[..., inv]
  68. return c.reshape(E, N, G * 32).to(torch.int32)
  69. def repack_837_to_2p1(packed_837: torch.Tensor, K: int, code_order=None) -> torch.Tensor:
  70. """One-shot load-time repack, chunked over E to bound peak VRAM."""
  71. E, N, _ = packed_837.shape
  72. out = torch.empty(E, N, K // 32, 3, dtype=torch.uint32, device=packed_837.device)
  73. step = max(1, min(E, 8))
  74. for e0 in range(0, E, step):
  75. codes = unpack_837_codes(packed_837[e0:e0 + step], K)
  76. out[e0:e0 + step] = pack_2p1(codes, code_order)
  77. return out
  78. def hybrid_expert_permutation(kept_global_ids: torch.Tensor, n_experts: int = 256) -> torch.Tensor:
  79. """perm[new_id] = old_id with kept (NVFP4) experts first: new 0..63 = kept, 64.. = NF3.
  80. Apply to: expert weight order, gate.weight rows, e_score_correction_bias entries.
  81. n_group==1/topk_group==1 on GLM-5.2 -> renumbering is routing-neutral.
  82. """
  83. kept = kept_global_ids.to(torch.long)
  84. mask = torch.ones(n_experts, dtype=torch.bool)
  85. mask[kept] = False
  86. rest = torch.arange(n_experts)[mask]
  87. return torch.cat([kept, rest])
  88. def selector_build_reference(lo8: int, hi4: int) -> int:
  89. """Host-side reference of the in-kernel selector build (4 codes -> 4 prmt nibbles)."""
  90. t = (lo8 | (lo8 << 4)) & 0x0F0F
  91. s2 = (t | (t << 2)) & 0x3333
  92. u = (hi4 | (hi4 << 6)) & 0x0303
  93. s1 = (u | (u << 3)) & 0x1111
  94. return s2 | (s1 << 2)
  95. if __name__ == "__main__":
  96. torch.manual_seed(0)
  97. dev = "cuda" if torch.cuda.is_available() else "cpu"
  98. import sys
  99. sys.path.insert(0, "/Users/phantom/Downloads/glm52-hybrid")
  100. from nf3_kernel import pack_nf3_837
  101. E, N, K = 3, 8, 128
  102. codes = torch.randint(0, 8, (E, N, K), device=dev)
  103. p837 = pack_nf3_837(codes)
  104. rt = unpack_837_codes(p837, K)
  105. assert torch.equal(rt, codes.to(torch.int32)), "837 unpack mismatch"
  106. planes = pack_2p1(codes)
  107. rt2 = unpack_2p1(planes)
  108. assert torch.equal(rt2, codes.to(torch.int32)), "2+1 round-trip mismatch"
  109. planes2 = repack_837_to_2p1(p837, K)
  110. assert torch.equal(planes2, planes), "837->2p1 repack mismatch"
  111. print("PACK ROUND-TRIP: PASS")
  112. # selector build: exhaustive over all 4-code combos (8^4 = 4096)
  113. for c0 in range(8):
  114. for c1 in range(8):
  115. for c2 in range(8):
  116. for c3 in range(8):
  117. quad = [c0, c1, c2, c3]
  118. lo8 = sum((c & 3) << (2 * j) for j, c in enumerate(quad))
  119. hi4 = sum(((c >> 2) & 1) << j for j, c in enumerate(quad))
  120. sel = selector_build_reference(lo8, hi4)
  121. for j in range(4):
  122. got = (sel >> (4 * j)) & 0xF
  123. assert got == quad[j], (quad, j, got)
  124. print("SELECTOR BUILD: PASS (exhaustive 8^4)")
  125. # cross-check selector inputs against packed planes
  126. g = 3 # arbitrary group-of-4 within the 32-code group
  127. grp = codes[0, 0, :32]
  128. lo64 = int(planes[0, 0, 0, 0]) | (int(planes[0, 0, 0, 1]) << 32)
  129. hiw = int(planes[0, 0, 0, 2])
  130. lo8 = (lo64 >> (8 * g)) & 0xFF
  131. hi4 = (hiw >> (4 * g)) & 0xF
  132. sel = selector_build_reference(lo8, hi4)
  133. for j in range(4):
  134. assert (sel >> (4 * j)) & 0xF == int(grp[4 * g + j]), "plane->selector mismatch"
  135. print("PLANE->SELECTOR: PASS")
  136. kept = torch.randperm(256)[:64]
  137. perm = hybrid_expert_permutation(kept)
  138. assert perm.shape == (256,) and len(set(perm.tolist())) == 256
  139. assert torch.equal(perm[:64], kept.to(torch.long))
  140. print("PERMUTATION: PASS")