draft_main_proj.py 5.6 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120
  1. """The DSpark draft's main_proj, column-split over the TP group. Default OFF.
  2. DSV41_DRAFT_MAIN_PROJ_SPLIT=1. stage0.main_proj (hidden*4 -> hidden, MXFP8) is a ReplicatedLinear:
  3. every rank streams the whole 105 MB weight to compute the same [M, 5120] output, 430 us a step.
  4. Here each rank keeps only its 1280 output columns (e4m3 with the original per-row 32-element
  5. power-of-two scales, so the weight bytes are the same values) and the columns are all-gathered
  6. (RoCE at small M). Differences to the stock path: the activation stays bf16 instead of MXFP8, and
  7. the reduction order is the Triton kernel's. Only the DRAFT changes, so every accepted token stays
  8. exact; acceptance is what to watch. Rows above MAX_M (prefill extends) take the stock path.
  9. The dequantized weight is read out of the layer itself (main_proj applied to identity rows is
  10. exact: a lone 1.0 per 32-block quantizes to MXFP8 without error and each output is one product),
  11. so nothing depends on the MXFP8 scale layout of the installed backend.
  12. """
  13. import os
  14. import torch
  15. import triton
  16. import triton.language as tl
  17. ENABLED = os.environ.get("DSV41_DRAFT_MAIN_PROJ_SPLIT", "0").strip() not in ("0", "", "off", "false")
  18. MAX_M = int(os.environ.get("DSV41_DRAFT_MAIN_PROJ_MAX_M", "16"))
  19. _S = {"twin": None}
  20. @triton.jit
  21. def _gemv_rowscale_kernel(X, W8, S, Y, M, N, K: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
  22. n = tl.program_id(0) * BN + tl.arange(0, BN)
  23. m = tl.program_id(1) * 16 + tl.arange(0, 16)
  24. nmask = n < N
  25. acc = tl.zeros((16, BN), tl.float32)
  26. for k0 in range(0, K, BK):
  27. k = k0 + tl.arange(0, BK)
  28. x = tl.load(X + m[:, None] * K + k[None, :], m[:, None] < M, 0.0)
  29. w8 = tl.load(W8 + n[None, :] * K + k[:, None], nmask[None, :], 0.0)
  30. e = tl.load(S + n[None, :] * (K // 32) + k[:, None] // 32, nmask[None, :], 127)
  31. w = (w8.to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)).to(tl.bfloat16)
  32. acc += tl.dot(x, w)
  33. tl.store(Y + m[:, None] * N + n[None, :], acc.to(tl.bfloat16), (m[:, None] < M) & nmask[None, :])
  34. def gemv(x, twin):
  35. w8, s = twin
  36. m, k = x.shape
  37. n = w8.shape[0]
  38. y = torch.empty((m, n), dtype=torch.bfloat16, device=x.device)
  39. # GB10, N=1280 K=20480: BN 8 / BK 256 / 2 warps = 124 us at M=6 (stock replicated MXFP8: 430 us)
  40. _gemv_rowscale_kernel[(triton.cdiv(n, 8), triton.cdiv(m, 16))](x.contiguous(), w8, s, y, m, n, k, 8, 256,
  41. num_warps=2, num_stages=3)
  42. return y
  43. def quant_rowscale(w):
  44. """bf16 [N, K] (values already on an e4m3 x 2^e grid per 32-block) -> (e4m3, biased exponent)."""
  45. n, k = w.shape
  46. wf = w.float().view(n, k // 32, 32)
  47. amax = wf.abs().amax(-1).clamp_min(2.0 ** -126)
  48. e = torch.ceil(torch.log2(amax / 448.0)).clamp(-127, 127)
  49. q = (wf / torch.exp2(e)[..., None]).to(torch.float8_e4m3fn)
  50. return q.view(n, k).contiguous(), (e + 127).to(torch.uint8).contiguous()
  51. def _build(layer, group):
  52. k = layer.input_size if hasattr(layer, "input_size") else None
  53. w = getattr(layer, "weight", None)
  54. n = layer.output_size if hasattr(layer, "output_size") else w.shape[0]
  55. k = k or w.shape[1]
  56. world, rank = group.world_size, group.rank_in_group
  57. if n % world or k % 32:
  58. raise RuntimeError(f"main_proj {n}x{k} does not split over {world} ranks")
  59. dev = torch.device("cuda", torch.cuda.current_device())
  60. per = n // world
  61. cols = []
  62. step = 1024
  63. for k0 in range(0, k, step):
  64. c = min(step, k - k0)
  65. eye = torch.zeros((c, k), dtype=torch.bfloat16, device=dev)
  66. eye[torch.arange(c, device=dev), k0 + torch.arange(c, device=dev)] = 1.0
  67. out = layer(eye)
  68. out = out[0] if isinstance(out, tuple) else out
  69. cols.append(out[:, rank * per:(rank + 1) * per].to(torch.bfloat16)) # W[cols, k0:k0+c]^T
  70. del eye, out
  71. local = torch.cat(cols, 0).t().contiguous() # [per, k]
  72. del cols
  73. twin = quant_rowscale(local)
  74. back = (twin[0].float().view(per, k // 32, 32) * torch.exp2(twin[1].float() - 127)[..., None]).view(per, k)
  75. exact = torch.equal(back.to(torch.bfloat16), local)
  76. del local, back
  77. torch.cuda.empty_cache()
  78. print(f"[draft_main_proj] rank {rank}: column shard {per}x{k} fp8 (weight bytes exact: {exact})", flush=True)
  79. return twin
  80. def install(dspark_module):
  81. """sglang.srt.models.deepseek_v4_dspark: DeepseekV4ForCausalLMDSpark.project_target_hidden."""
  82. if not ENABLED:
  83. return
  84. cls = dspark_module.DeepseekV4ForCausalLMDSpark
  85. if not hasattr(cls, "project_target_hidden"):
  86. raise RuntimeError("DSV41_DRAFT_MAIN_PROJ_SPLIT: project_target_hidden is gone; engine drifted")
  87. orig = cls.project_target_hidden
  88. def project_target_hidden(self, main_hidden):
  89. m = main_hidden.shape[0]
  90. if m == 0 or m > MAX_M or main_hidden.dim() != 2:
  91. return orig(self, main_hidden)
  92. from sglang.srt.distributed import get_tp_group
  93. group = get_tp_group()
  94. if _S["twin"] is None:
  95. if torch.cuda.is_current_stream_capturing():
  96. return orig(self, main_hidden)
  97. # no per-rank fallback: a rank on the stock path would skip the all-gather and hang the
  98. # others, so a failed build fails the boot
  99. _S["twin"] = _build(self.stages[0].main_proj, group)
  100. local = gemv(main_hidden.to(torch.bfloat16), _S["twin"])
  101. projected = group.all_gather(local, dim=-1) if group.world_size > 1 else local
  102. return self.stages[0].main_norm(projected)
  103. cls.project_target_hidden = project_target_hidden
  104. print(f"[draft_main_proj] armed (rows <= {MAX_M})", flush=True)