From 7ea8908d83827dd8d82c34ba6a60b2beaa8057d6 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 12:58:06 +0800 Subject: [PATCH 01/25] [Kernel] Fuse SM70 adapter projection epilogues Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- csrc/sm70_turbomind/ops/diffusion_epilogue.h | 99 ++++++++++++ csrc/sm70_turbomind/ops/h3_w8a16.cu | 4 + docs/design/minimax_h3/SM70_EPILOGUES.md | 50 +++++++ tests/video/test_h3_prepared_linear.py | 9 +- tests/video/test_sm70_scaled_add.py | 141 ++++++++++++++++++ vllm/model_executor/layers/sm70_diffusion.py | 34 +++++ vllm/model_executor/models/minimax_h3/lora.py | 31 +++- 7 files changed, 359 insertions(+), 9 deletions(-) create mode 100644 csrc/sm70_turbomind/ops/diffusion_epilogue.h create mode 100644 docs/design/minimax_h3/SM70_EPILOGUES.md create mode 100644 tests/video/test_sm70_scaled_add.py diff --git a/csrc/sm70_turbomind/ops/diffusion_epilogue.h b/csrc/sm70_turbomind/ops/diffusion_epilogue.h new file mode 100644 index 0000000000..9e53dda799 --- /dev/null +++ b/csrc/sm70_turbomind/ops/diffusion_epilogue.h @@ -0,0 +1,99 @@ +// SPDX-License-Identifier: Apache-2.0 +// SPDX-FileCopyrightText: Copyright contributors to the vLLM project +#pragma once + +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +namespace sm70_diffusion { +template +__global__ void scaled_add_rows(Output* output, const float* delta, + const float* scales, int64_t count, + int64_t width, int64_t output_width, + int64_t offset, float alpha) { + for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x; + index < count; index += int64_t(gridDim.x) * blockDim.x) { + const int64_t row = index / width; + const int64_t col = index - row * width; + const int64_t destination = row * output_width + offset + col; + // Retain the explicit FP32 scale-restoration boundary before addition. + const float restored = + scales ? __fmul_rn(delta[index], scales[row]) : delta[index]; + float base; + if constexpr (std::is_same_v) { + base = __half2float(output[destination]); + } else { + base = output[destination]; + } + const float result = __fmaf_rn(alpha, restored, base); + if constexpr (std::is_same_v) { + output[destination] = __float2half_rn(result); + } else { + output[destination] = result; + } + } +} + +inline torch::Tensor scaled_add(torch::Tensor output, torch::Tensor delta, + std::optional scales, + double alpha, int64_t offset) { + TORCH_CHECK(output.is_cuda() && output.dim() == 2 && output.is_contiguous(), + "SM70 scaled addition requires contiguous CUDA [M,N] output"); + TORCH_CHECK(output.scalar_type() == torch::kFloat16 || + output.scalar_type() == torch::kFloat32, + "SM70 scaled addition output must be FP16 or FP32"); + TORCH_CHECK(delta.device() == output.device() && delta.dim() == 2 && + delta.is_contiguous() && + delta.scalar_type() == torch::kFloat32 && + delta.size(0) == output.size(0), + "SM70 scaled addition requires matching FP32 [M,K] delta"); + TORCH_CHECK(offset >= 0 && offset <= output.size(1) && + delta.size(1) <= output.size(1) - offset, + "SM70 scaled addition slice is outside output"); + TORCH_CHECK(std::isfinite(alpha) && + std::abs(alpha) <= std::numeric_limits::max(), + "SM70 scaled addition alpha must be finite FP32"); + TORCH_CHECK(!output.requires_grad() && !delta.requires_grad(), + "SM70 scaled addition is inference-only"); + at::assert_no_overlap(output, delta); + const c10::cuda::CUDAGuard guard(output.device()); + const auto* properties = at::cuda::getCurrentDeviceProperties(); + TORCH_CHECK(properties->major == 7 && properties->minor == 0, + "SM70 scaled addition requires SM70"); + const float* scale_data = nullptr; + if (scales.has_value()) { + TORCH_CHECK( + scales->device() == output.device() && scales->is_contiguous() && + scales->scalar_type() == torch::kFloat32 && + scales->numel() == output.size(0) && !scales->requires_grad(), + "SM70 scaled addition needs one FP32 scale per row"); + at::assert_no_overlap(output, *scales); + scale_data = scales->data_ptr(); + } + if (!delta.numel()) return output; + const int blocks = std::min((delta.numel() + 255) / 256, 65535); + const auto stream = at::cuda::getCurrentCUDAStream(); + if (output.scalar_type() == torch::kFloat16) { + scaled_add_rows<<>>( + reinterpret_cast(output.data_ptr()), + delta.data_ptr(), scale_data, delta.numel(), delta.size(1), + output.size(1), offset, float(alpha)); + } else { + scaled_add_rows<<>>( + output.data_ptr(), delta.data_ptr(), scale_data, + delta.numel(), delta.size(1), output.size(1), offset, float(alpha)); + } + C10_CUDA_KERNEL_LAUNCH_CHECK(); + return output; +} +} // namespace sm70_diffusion diff --git a/csrc/sm70_turbomind/ops/h3_w8a16.cu b/csrc/sm70_turbomind/ops/h3_w8a16.cu index 27a4aeec5a..72383c9387 100644 --- a/csrc/sm70_turbomind/ops/h3_w8a16.cu +++ b/csrc/sm70_turbomind/ops/h3_w8a16.cu @@ -8,6 +8,7 @@ #include #include "h3_column_major_gemm.h" +#include "diffusion_epilogue.h" namespace { __global__ void prepare_fp16_rows(const float* input, half* output, @@ -256,6 +257,9 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.def("prepare_fp16", &h3_prepare_fp16); m.def("dequantize", &h3_dequantize); m.def("rotate", &h3_rotate); + m.def("scaled_add_", &sm70_diffusion::scaled_add, pybind11::arg("output"), + pybind11::arg("delta"), pybind11::arg("scales"), pybind11::arg("alpha"), + pybind11::arg("offset") = 0); m.def("gemm", &h3_fp16_gemm, pybind11::arg("input"), pybind11::arg("weight"), pybind11::arg("output_fp32") = false); } diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md new file mode 100644 index 0000000000..6163a40ed1 --- /dev/null +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -0,0 +1,50 @@ +# Shared SM70 projection epilogues + +This development branch is stacked on the common prepared execution (#571) +and workflow accounting (#578) branches. No configuration has passed the +campaign's >80 useful TFLOP/s/card and complete official quality gates. + +## Measured problem and implementation + +The matching four-step FA denoise profile spends 6.890 seconds in miscellaneous +elementwise kernels, including FP32 row-scale restoration, adapter additions +and output casts. Attention, GEMM and communication separately consume +31.439, 16.041 and 8.453 seconds. These are profiled service diagnostics; +the unprofiled audited baseline is 47.091839–47.091855 useful TFLOP/s/card. + +`sm70_diffusion.fp16_linear_add` prepares the projection input at the existing +FP16 boundary and keeps GEMM accumulation in FP32. A shared CUDA epilogue +restores row scales with a rounded FP32 multiplication, adds the scaled delta +with the same FP32 fused multiply-add as PyTorch, and stores FP16 or FP32. +It writes only the requested output slice. Delta/output and scale/output +storage overlap is rejected, including differently typed views of one buffer. + +The H3 adapter uses this interface when output hardware, precision and layout +support it and adapter slices are disjoint. Overlapping contributions retain +the original FP32 buffer until the final cast. Wheels without the new extension +ABI retain the ordinary path. No quantization label or adapter name controls +dispatch. The first LoRA projection still uses unrotated inputs; row-parallel +increments still join the partial result before the collective. + +## Development evidence + +Environment: Python 3.12.13, Torch 2.10.0+cu128, CUDA toolkit 12.8.93, +V100 SXM2 32GB. All GPU tests use an owned native lease. Evidence root: +`/data/minimax-h3/sm70-general-20260909/`. + +- `epilogue-integration-v1.log`: 43 checks pass, including explicit comparison + with the old unfused adapter path for original/W8A16 bases, FP16/FP32 output, + prepared/ordinary inputs, negative scales, wide intermediates, untouched + slices, overlap fallback, unaligned storage and CUDA Graph replay. +- `epilogue-cpu-v1.log`: 94 adapter, workflow and strict acceptance checks pass. +- `epilogue-micro.json`: paired postprocessing-only medians for 34,560 rows, + output width 5,376: three FP16 QKV slices 7.524352 -> 2.015232 ms; one FP32 + row projection 4.562944 -> 3.094528 ms. Results match bitwise. These timings + exclude GEMM and do not establish complete-request speedup. +- `epilogue-binaries.json` retains the first tested binary. The strengthened + alias check is in `epilogue-binaries-v3.json`; validation is recorded separately. + +Complete four-step latent/RGB/PCM comparison and one full warmup plus three +unprofiled requests remain required before this change can be promoted. +Independent official reference and full audiovisual review remain required +even if frozen-mainline preservation passes. diff --git a/tests/video/test_h3_prepared_linear.py b/tests/video/test_h3_prepared_linear.py index 54d3699d7d..6c5a3cd4bc 100644 --- a/tests/video/test_h3_prepared_linear.py +++ b/tests/video/test_h3_prepared_linear.py @@ -68,7 +68,10 @@ def test_dense_layout_keeps_logical_weights_and_wide_output(layout): @pytest.mark.skipif(not torch.cuda.is_available(), reason="requires SM70") @pytest.mark.parametrize("quantized", [False, True]) @pytest.mark.parametrize("adapter_scale", [0.0, 0.75, -0.5]) -def test_gpu_prepared_adapter_preserves_wide_intermediates(quantized, adapter_scale): +def test_gpu_prepared_adapter_preserves_wide_intermediates( + quantized, adapter_scale, monkeypatch +): + import vllm.model_executor.models.minimax_h3.lora as adapter from vllm.model_executor.models.minimax_h3.cuda_ops import w8a16_extension torch.manual_seed(42) @@ -93,7 +96,9 @@ def test_gpu_prepared_adapter_preserves_wide_intermediates(quantized, adapter_sc values, scale = fp16_gemm_input(x) token = lora_scale.set(adapter_scale) try: - expected = method.apply(layer, x) + with monkeypatch.context() as context: + context.setattr(adapter, "supports_fused_scaled_add", lambda _: False) + expected = method.apply(layer, x) actual = method.apply_prepared(layer, values, scale) torch.testing.assert_close(actual, expected, atol=0, rtol=0) assert actual.dtype == torch.float32 and torch.isfinite(actual).all() diff --git a/tests/video/test_sm70_scaled_add.py b/tests/video/test_sm70_scaled_add.py new file mode 100644 index 0000000000..d99e93f37e --- /dev/null +++ b/tests/video/test_sm70_scaled_add.py @@ -0,0 +1,141 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +from types import SimpleNamespace + +import pytest +import torch + +from vllm.model_executor.layers.sm70_diffusion import sm70_extension + +pytestmark = pytest.mark.skipif( + not torch.cuda.is_available(), reason="requires a leased SM70 GPU" +) + + +@pytest.mark.parametrize("dtype", [torch.float16, torch.float32]) +@pytest.mark.parametrize("alpha", [0.0, 0.0625, 0.75, -0.5, 1.0]) +@pytest.mark.parametrize("scaled", [False, True]) +def test_scaled_add_retains_fp32_rounding_and_untouched_slices(dtype, alpha, scaled): + torch.manual_seed(610) + rows, columns, width, offset = 17, 43, 33, 3 + storage = torch.randn(rows * columns + 1, device="cuda", dtype=dtype) + actual = storage[1:].reshape(rows, columns) + base = actual.clone() + delta = torch.randn(rows, width, device="cuda") * 100 + scales = ( + torch.ldexp( + torch.ones(rows, 1, device="cuda"), + (torch.arange(rows, device="cuda") - 8)[:, None], + ) + if scaled + else None + ) + restored = delta if scales is None else delta * scales + expected = base.float().clone() + expected[:, offset : offset + width].add_(restored, alpha=alpha) + result = sm70_extension().scaled_add_(actual, delta, scales, alpha, offset) + assert result.data_ptr() == actual.data_ptr() + torch.testing.assert_close(actual, expected.to(dtype), rtol=0, atol=0) + torch.testing.assert_close(actual[:, :offset], base[:, :offset], rtol=0, atol=0) + torch.testing.assert_close( + actual[:, offset + width :], base[:, offset + width :], rtol=0, atol=0 + ) + + +def test_scaled_add_rejects_unsafe_memory_aliases_and_bad_bounds(): + data = torch.zeros(65, device="cuda") + output = data[:-1].reshape(8, 8) + overlapping = data[1:].reshape(8, 8) + op = sm70_extension().scaled_add_ + with pytest.raises(RuntimeError): + op(output, overlapping, None, 1.0, 0) + with pytest.raises(RuntimeError): + op(output.view(torch.float16), output, None, 1.0, 0) + with pytest.raises(RuntimeError): + op(output, torch.ones_like(output), data[:8], 1.0, 0) + with pytest.raises(RuntimeError, match="outside output"): + op(output, torch.ones_like(output), None, 1.0, 1) + with pytest.raises(RuntimeError, match="finite FP32"): + op(output, torch.ones_like(output), None, float("nan"), 0) + + +def test_scaled_add_graph_reads_current_operands(): + output = torch.zeros(16, 192, device="cuda", dtype=torch.float16) + delta = torch.randn(16, 192, device="cuda") + scales = torch.ones(16, 1, device="cuda") * 8 + op = sm70_extension().scaled_add_ + stream = torch.cuda.Stream() + stream.wait_stream(torch.cuda.current_stream()) + with torch.cuda.stream(stream): + op(output, delta, scales, 0.75, 0) + torch.cuda.current_stream().wait_stream(stream) + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + op(output, delta, scales, 0.75, 0) + for factor in (1.0, -2.0): + output.fill_(factor) + delta.mul_(factor) + expected = output.float().add(delta * scales, alpha=0.75).half() + graph.replay() + torch.testing.assert_close(output, expected, rtol=0, atol=0) + + +@pytest.mark.parametrize("quantized", [False, True]) +@pytest.mark.parametrize("output_fp32", [False, True]) +@pytest.mark.parametrize("overlapping", [False, True]) +def test_adapter_fusion_matches_unfused_rounding( + quantized, output_fp32, overlapping, monkeypatch +): + import vllm.model_executor.models.minimax_h3.lora as adapter + from vllm.model_executor.models.minimax_h3.quantization import ( + DiffusionInt8ConvRotConfig, + FP16LinearMethod, + FP32OutputLinearMethod, + Int8ConvRotLayerConfig, + Int8ConvRotLinearMethod, + ) + + torch.manual_seed(913) + if quantized: + base = Int8ConvRotLinearMethod( + DiffusionInt8ConvRotConfig(), Int8ConvRotLayerConfig(True), prefix="probe" + ) + weight = torch.randint(-7, 8, (384, 256), device="cuda", dtype=torch.int8) + else: + base = FP32OutputLinearMethod() if output_fp32 else FP16LinearMethod() + weight = torch.randn(384, 256, device="cuda", dtype=torch.float16) * 0.03 + layer = SimpleNamespace( + weight=weight, + weight_scale=torch.full((384,), 0.01, device="cuda"), + h3_output_fp32=output_fp32, + ) + parts = [] + for i in range(3): + setattr(layer, f"h3_lora_a_{i}", torch.randn(8, 256, device="cuda").half()) + setattr(layer, f"h3_lora_b_{i}", torch.randn(96, 8, device="cuda").half()) + parts.append((i, 32 if overlapping else 16 + i * 112, 96)) + method = adapter.TurboLinearMethod(base, parts, 0.125) + inputs = torch.randn(33, 256, device="cuda", dtype=torch.float16) + fused_add = adapter.fp16_linear_add + calls = [] + + def record(*args, **kwargs): + calls.append(kwargs["offset"]) + return fused_add(*args, **kwargs) + + monkeypatch.setattr(adapter, "fp16_linear_add", record) + token = adapter.lora_scale.set(-0.75) + try: + with monkeypatch.context() as context: + context.setattr(adapter, "supports_fused_scaled_add", lambda _: False) + expected = method.apply(layer, inputs) + actual = method.apply(layer, inputs) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + assert actual.dtype == (torch.float32 if output_fp32 else torch.float16) + assert calls == ([] if overlapping else [16, 128, 240]) + calls.clear() + actual = method.apply_prepared(layer, inputs, None) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + assert calls == ([] if overlapping else [16, 128, 240]) + finally: + adapter.lora_scale.reset(token) diff --git a/vllm/model_executor/layers/sm70_diffusion.py b/vllm/model_executor/layers/sm70_diffusion.py index 68ecee48d9..6a6f33db08 100644 --- a/vllm/model_executor/layers/sm70_diffusion.py +++ b/vllm/model_executor/layers/sm70_diffusion.py @@ -106,3 +106,37 @@ def fp16_linear_prepared(values, weight, scale=None, *, output_fp32=False): if scale is not None: output = output * scale return output.reshape(*values.shape[:-1], weight.shape[0]) + + +def supports_fused_scaled_add(output): + """Old wheels and unsupported output layouts keep ordinary epilogues.""" + return ( + output.is_cuda + and output.dtype in (torch.float16, torch.float32) + and output.is_contiguous() + and torch.cuda.get_device_capability(output.device) == (7, 0) + and hasattr(sm70_extension(), "scaled_add_") + ) + + +def fp16_linear_add(x, weight, output, *, alpha, offset=0): + """Add a scaled FP16 projection into a contiguous output slice in place. + + GEMM and row-scale restoration retain FP32 boundaries. FP16 output is + rounded after this contribution, so callers combining overlapping deltas + must retain an FP32 accumulation buffer until their last contribution. + """ + if not output.is_contiguous(): + raise ValueError("SM70 projection addition requires contiguous output") + values, scale = fp16_gemm_input(x) + delta = fp16_gemm(values, weight, output_fp32=True) + flat = output.view(-1, output.shape[-1]) + if supports_fused_scaled_add(output): + sm70_extension().scaled_add_(flat, delta, scale, alpha, offset) + else: + if scale is not None: + delta = delta * scale + target = flat[:, offset : offset + weight.shape[0]] + result = target.float().add(delta, alpha=alpha).to(output.dtype) + target.copy_(result) + return output diff --git a/vllm/model_executor/models/minimax_h3/lora.py b/vllm/model_executor/models/minimax_h3/lora.py index 2b42e04b2f..af379f1a47 100644 --- a/vllm/model_executor/models/minimax_h3/lora.py +++ b/vllm/model_executor/models/minimax_h3/lora.py @@ -21,7 +21,11 @@ from vllm.logger import init_logger from vllm.model_executor.layers.linear import LinearMethodBase -from vllm.model_executor.layers.sm70_diffusion import fp16_linear_prepared +from vllm.model_executor.layers.sm70_diffusion import ( + fp16_linear_add, + fp16_linear_prepared, + supports_fused_scaled_add, +) from .config import H3InputError from .fasth3 import FASTH3_FILENAME, FastH3Spec @@ -273,6 +277,8 @@ def __init__(self, base, parts, alpha_over_rank): self.base = base self.parts = parts self.alpha_over_rank = alpha_over_rank + spans = sorted((offset, offset + width) for _, offset, width in parts) + self.disjoint_parts = all(a[1] <= b[0] for a, b in zip(spans, spans[1:])) def create_weights(self, *args, **kwargs): raise RuntimeError("Turbo is installed after base checkpoint loading") @@ -304,7 +310,9 @@ def apply_prepared( return output original = original_input if input_is_rotated else values dtype = output.dtype - output = output.float() + fused = self.disjoint_parts and supports_fused_scaled_add(output) + if not fused: + output = output.float() for index, offset, width in self.parts: a = getattr(layer, f"h3_lora_a_{index}") b = getattr(layer, f"h3_lora_b_{index}") @@ -313,8 +321,11 @@ def apply_prepared( intermediate = fp16_linear_prepared( original, a, input_scale, output_fp32=True ) - delta = _linear_fp32(intermediate, b) - output[..., offset : offset + width].add_(delta, alpha=scale) + if fused: + fp16_linear_add(intermediate, b, output, alpha=scale, offset=offset) + else: + delta = _linear_fp32(intermediate, b) + output[..., offset : offset + width].add_(delta, alpha=scale) return output.to(dtype) def apply(self, layer, x, bias=None): @@ -323,12 +334,18 @@ def apply(self, layer, x, bias=None): if scale == 0: return output dtype = output.dtype - output = output.float() + fused = self.disjoint_parts and supports_fused_scaled_add(output) + if not fused: + output = output.float() for index, offset, width in self.parts: a = getattr(layer, f"h3_lora_a_{index}") b = getattr(layer, f"h3_lora_b_{index}") - delta = _linear_fp32(_linear_fp32(x, a), b) - output[..., offset : offset + width].add_(delta, alpha=scale) + intermediate = _linear_fp32(x, a) + if fused: + fp16_linear_add(intermediate, b, output, alpha=scale, offset=offset) + else: + delta = _linear_fp32(intermediate, b) + output[..., offset : offset + width].add_(delta, alpha=scale) return output.to(dtype) From 90f1fcc8a6071221273b5d5519cfa81cdecd9951 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 13:17:00 +0800 Subject: [PATCH 02/25] [Doc] Record complete SM70 epilogue quality and timing controls Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 31 +++++++++++++++++++++--- 1 file changed, 27 insertions(+), 4 deletions(-) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 6163a40ed1..957dafcf5c 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -44,7 +44,30 @@ V100 SXM2 32GB. All GPU tests use an owned native lease. Evidence root: - `epilogue-binaries.json` retains the first tested binary. The strengthened alias check is in `epilogue-binaries-v3.json`; validation is recorded separately. -Complete four-step latent/RGB/PCM comparison and one full warmup plus three -unprofiled requests remain required before this change can be promoted. -Independent official reference and full audiovisual review remain required -even if frozen-mainline preservation passes. +The final alias guard passed on GPU0 (`epilogue-alias-v3-gpu0.log`). An earlier +GPU4 attempt was refused by an existing lease before starting the test. + +## Complete four-step control and measurements + +Source `7ea8908d83827dd8d82c34ba6a60b2beaa8057d6`, TP4, LightX2V four-step v1.2, +W8A16, FA, exact residual sharding, no persistent FP16 weight cache, and the +original 1280x736/124-frame internal canvas for the five-second sample: + +- `epilogue-quality.json`: final video/audio latents and all pre-encoding RGB + frames match frozen mainline bitwise; video SSIM 1, spectral cosine 1 and RMS + ratio 1. This is numerical preservation, not independent official acceptance. +- `epilogue-720p-three-runs/performance.json`: one full warmup (64.373997 seconds + denoise) followed by three complete requests without profiler or captures. + Denoise times are 62.245159 / 62.183194 / 62.162648 seconds; CV 0.056387%. + Every rank reports median **49.905491–49.905510 useful TFLOP/s**. The declared + >80 gate fails and remains incomplete. +- Complete request times are 84.363265 / 91.992465 / 88.620830 seconds. Peak + allocation remains 19,501,498,880 bytes per card. Exact source/kernel hashes, + per-step records and NVML samples are retained beside each run. +- The 62.183194-second median is 5.64% below the audited original FA baseline + (65.898529 seconds). This measures the **combined** prepared/residual/epilogue + changes, not an isolated attribution to this CUDA epilogue. A separate matched + profile is necessary for attribution. + +Independent official reference, full audiovisual review, other adapters and +the complete shape/TP matrix remain required. No AUTO selection is qualified. From 6b39f1c23ca6834e9beead89b4cdf57548101e97 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 13:32:34 +0800 Subject: [PATCH 03/25] [Kernel] Add explicit SM70 attention query geometry Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 17 +++++ flash-attention-v100/kernel/h3/forward.cu | 73 ++++++++++++------- tests/video/test_h3_flashattn.py | 29 +++++--- vllm/entrypoints/cli/video.py | 4 + .../models/minimax_h3/attention.py | 12 ++- .../models/minimax_h3/config.py | 8 ++ .../models/minimax_h3/pipeline.py | 3 + 7 files changed, 107 insertions(+), 39 deletions(-) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 957dafcf5c..c4ac81af5d 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -71,3 +71,20 @@ original 1280x736/124-frame internal canvas for the five-second sample: Independent official reference, full audiovisual review, other adapters and the complete shape/TP matrix remain required. No AUTO selection is qualified. + +## Explicit attention query geometry + +`attention_query_tile=128` / `--attention-query-tile 128` opts into a 128-query +FlashAttention-V100 CTA. The default remains 64 and retains the previous call +ABI. Both sizes use the same 32x64 warp arithmetic and key-tile selection. +The option applies independently of model weights and adapters; other attention +backends reject this explicit tiling option rather than ignoring it. + +The separate prototype retains exact outputs at nine boundary lengths and +the actual 34,551-token Q/K/V capture. Full four-step video/audio latents also +match frozen mainline bitwise (`q128-profile-quality.json`). A matching pair of +full-denoise profiles is retained in `epilogue-profile-breakdown/` and +`q128-epilogue-profile-breakdown/`; profiler timings are not acceptance results. +The public kernel/CLI implementation additionally passes 69 GPU tail, storage, +cross-attention-length and graph checks. Complete native API quality and three +unprofiled measurements of this explicit option are still pending. diff --git a/flash-attention-v100/kernel/h3/forward.cu b/flash-attention-v100/kernel/h3/forward.cu index acd5f96212..74218bd055 100644 --- a/flash-attention-v100/kernel/h3/forward.cu +++ b/flash-attention-v100/kernel/h3/forward.cu @@ -17,29 +17,30 @@ namespace { using Half = cutlass::half_t; constexpr int kHeadDim = 128; -constexpr int kQueries = 64; -template -using KernelFor = typename cutlass::gemm::kernel::H3FMHA< - Half, cutlass::arch::Sm70, true, kQueries, Keys, kHeadDim>::FMHAKernel; +template +using KernelFor = + typename cutlass::gemm::kernel::H3FMHA::FMHAKernel; -template -__global__ __launch_bounds__(128, 1) void h3_flash_v100_d128( - typename KernelFor::DirectParams params) { +template +__global__ __launch_bounds__(Queries * 2, 1) void h3_flash_v100_d128( + typename KernelFor::DirectParams params) { extern __shared__ __align__(16) unsigned char storage[]; if constexpr (Fixed) { params.heads = 14; params.queries = 12323; params.keys = 12323; } - KernelFor kernel; + KernelFor kernel; kernel(params, - *reinterpret_cast::SharedStorage*>(storage)); + *reinterpret_cast::SharedStorage*>( + storage)); } -template +template void launch_attention(at::Tensor const& q, at::Tensor const& k, at::Tensor const& v, at::Tensor& output, float scale) { - using Kernel = KernelFor; + using Kernel = KernelFor; typename Kernel::DirectParams params{ reinterpret_cast(q.data_ptr()), reinterpret_cast(k.data_ptr()), @@ -50,18 +51,38 @@ void launch_attention(at::Tensor const& q, at::Tensor const& k, int(q.size(2)), scale}; if constexpr (Keys == 128) { - // 34,304 bytes/block: allow two resident blocks without extra global - // storage. + // The 64-query tile uses 34,304 bytes/block. Prefer full shared-memory + // capacity for both query geometries without extra global storage. C10_CUDA_CHECK(cudaFuncSetAttribute( - h3_flash_v100_d128, + h3_flash_v100_d128, cudaFuncAttributePreferredSharedMemoryCarveout, 100)); } - h3_flash_v100_d128 - <<, + cudaFuncAttributeMaxDynamicSharedMemorySize, + sizeof(typename Kernel::SharedStorage))); + } + h3_flash_v100_d128 + <<>>(params); } +template +void dispatch_attention(at::Tensor const& q, at::Tensor const& k, + at::Tensor const& v, at::Tensor& output, float scale, + int selected) { + if (selected == 128) { + if (q.size(1) == 12323 && k.size(1) == 12323 && q.size(2) == 14) + launch_attention(q, k, v, output, scale); + else + launch_attention(q, k, v, output, scale); + } else { + launch_attention(q, k, v, output, scale); + } +} + at::Tensor aligned_contiguous(const at::Tensor& tensor) { auto result = tensor.contiguous(); // contiguous() may preserve a contiguous view with an unaligned offset. @@ -72,7 +93,8 @@ at::Tensor aligned_contiguous(const at::Tensor& tensor) { } // namespace at::Tensor h3_flash_attention_forward(at::Tensor q, at::Tensor k, at::Tensor v, - double scale, int key_tile) { + double scale, int key_tile, + int query_tile) { TORCH_CHECK(q.is_cuda() && q.dim() == 4 && q.scalar_type() == at::kHalf, "H3 FlashAttention-V100 requires CUDA FP16 BSND tensors"); TORCH_CHECK(k.device() == q.device() && v.device() == q.device() && @@ -91,12 +113,14 @@ at::Tensor h3_flash_attention_forward(at::Tensor q, at::Tensor k, at::Tensor v, "H3 FlashAttention-V100 is an inference-only operator"); TORCH_CHECK(key_tile == 0 || key_tile == 64 || key_tile == 128, "H3 attention key tile must be 0, 64 or 128"); + TORCH_CHECK(query_tile == 64 || query_tile == 128, + "H3 attention query tile must be 64 or 128"); const c10::cuda::CUDAGuard guard(q.device()); auto* properties = at::cuda::getCurrentDeviceProperties(); TORCH_CHECK(properties->major == 7 && properties->minor == 0, "H3 FlashAttention-V100 requires SM70"); int64_t groups64 = q.size(0) * q.size(2); - int64_t blocks64 = ((q.size(1) + kQueries - 1) / kQueries) * groups64; + int64_t blocks64 = ((q.size(1) + query_tile - 1) / query_tile) * groups64; TORCH_CHECK(q.size(1) <= INT_MAX && k.size(1) <= INT_MAX && groups64 <= 65535 && blocks64 <= INT_MAX, "H3 attention shape exceeds kernel index limits"); @@ -108,13 +132,10 @@ at::Tensor h3_flash_attention_forward(at::Tensor q, at::Tensor k, at::Tensor v, // self-attention reuses each Q fragment across twice as many keys. int selected = key_tile ? key_tile : (q.size(1) >= 1024 && k.size(1) >= 1024 ? 128 : 64); - if (selected == 128) { - if (q.size(1) == 12323 && k.size(1) == 12323 && q.size(2) == 14) - launch_attention<128, true>(q, k, v, output, float(scale)); - else - launch_attention<128>(q, k, v, output, float(scale)); - } else - launch_attention<64>(q, k, v, output, float(scale)); + if (query_tile == 128) + dispatch_attention<128>(q, k, v, output, float(scale), selected); + else + dispatch_attention<64>(q, k, v, output, float(scale), selected); C10_CUDA_KERNEL_LAUNCH_CHECK(); return output; } @@ -122,5 +143,5 @@ at::Tensor h3_flash_attention_forward(at::Tensor q, at::Tensor k, at::Tensor v, PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.def("forward", &h3_flash_attention_forward, pybind11::arg("q"), pybind11::arg("k"), pybind11::arg("v"), pybind11::arg("scale"), - pybind11::arg("key_tile") = 0); + pybind11::arg("key_tile") = 0, pybind11::arg("query_tile") = 64); } diff --git a/tests/video/test_h3_flashattn.py b/tests/video/test_h3_flashattn.py index b05ad2f142..392eb74a85 100644 --- a/tests/video/test_h3_flashattn.py +++ b/tests/video/test_h3_flashattn.py @@ -26,7 +26,8 @@ def test_flashattn_rejects_batch_head_grid_overflow(): "length", [1, 31, 32, 33, 63, 64, 65, 96, 97, 127, 128, 129, 12323] ) @pytest.mark.parametrize("key_tile", [64, 128]) -def test_flashattn_d128_mha_tails_and_online_rescaling(length, key_tile): +@pytest.mark.parametrize("query_tile", [64, 128]) +def test_flashattn_d128_mha_tails_and_online_rescaling(length, key_tile, query_tile): torch.manual_seed(42) q, k, v = [ torch.randn(2, length, 2, 128, device="cuda", dtype=torch.float16) @@ -37,14 +38,18 @@ def test_flashattn_d128_mha_tails_and_online_rescaling(length, key_tile): k[:, length // 2 :] *= 4 rows = torch.linspace(0, length - 1, min(length, 65), device="cuda").long() expected = chunked_attention_reference(q[:, rows], k, v, scale=128**-0.5) - actual = flashattn_extension().forward(q, k, v, 128**-0.5, key_tile) + actual = flashattn_extension().forward(q, k, v, 128**-0.5, key_tile, query_tile) assert torch.isfinite(actual).all() torch.testing.assert_close(actual[:, rows], expected, atol=0.002, rtol=0.03) + if query_tile == 128: + control = flashattn_extension().forward(q, k, v, 128**-0.5, key_tile, 64) + torch.testing.assert_close(actual, control, atol=0, rtol=0) @pytest.mark.parametrize("layout", ["offset", "strided"]) @pytest.mark.parametrize("key_tile", [64, 128]) -def test_flashattn_d128_storage_and_different_q_k_lengths(layout, key_tile): +@pytest.mark.parametrize("query_tile", [64, 128]) +def test_flashattn_d128_storage_and_different_q_k_lengths(layout, key_tile, query_tile): torch.manual_seed(42) tensors = [] for length, offset in [(33, 1), (65, 3), (65, 5)]: @@ -59,7 +64,7 @@ def test_flashattn_d128_storage_and_different_q_k_lengths(layout, key_tile): tensors.append(value) q, k, v = tensors expected = chunked_attention_reference(q, k, v, scale=128**-0.5) - actual = flashattn_extension().forward(q, k, v, 128**-0.5, key_tile) + actual = flashattn_extension().forward(q, k, v, 128**-0.5, key_tile, query_tile) torch.testing.assert_close(actual, expected, atol=0.002, rtol=0.03) @@ -98,7 +103,8 @@ def forward(self, q, k, v, scale): @pytest.mark.parametrize("key_tile", [64, 128]) -def test_flashattn_graph_replay_uses_new_values(key_tile): +@pytest.mark.parametrize("query_tile", [64, 128]) +def test_flashattn_graph_replay_uses_new_values(key_tile, query_tile): ops = flashattn_extension() q, k, v = [ torch.randn(1, 129, 2, 128, device="cuda", dtype=torch.float16) @@ -108,11 +114,11 @@ def test_flashattn_graph_replay_uses_new_values(key_tile): stream.wait_stream(torch.cuda.current_stream()) with torch.cuda.stream(stream): for _ in range(3): - ops.forward(q, k, v, 128**-0.5, key_tile) + ops.forward(q, k, v, 128**-0.5, key_tile, query_tile) torch.cuda.current_stream().wait_stream(stream) graph = torch.cuda.CUDAGraph() with torch.cuda.graph(graph): - actual = ops.forward(q, k, v, 128**-0.5, key_tile) + actual = ops.forward(q, k, v, 128**-0.5, key_tile, query_tile) for _ in range(2): q.normal_() k.normal_() @@ -122,7 +128,8 @@ def test_flashattn_graph_replay_uses_new_values(key_tile): torch.testing.assert_close(actual, expected, atol=0.002, rtol=0.03) -def test_flashattn_fixed_shape_graph_replay(): +@pytest.mark.parametrize("query_tile", [64, 128]) +def test_flashattn_fixed_shape_graph_replay(query_tile): ops = flashattn_extension() q, k, v = [ torch.randn(1, 12323, 14, 128, device="cuda", dtype=torch.float16) @@ -132,11 +139,11 @@ def test_flashattn_fixed_shape_graph_replay(): stream.wait_stream(torch.cuda.current_stream()) with torch.cuda.stream(stream): for _ in range(3): - ops.forward(q, k, v, 128**-0.5) + ops.forward(q, k, v, 128**-0.5, 0, query_tile) torch.cuda.current_stream().wait_stream(stream) graph = torch.cuda.CUDAGraph() with torch.cuda.graph(graph): - actual = ops.forward(q, k, v, 128**-0.5) + actual = ops.forward(q, k, v, 128**-0.5, 0, query_tile) rows = torch.tensor([0, 31, 32, 63, 64, 127, 128, 12322], device="cuda") for _ in range(2): v.normal_() @@ -155,3 +162,5 @@ def test_flashattn_rejects_changed_head_dimension_and_scale(): ops.forward(q, q, q, float("nan")) with pytest.raises(RuntimeError, match="key tile"): ops.forward(q, q, q, 128**-0.5, 32) + with pytest.raises(RuntimeError, match="query tile"): + ops.forward(q, q, q, 128**-0.5, 64, 32) diff --git a/vllm/entrypoints/cli/video.py b/vllm/entrypoints/cli/video.py index cb5bbe8e96..b66255209b 100644 --- a/vllm/entrypoints/cli/video.py +++ b/vllm/entrypoints/cli/video.py @@ -40,6 +40,9 @@ def subparser_init(self, subparsers): default="FLASH_ATTN_V100", ) mode.add_argument("--fp16-weight-cache-gib", type=float, default=0) + mode.add_argument( + "--attention-query-tile", type=int, choices=(64, 128), default=64 + ) mode.add_argument( "--disable-host-weight-pinning", dest="host_weight_pin_memory", @@ -114,6 +117,7 @@ def cmd(args): transformer_path=args.transformer_path, tensor_parallel_size=args.tensor_parallel_size, attention_backend=args.attention_backend, + attention_query_tile=args.attention_query_tile, fp16_weight_cache_gib=args.fp16_weight_cache_gib, fp16_cache_layers=tuple(args.fp16_cache_layer), lora_path=args.lora_path, diff --git a/vllm/model_executor/models/minimax_h3/attention.py b/vllm/model_executor/models/minimax_h3/attention.py index 4cb91e208e..528491aae3 100644 --- a/vllm/model_executor/models/minimax_h3/attention.py +++ b/vllm/model_executor/models/minimax_h3/attention.py @@ -81,6 +81,7 @@ def __init__( self.backend = attention_backend.get() self.scale = softmax_scale self.head_size = head_size + self.query_tile = 64 @property def attn_backend(self): @@ -107,9 +108,14 @@ def forward(self, q, k, v, metadata): if self.backend == "FLASH_ATTN_V100": from .cuda_ops import flashattn_extension - attended = flashattn_extension().forward( - q_valid, k_valid, v_valid, self.scale - ) + if self.query_tile == 64: + attended = flashattn_extension().forward( + q_valid, k_valid, v_valid, self.scale + ) + else: + attended = flashattn_extension().forward( + q_valid, k_valid, v_valid, self.scale, 0, self.query_tile + ) elif self.backend == "FLASHINFER_SM70": from .cuda_ops import flashinfer_extension diff --git a/vllm/model_executor/models/minimax_h3/config.py b/vllm/model_executor/models/minimax_h3/config.py index 7861ac6da3..7fcde8c0c1 100644 --- a/vllm/model_executor/models/minimax_h3/config.py +++ b/vllm/model_executor/models/minimax_h3/config.py @@ -38,6 +38,7 @@ class H3Config: transformer_path: str | None = None tensor_parallel_size: int = 4 attention_backend: str = "FLASH_ATTN_V100" + attention_query_tile: Literal[64, 128] = 64 fp16_weight_cache_gib: float = 0.0 fp16_cache_layers: tuple[str, ...] = () lora_path: str | None = None @@ -48,6 +49,13 @@ class H3Config: video_encoder: Literal["libx264", "h264_nvenc"] = "libx264" def __post_init__(self) -> None: + if self.attention_query_tile not in (64, 128): + raise H3InputError("Attention query tile must be 64 or 128") + if ( + self.attention_query_tile != 64 + and self.attention_backend != "FLASH_ATTN_V100" + ): + raise H3InputError("Explicit query tiling requires FLASH_ATTN_V100") if not isinstance(self.host_weight_pin_memory, bool): raise H3InputError("host weight pinning must be a boolean") if self.video_encoder not in ("libx264", "h264_nvenc"): diff --git a/vllm/model_executor/models/minimax_h3/pipeline.py b/vllm/model_executor/models/minimax_h3/pipeline.py index b5d96d7be4..b71fa20d49 100644 --- a/vllm/model_executor/models/minimax_h3/pipeline.py +++ b/vllm/model_executor/models/minimax_h3/pipeline.py @@ -497,6 +497,9 @@ def __init__(self, config: H3Config): ) finally: attention_backend.reset(token) + for module in self.transformer.modules(): + if isinstance(module, Attention): + module.query_tile = config.attention_query_tile weights = iter_checkpoint_weights(transformer_path) if restore_adaln: weights = restore_dense_adaln_weights(weights, path / "transformer") From bd5e1898265eb1783fcc413de321125230fbe594 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 14:07:11 +0800 Subject: [PATCH 04/25] [Doc] Record complete query tiling quality and performance results Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 24 ++++++++++++++++++++++-- 1 file changed, 22 insertions(+), 2 deletions(-) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index c4ac81af5d..dc67e8af1b 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -86,5 +86,25 @@ match frozen mainline bitwise (`q128-profile-quality.json`). A matching pair of full-denoise profiles is retained in `epilogue-profile-breakdown/` and `q128-epilogue-profile-breakdown/`; profiler timings are not acceptance results. The public kernel/CLI implementation additionally passes 69 GPU tail, storage, -cross-attention-length and graph checks. Complete native API quality and three -unprofiled measurements of this explicit option are still pending. +cross-attention-length and graph checks; 53 strengthened comparisons also +require exact equality between query geometries and reject invalid query tiles. + +The native option at source `6b39f1c23ca6834e9beead89b4cdf57548101e97` passes +the complete latent/RGB/PCM comparison (`query-tile-quality.json`): both final +latents and all 124 frames match frozen mainline bitwise; SSIM 1 and all audio +gates pass. The same explicit configuration completed one full warmup plus +three unprofiled requests (`query-tile-720p-three-runs/performance.json`): + +| Configuration | Median denoise seconds | Useful TFLOP/s/card | Denoise CV | +| --- | ---: | ---: | ---: | +| Audited original FA baseline | 65.898529 | 47.091839–47.091855 | 0.055668% | +| Prepared/residual/epilogue, query tile 64 | 62.183194 | 49.905491–49.905510 | 0.056387% | +| Same path, explicit query tile 128 | 59.748563 | 51.939038–51.939057 | 0.003793% | + +The last three denoise times are 59.743807 / 59.748563 / 59.748666 seconds. +Complete request times are 81.805001 / 82.780855 / 84.272880 seconds and peak +allocation remains 19,501,498,880 bytes/card. The 128-query setting reduces +median denoise by 3.92% relative to the same prepared 64-query configuration, +and the combined change reduces it by 9.33% relative to the original baseline. +These measurements cover one five-second workflow only. **The >80 gate still +fails; official reference, human review and the wider matrix remain pending.** From 4ccd45428e33749e1becb1b85912a86187e879e9 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 15:55:13 +0800 Subject: [PATCH 05/25] [Doc] Record eight-step H3 numerical preservation Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 28 ++++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index dc67e8af1b..8725af4da9 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -108,3 +108,31 @@ median denoise by 3.92% relative to the same prepared 64-query configuration, and the combined change reduces it by 9.33% relative to the original baseline. These measurements cover one five-second workflow only. **The >80 gate still fails; official reference, human review and the wider matrix remain pending.** + +## Eight-step numerical preservation + +The same explicit query-128/prepared/residual/epilogue configuration at +`bd5e1898265eb1783fcc413de321125230fbe594` also completed a matched TP4 +LightX2V eight-step FL2V v1.0_768p comparison with frozen mainline `4f19ef7`. +Both runs use W8A16, seed 42, the same five-second request, nine sigma points, +flow shift 6 and audio flow shift 3. The official adapter SHA256 is +`9b0efe3613b43a84e30febaa43af27432ea9d0711eac7bba904b2556b175f6d4`. + +`light8-720p-quality.json` passes every declared numerical gate: both final +latents, all 124 RGB frames and decoded PCM match bitwise. This extends +preservation evidence beyond the four-step adapter, using the same shared +operators without an adapter-specific dispatch exception. It remains a frozen +native control, not independent official-model acceptance. + +The captured cold requests took 141.160054 and 121.541891 seconds in denoise; +complete request times were 174.840862 and 203.651622 seconds respectively. +These single captured runs have different staging conditions and no full +warmup, so they do not establish a formal performance result or an overall +request speedup. Source hashes, binary manifests and commands are recorded +in `light8-pair.json`. The eight-step >80 gate remains unmeasured. + +An additional 16-row warp experiment (`attention-warp16/hypothesis.json`) +was rejected at compilation: CUTLASS Volta MMA requires a multiple of its +interleaved tile shape. No GPU run or production change followed. Supporting +that geometry requires new MMA and accumulator iterators, not another +configuration-only benchmark of the rejected shape. From 9d2489fc4e2f32ea500ec13478bd68ef9000c1cc Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 16:28:53 +0800 Subject: [PATCH 06/25] [Doc] Record exact mixed-reference H3 eight-step control Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 34 ++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 8725af4da9..0b8686857e 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -136,3 +136,37 @@ was rejected at compilation: CUTLASS Volta MMA requires a multiple of its interleaved tile shape. No GPU run or production change followed. Supporting that geometry requires new MMA and accumulator iterators, not another configuration-only benchmark of the rejected shape. + +## Mixed-reference eight-step control + +Source `2063b09f2d75c4a63af3f90a3b7803744ffc6e02` completes Ref2VA with the +official eight-step v1.0_768p adapter, W8A16 Ref2VA base, seed 42 and one image, +one 2.5-second reference video plus one standalone audio reference. The video +start time is zero. The output remains 1280x736/124 internal frames for the +five-second request. This control uses 69,325 valid DiT tokens and a 10,273-token +Qwen presentation, exercising mixed reference indices and padded residual rows. + +`ref8-mixed-quality.json` passes all gates against frozen native mainline: +video/audio latents, all RGB frames and PCM are bitwise equal; PSNR infinity, +SSIM 1, RMS ratio 1. Candidate settings include prepared execution, exact +residual sharding, query tile 128 and explicit shared pageable VAE host weights. +The frozen control uses its ordinary query-64 path and pinned host masters. +Host residency changes byte ownership/transfers, not GPU arithmetic. + +Single captured cold denoise times are 421.173489 seconds for the frozen +control and 366.799932 seconds for the candidate; request times are +586.794155 and 454.461091 seconds. Peak GPU allocation is 19,623,684,096 and +19,615,279,104 bytes/card respectively. The candidate reports corrected useful +throughput 52.919225–52.919231 TFLOP/s/card. These runs lack full warmup and +three measurements and have different host staging policies, so they do not +qualify either performance or attribution to an individual optimization. +Independent official quality, audiovisual/reference review and other reference +combinations remain pending. Raw contracts and media are retained in +`/home/ymzx/h3-sm70-artifacts-20260909/runs/ref8-mixed-720p-{baseline,candidate}/`. + +Two additional CTA-barrier coalescing candidates preserve bitwise edge and +34,551-token results. The second passes 20 synccheck and racecheck geometries +with zero errors/hazards. Their paired operator gains are only 0.3–0.6%, so +neither is retained or promoted to a full-request performance claim. Evidence: +`attention-barrier-coalesce/`, `attention-barrier-coalesce-v2/` and the +`barrier-v2-*.log` files under the campaign root. From e8d0185b0573ca1b84651c729ebd08d9a1ab4d8d Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 17:03:45 +0800 Subject: [PATCH 07/25] [Doc] Record full H3 243-frame and 15-second compatibility Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 25 ++++++++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 0b8686857e..76ed4b97d2 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -170,3 +170,28 @@ with zero errors/hazards. Their paired operator gains are only 0.3–0.6%, so neither is retained or promoted to a full-request performance claim. Evidence: `attention-barrier-coalesce/`, `attention-barrier-coalesce-v2/` and the `barrier-v2-*.log` files under the campaign root. + +## Larger canvas and duration compatibility + +Source `9d2489fc4e2f32ea500ec13478bd68ef9000c1cc` completes two additional +TP4 W8A16 LightX2V four-step requests with prepared execution, exact residual +sharding, query tile 128 and shared pageable VAE masters. Both use seed 42, +the original paper-boat prompt, five sigma points and flow shifts 6/3. + +| Requested shape | Actual frames | Denoise seconds | Request seconds | Peak GPU allocation bytes/card | Useful TFLOP/s/card | +| --- | ---: | ---: | ---: | ---: | ---: | +| 1344x768, 243 frames | 243 | 203.431920 | 254.138019 | 24,341,115,904 | 52.603224–52.603229 | +| 1344x768, 15 seconds | 362 | 401.026288 | 469.032456 | 28,777,495,040 | 53.561244–53.561247 | + +The 15-second request resolves to the model's 362-frame aligned output; it is +not claimed to be an exactly 15.000-second encoded clip. Both complete native +media validation and strict actual-work checks, and remove their owned shared +weight directories after shutdown. Full captures, source/binary manifests, +per-rank stages and NVML samples are retained under +`/home/ymzx/h3-sm70-artifacts-20260909/runs/official-243-frame/` and +`boundary-15-second/`; `large-canvas-summary.json` summarizes the evidence. + +These first captured requests establish shape and memory compatibility only. +They have no matched quality reference, full warmup or three post-warmup +measurements. Both are below 80 and remain unqualified. The independent +official reference and human review gates are also pending. From be89a26d1c813950836f1f2b11c123667cd8931e Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 17:52:49 +0800 Subject: [PATCH 08/25] [Core] Expose shared SM70 noncausal attention operators Both H3 and other DiT callers use explicit native FP16 attention without model-specific dispatch gates. Twelve GPU interface, non-H3 shape, padding and graph checks pass; 18 affected CPU checks pass. Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SHARED_ATTENTION.md | 47 ++++++++++ tests/video/test_h3_flashattn.py | 2 +- tests/video/test_sm70_attention.py | 92 +++++++++++++++++++ vllm/model_executor/layers/sm70_attention.py | 86 +++++++++++++++++ .../models/minimax_h3/attention.py | 26 ++---- .../models/minimax_h3/cuda_ops.py | 66 ++----------- 6 files changed, 246 insertions(+), 73 deletions(-) create mode 100644 docs/design/minimax_h3/SHARED_ATTENTION.md create mode 100644 tests/video/test_sm70_attention.py create mode 100644 vllm/model_executor/layers/sm70_attention.py diff --git a/docs/design/minimax_h3/SHARED_ATTENTION.md b/docs/design/minimax_h3/SHARED_ATTENTION.md new file mode 100644 index 0000000000..281fd34c71 --- /dev/null +++ b/docs/design/minimax_h3/SHARED_ATTENTION.md @@ -0,0 +1,47 @@ +# Shared SM70 attention interface + +`vllm.model_executor.layers.sm70_attention.noncausal_attention` exposes both +native SM70 dense attention implementations without importing the H3 model +package. H3's dense facade uses the same function. Historical extension ABI +names and H3 loader exports remain compatible with existing wheels; a default +call still invokes the original four-argument entrypoint. + +```python +from vllm.model_executor.layers.sm70_attention import noncausal_attention + +output = noncausal_attention( + q, k, v, scale=128**-0.5, backend="FLASH_ATTN_V100", query_tile=128 +) +``` + +Inputs are CUDA FP16 BSHD MHA tensors with head dimension 128; accumulation and +softmax remain FP32. No model name, quantization label or adapter identity is +required. FlashAttention supports different Q/K lengths, strided storage and +explicit query/key tiles. FlashInfer requires matching lengths and receives +contiguous inputs. Invalid scales, including values overflowing FP32, fail +before native loading. Unsupported dtype, hardware and shapes fail at the +CUDA entrypoint. There is no silent BF16/FP32 conversion or backend substitution. + +Callers own masks, suffix padding and sparse geometry. The H3 facade still +slices valid tokens before dispatch and restores zero suffix padding; VSA +keeps its separate prefix, selected-block and learned-gate implementation. +AUTO qualification is a separate unfinished campaign requirement. + +## Validation + +Environment: Python 3.12.13, Torch 2.10.0+cu128, CUDA 12.8.93, V100 SXM2 32GB. +No CUDA source or binary changes accompany this interface extraction. + +- `shared-attention-gpu.log`: 12 GPU checks pass. Both native backends and both + FA query geometries preserve direct-entrypoint results bitwise for BSHD + shapes `(1, 1537, 24, 128)` and `(2, 65, 8, 128)`, including strided inputs + and a later increase in the online maximum. Sampled FP32 relative L2 is + below 0.001. Additional checks cover unequal Q/K lengths, dtype rejection, + H3 poisoned suffix padding and CUDA Graph replay with new input values. +- `shared-attention-cpu.log`: 18 scale, deployment, service and residency + checks pass; nine GPU cases are deselected. + +These tests exercise non-H3 DiT operator shapes, not a second complete model. +They establish interface preservation, not a new end-to-end speed result or +independent official quality acceptance. Raw logs and binary hashes are under +`/data/minimax-h3/sm70-general-20260909/`. diff --git a/tests/video/test_h3_flashattn.py b/tests/video/test_h3_flashattn.py index 392eb74a85..6bdeb9ae5f 100644 --- a/tests/video/test_h3_flashattn.py +++ b/tests/video/test_h3_flashattn.py @@ -69,7 +69,7 @@ def test_flashattn_d128_storage_and_different_q_k_lengths(layout, key_tile, quer def test_flashattn_dispatch_slices_poisoned_padding(monkeypatch): - from vllm.model_executor.models.minimax_h3 import cuda_ops + from vllm.model_executor.layers import sm70_attention as cuda_ops native = flashattn_extension() calls = [] diff --git a/tests/video/test_sm70_attention.py b/tests/video/test_sm70_attention.py new file mode 100644 index 0000000000..cedad0e16f --- /dev/null +++ b/tests/video/test_sm70_attention.py @@ -0,0 +1,92 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Shared attention contracts and non-H3 DiT shapes, without model imports.""" + +import pytest +import torch + +from vllm.model_executor.layers import sm70_attention as ops + + +@pytest.mark.parametrize("scale", [0, -1, float("inf"), float("nan"), 1e100]) +def test_invalid_scale_does_not_load_native_code(monkeypatch, scale): + def unexpected_load(): + pytest.fail("invalid scale must fail before extension loading") + + monkeypatch.setattr(ops, "flashattn_extension", unexpected_load) + monkeypatch.setattr(ops, "flashinfer_extension", unexpected_load) + for backend in ("FLASH_ATTN_V100", "FLASHINFER_SM70"): + with pytest.raises(ValueError, match="scale"): + ops.noncausal_attention(None, None, None, scale=scale, backend=backend) + + +def test_explicit_backend_and_geometry_contract(): + for backend, tile in (("AUTO", 64), ("FLASHINFER_SM70", 128)): + with pytest.raises(ValueError): + ops.noncausal_attention( + None, None, None, scale=0.1, backend=backend, query_tile=tile + ) + + +@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires SM70 GPU") +@pytest.mark.parametrize( + "backend,tile", + [("FLASH_ATTN_V100", 64), ("FLASH_ATTN_V100", 128), ("FLASHINFER_SM70", 64)], +) +@pytest.mark.parametrize("batch,length,heads", [(1, 1537, 24), (2, 65, 8)]) +def test_gpu_non_h3_shapes_preserve_native_output(backend, tile, batch, length, heads): + torch.manual_seed(206) + torch.backends.cuda.matmul.allow_tf32 = False + torch.backends.cuda.matmul.allow_fp16_reduced_precision_reduction = False + # Strided token storage and non-H3 head counts exercise the shared contract. + q, k, v = [ + torch.randn(batch, length * 2, heads, 128, device="cuda", dtype=torch.float16)[ + :, ::2 + ] + for _ in range(3) + ] + k[:, length // 2 :] *= 4 + scale = 128**-0.5 + actual = ops.noncausal_attention( + q, k, v, scale=scale, backend=backend, query_tile=tile + ) + native = ( + ops.flashattn_extension() + if backend == "FLASH_ATTN_V100" + else ops.flashinfer_extension() + ) + args = (scale, 0, tile) if tile != 64 else (scale,) + expected = native.forward(q.contiguous(), k.contiguous(), v.contiguous(), *args) + torch.testing.assert_close(actual, expected, atol=0, rtol=0) + rows = torch.linspace(0, length - 1, min(length, 33), device="cuda").long() + qh, kh, vh = (x.transpose(1, 2).float() for x in (q[:, rows], k, v)) + reference = (((qh @ kh.transpose(-1, -2)) * scale).softmax(-1) @ vh).transpose(1, 2) + assert torch.isfinite(actual).all() + relative_l2 = (actual[:, rows].float() - reference).norm() / reference.norm() + assert relative_l2 < 0.001 + with pytest.raises(RuntimeError, match="FP16"): + ops.noncausal_attention( + q.float(), + k.float(), + v.float(), + scale=scale, + backend=backend, + query_tile=tile, + ) + + +@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires SM70 GPU") +def test_gpu_cross_attention_contract(): + torch.manual_seed(207) + q = torch.randn(2, 33, 8, 128, device="cuda", dtype=torch.float16) + k, v = [ + torch.randn(2, 129, 8, 128, device="cuda", dtype=torch.float16) + for _ in range(2) + ] + actual = ops.noncausal_attention( + q, k, v, scale=0.1, backend="FLASH_ATTN_V100", query_tile=128 + ) + expected = ops.flashattn_extension().forward(q, k, v, 0.1, 0, 128) + torch.testing.assert_close(actual, expected, atol=0, rtol=0) + with pytest.raises(RuntimeError, match="matching"): + ops.noncausal_attention(q, k, v, scale=0.1, backend="FLASHINFER_SM70") diff --git a/vllm/model_executor/layers/sm70_attention.py b/vllm/model_executor/layers/sm70_attention.py new file mode 100644 index 0000000000..ecf654d838 --- /dev/null +++ b/vllm/model_executor/layers/sm70_attention.py @@ -0,0 +1,86 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Explicit model-independent SM70 FP16 non-causal D128 attention. + +The extension ABI retains historical H3 names. Inputs are BSHD; no model, +quantization, adapter or sampler identity participates in dispatch. +""" + +import math +import os +from functools import lru_cache +from importlib import import_module +from pathlib import Path + + +@lru_cache(maxsize=1) +def flashinfer_extension(): + try: + return import_module("vllm._h3_flashinfer_C") + except ImportError: + pass + from torch.utils.cpp_extension import load + + root = Path(__file__).resolve().parents[3] + return load( + name="onecat_h3_flashinfer_sm70", + sources=[str(root / "flashinfer-sm70/csrc/h3_noncausal_sm70.cu")], + extra_include_paths=[str(root / "flashinfer-sm70/include")], + extra_cuda_cflags=["-O3", "-gencode=arch=compute_70,code=sm_70"], + verbose=False, + ) + + +@lru_cache(maxsize=1) +def flashattn_extension(): + try: + return import_module("vllm._h3_flashattn_C") + except ImportError: + pass + from torch.utils.cpp_extension import load + + root = Path(__file__).resolve().parents[3] + cutlass_path = os.environ.get("VLLM_CUTLASS_SRC_DIR") + if not cutlass_path: + raise RuntimeError( + "Build the H3 FlashAttention-V100 extension (_h3_flashattn_C), or " + "set VLLM_CUTLASS_SRC_DIR to CUTLASS v4.4.2 for source development" + ) + cutlass = Path(cutlass_path) + return load( + name="onecat_h3_flashattn_sm70", + sources=[str(root / "flash-attention-v100/kernel/h3/forward.cu")], + extra_include_paths=[ + str(cutlass / "include"), + str(cutlass / "examples/41_fused_multi_head_attention"), + ], + extra_cuda_cflags=["-O3", "-gencode=arch=compute_70,code=sm_70"], + verbose=False, + ) + + +def noncausal_attention(q, k, v, *, scale, backend, query_tile=64, key_tile=0): + """Run an explicitly selected SM70 implementation without changing precision. + + CUDA entrypoints validate device, FP16 dtype, D128 heads, index limits and + shapes. FlashAttention supports unequal Q/K lengths and strided inputs; + FlashInfer supports matching lengths and receives contiguous BSHD tensors. + Padding/masking and model-specific sparse selection belong to callers. + The default four-argument call remains compatible with existing wheels. + """ + if not math.isfinite(scale) or not 0 < scale <= 3.4028234663852886e38: + raise ValueError("Attention scale must be positive and finite in FP32") + if backend == "FLASH_ATTN_V100": + if query_tile not in (64, 128) or key_tile not in (0, 64, 128): + raise ValueError("Unsupported SM70 attention tile") + ops = flashattn_extension() + if query_tile == 64 and key_tile == 0: + return ops.forward(q, k, v, scale) + return ops.forward(q, k, v, scale, key_tile, query_tile) + if backend == "FLASHINFER_SM70": + if query_tile != 64 or key_tile != 0: + raise ValueError("Explicit attention tiles require FLASH_ATTN_V100") + return flashinfer_extension().forward( + q.contiguous(), k.contiguous(), v.contiguous(), scale + ) + raise ValueError(f"Unsupported SM70 attention backend: {backend}") diff --git a/vllm/model_executor/models/minimax_h3/attention.py b/vllm/model_executor/models/minimax_h3/attention.py index 528491aae3..5469cb3b44 100644 --- a/vllm/model_executor/models/minimax_h3/attention.py +++ b/vllm/model_executor/models/minimax_h3/attention.py @@ -105,22 +105,16 @@ def forward(self, q, k, v, metadata): if not 0 < used <= q.shape[1] or k.shape != v.shape: raise ValueError("invalid packed H3 attention lengths") q_valid, k_valid, v_valid = (x[:, :used].contiguous() for x in (q, k, v)) - if self.backend == "FLASH_ATTN_V100": - from .cuda_ops import flashattn_extension - - if self.query_tile == 64: - attended = flashattn_extension().forward( - q_valid, k_valid, v_valid, self.scale - ) - else: - attended = flashattn_extension().forward( - q_valid, k_valid, v_valid, self.scale, 0, self.query_tile - ) - elif self.backend == "FLASHINFER_SM70": - from .cuda_ops import flashinfer_extension - - attended = flashinfer_extension().forward( - q_valid, k_valid, v_valid, self.scale + if self.backend in ("FLASH_ATTN_V100", "FLASHINFER_SM70"): + from vllm.model_executor.layers.sm70_attention import noncausal_attention + + attended = noncausal_attention( + q_valid, + k_valid, + v_valid, + scale=self.scale, + backend=self.backend, + query_tile=self.query_tile, ) elif self.backend == "TORCH_SDPA": attended = chunked_attention_reference( diff --git a/vllm/model_executor/models/minimax_h3/cuda_ops.py b/vllm/model_executor/models/minimax_h3/cuda_ops.py index 0f9b6cc69e..c99e8dcbac 100644 --- a/vllm/model_executor/models/minimax_h3/cuda_ops.py +++ b/vllm/model_executor/models/minimax_h3/cuda_ops.py @@ -1,67 +1,21 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Lazy build of native H3 SM70 extensions for source development.""" - -import os -from functools import lru_cache -from pathlib import Path +"""Compatibility exports for model-independent SM70 diffusion operators.""" +from vllm.model_executor.layers.sm70_attention import ( + flashattn_extension as flashattn_extension, +) +from vllm.model_executor.layers.sm70_attention import ( + flashinfer_extension as flashinfer_extension, +) from vllm.model_executor.layers.sm70_diffusion import ( _column_major_plan as _column_major_plan, ) from vllm.model_executor.layers.sm70_diffusion import ( fp16_gemm as fp16_gemm, ) -from vllm.model_executor.layers.sm70_diffusion import sm70_extension +from vllm.model_executor.layers.sm70_diffusion import ( + sm70_extension, +) w8a16_extension = sm70_extension - - -@lru_cache(maxsize=1) -def flashinfer_extension(): - try: - from vllm import _h3_flashinfer_C - - return _h3_flashinfer_C - except ImportError: - pass - from torch.utils.cpp_extension import load - - root = Path(__file__).resolve().parents[4] - return load( - name="onecat_h3_flashinfer_sm70", - sources=[str(root / "flashinfer-sm70/csrc/h3_noncausal_sm70.cu")], - extra_include_paths=[str(root / "flashinfer-sm70/include")], - extra_cuda_cflags=["-O3", "-gencode=arch=compute_70,code=sm_70"], - verbose=False, - ) - - -@lru_cache(maxsize=1) -def flashattn_extension(): - try: - from vllm import _h3_flashattn_C - - return _h3_flashattn_C - except ImportError: - pass - from torch.utils.cpp_extension import load - - root = Path(__file__).resolve().parents[4] - cutlass_path = os.environ.get("VLLM_CUTLASS_SRC_DIR") - if not cutlass_path: - raise RuntimeError( - "Build the H3 FlashAttention-V100 extension (_h3_flashattn_C), or " - "set VLLM_CUTLASS_SRC_DIR to CUTLASS v4.4.2 for source development" - ) - cutlass = Path(cutlass_path) - return load( - name="onecat_h3_flashattn_sm70", - sources=[str(root / "flash-attention-v100/kernel/h3/forward.cu")], - extra_include_paths=[ - str(cutlass / "include"), - str(cutlass / "examples/41_fused_multi_head_attention"), - ], - extra_cuda_cflags=["-O3", "-gencode=arch=compute_70,code=sm_70"], - verbose=False, - ) From 3217459df708b99953d1ec17818cc3f1968ab991 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 18:42:47 +0800 Subject: [PATCH 09/25] [Doc] Record full quality controls for all FL2V Turbo artifacts Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 34 ++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 125bff0391..10c6d89577 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -210,3 +210,37 @@ These first captured requests establish shape and memory compatibility only. They have no matched quality reference, full warmup or three post-warmup measurements. Both are below 80 and remain unqualified. The independent official reference and human review gates are also pending. + +## Remaining FL2V Turbo versions + +Source `be89a26d1c` completes four more matched frozen-mainline comparisons on +TP4 at the same 1280x736/124-frame internal canvas for a five-second request. +All use W8A16, seed 42, pageable host masters and no fixed FP16 weight cache. +The candidate uses prepared execution, exact residual sharding, shared LoRA +epilogues and query tile 128. The frozen `4f19ef7` control uses its ordinary +query-64 path. Each artifact retains its official alpha and flow shift. + +| Official artifact | Intervals / sigma points | Video shift / alpha | Candidate denoise seconds | Request seconds | +| --- | ---: | ---: | ---: | ---: | +| FL2V four-step v1.0_768p | 4 / 5 | 6 / 128 | 61.263482 | 96.285706 | +| FL2V four-step v1.1_768p | 4 / 5 | 6 / 128 | 61.018722 | 96.491105 | +| FL2V four-step v0.1 | 4 / 5 | 12 / 8 | 61.000133 | 94.938178 | +| FL2V eight-step v1.0 (non-768p) | 8 / 9 | 12 / 8 | 120.478234 | 156.924670 | + +For every pair, final video/audio latents, all 124 pre-encoding RGB frames and +decoded PCM match bitwise. SSIM and RMS ratio are 1. Strict per-rank workload +validators pass, including actual intervals, block counts and duplicate-work +exclusion. Peak allocation is 19,501,498,880 bytes/card for the four-step cases +and 19,502,023,168 bytes/card for eight-step. Single-request useful throughput +is approximately 50.65–51.52 TFLOP/s/card, below 80. + +These are captured cold quality controls, not warmed three-run performance. +Together with the existing v1.2 four-step and v1.0_768p eight-step controls, +all six official FL2V Turbo artifacts now have a complete W8A16 T2VA numerical +preservation result. Original-weight and keyframe combinations remain separate +pending coverage; Ref2V four-step is also still awaiting its full control. +No independent official quality or human acceptance is inferred. + +Evidence: `remaining-turbo-pairs.json`, `remaining-turbo-summary.json`, +`light4-v{10,11,01}-720p-quality.json`, `light8-v10-non768-720p-quality.json` +and the corresponding captured runs under the campaign's artifact root. From f8b85c681a402ad26aac13adc701687a342881d9 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 19:05:59 +0800 Subject: [PATCH 10/25] [Kernel] Keep SM70 attention probabilities in registers Preserve FP32 accumulation and the original K64 online reduction order while eliminating shared probability traffic and cross-warp softmax synchronization. Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- .../FLASHINFER_REGISTER_PROBABILITY.md | 76 +++++++ flashinfer-sm70/csrc/h3_noncausal_sm70.cu | 186 +++++++++--------- tests/video/test_h3_numerics.py | 8 +- 3 files changed, 171 insertions(+), 99 deletions(-) create mode 100644 docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md diff --git a/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md b/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md new file mode 100644 index 0000000000..dd16c95348 --- /dev/null +++ b/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md @@ -0,0 +1,76 @@ +# FlashInfer SM70 probability fragments in registers + +This change improves the explicitly selected FlashInfer backend. It does not +qualify an AUTO configuration or meet the campaign's >80 useful TFLOP/s/card +gate. The shared `sm70_attention.noncausal_attention` interface dispatches on +SM70 hardware, FP16 tensors and BSHD D128 layout, independently of model, +quantization and adapter identity. + +## Arithmetic and execution + +One warp owns 16 query rows and both logical 32-key halves of the existing +64-key tile. QK/PV accumulators, online maxima and denominators remain FP32. +Each original partial sum, XOR-2/XOR-8 reduction, left/right addition and +per-output MMA K order is retained. Probabilities cross the same FP16 rounding +boundary as before. + +Six 32-bit lane exchanges convert the rounded probability pairs into Volta A +fragments. PV reuses each fragment across eight output fragments. Probabilities +no longer traverse shared memory, and the warp owns its softmax state without +cross-warp barriers. CTA barriers still protect K/V staging and reuse. + +The CTA covers 192 queries with 384 threads. Cooperative prefetch rounds up +to three vectors per thread. Extra complete warps in the final prefetch group +skip rows outside the 64-key tile; the guard is warp-uniform before shuffles. +Valid-length padding, batches, heads and unaligned global-storage loads retain +their original handling. Explicit lane selects avoid addressable local arrays. + +CUDA 12.8 emits 168 registers/thread, no stack or spills, and 84,992 dynamic +shared bytes. These resource counts describe the implementation; they are not +throughput measurements. + +The FP16-accumulator layout shortcut in [FastAttention Appendix B](https://arxiv.org/html/2410.16663v1#A2) +is outside this campaign's precision contract. This implementation exchanges +already rounded probabilities while retaining FP32 accumulation. + +## Development evidence + +Environment: Python 3.12.13, Torch 2.10.0+cu128, CUDA toolkit 12.8.93, +V100 SXM2 32GB. Artifacts are under +`/data/minimax-h3/sm70-general-20260909/`. + +- `attention-fi-register-probability-q192/probe.json`: 17 boundary lengths + and the actual 34,551-token, 14-head H3 capture match the frozen FI binary + bitwise. Sampled FP32-reference relative L2 is 0.000374 on the actual input. +- Seven alternating operator timings on the same GPU give medians + 183.124985 -> 160.701447 ms, 12.245% lower latency. Observed clocks remain + 1425-1432 MHz. This is an operator measurement only. +- Compute Sanitizer 12.8 memcheck, racecheck and synccheck each report zero + errors on 12 boundary cases with two batches, three heads and storage offsets. +- The formatted native build passes 33 GPU numerical, unaligned-storage, + independent-query, graph and non-H3 shared-interface checks. The affected + CPU contracts and provenance/acceptance suite pass 46 checks, with seven + GPU checks skipped in the explicitly device-masked CPU invocation. +- `fi-general-control-quality.json`: the current common prepared/residual + path preserves the previously frozen FI final video/audio latents bitwise. +- `fi-register-denoise-summary.json`: one complete denoise warmup per + implementation, followed by one unprofiled measurement per implementation + in reverse order. Both use the same TP4 W8A16 LightX2V four-step v1.2 weights, + five sigma points, captured conditioning/noise, column layout, exact residual + sharding and zero persistent FP16 weight cache. All four final-latent pairs + match bitwise on every rank. + +| Complete-denoise control | Slowest-rank seconds | Useful TFLOP/s/card | +| --- | ---: | ---: | +| Previous FI kernel, common execution path | 67.303216 | 46.108983-46.109000 | +| Register-probability FI kernel, same path | 62.471266 | 49.675363-49.675382 | + +The isolated kernel change reduces this complete denoise by 7.179%. Candidate +steps take about 15.59-15.64 seconds. The control excludes encoder, VAE and +packaging, and includes only one measurement per implementation. It is not +the required full-request warmup-plus-three acceptance. Existing FA query-128 +remains faster in its separately recorded full-request measurements. + +Full native media preservation, formal repeated request measurements, +independent official references and human review remain separate gates. The +campaign remains incomplete. No precision gate or sampler setting is relaxed. diff --git a/flashinfer-sm70/csrc/h3_noncausal_sm70.cu b/flashinfer-sm70/csrc/h3_noncausal_sm70.cu index 00fffc314f..d086c3eabd 100644 --- a/flashinfer-sm70/csrc/h3_noncausal_sm70.cu +++ b/flashinfer-sm70/csrc/h3_noncausal_sm70.cu @@ -9,28 +9,20 @@ namespace fi = flashinfer::attention::sm70; namespace { constexpr int D = 128; -constexpr int BQ = 128; +constexpr int BQ = 192; constexpr int BK = 64; -// Two warps share 16 query rows. Each owns two K16 score fragments, so -// expanding BK does not double the CTA's thread count. -constexpr int KEY_WARPS = 2; +// One warp owns Q16 and both logical K32 halves, retaining FP32 arithmetic. +constexpr int KEY_WARPS = 1; constexpr int KEY_FRAGMENTS = BK / (KEY_WARPS * 16); constexpr int OUTPUT_FRAGMENTS = D / (KEY_WARPS * 16); constexpr int VLD = BK + 8; constexpr int THREADS = (BQ / 16) * KEY_WARPS * 32; -constexpr int PREFETCH_VECTORS = BK * D / (THREADS * 8); -static_assert(THREADS * PREFETCH_VECTORS * 8 == BK * D); -static_assert(BQ / 16 < 16); // Barrier 0 joins the CTA; 1..8 join query pairs. +constexpr int PREFETCH_VECTORS = (BK * D + THREADS * 8 - 1) / (THREADS * 8); +static_assert(THREADS * PREFETCH_VECTORS * 8 >= BK * D); +static_assert(KEY_WARPS == 1 && BK == 64); constexpr int shared_bytes() { - constexpr int QLD = D + 8, PLD = BK + 4; - return (BQ * D + BK * QLD + D * VLD + BQ * PLD) * 2 + - (BQ * KEY_WARPS * 2 + BQ * 2) * 4; -} - -// QK maxima and probability rows are consumed only by the matching pair. -// Keep CTA-wide barriers around K/V staging and tile consumption. -__device__ __forceinline__ void sync_query_pair(int query_group) { - asm volatile("bar.sync %0, 64;" ::"r"(query_group + 1) : "memory"); + constexpr int QLD = D + 8; + return (BQ * D + BK * QLD + D * VLD) * 2; } __device__ __forceinline__ int q_swizzle(int row) { @@ -49,19 +41,28 @@ __device__ __forceinline__ void load_q_fragment(fi::AFragment& fragment, values[1] = *reinterpret_cast(base + ((col + 8) ^ mask)); } -// A 68-half P stride makes accumulator pair stores conflict-free. Odd rows -// remain 8-byte aligned, so use 64-bit loads rather than WMMA's 128-bit loads. -__device__ __forceinline__ void load_p_fragment(fi::AFragment& fragment, - const half* source, int row, - int col) { +// Convert FP16-rounded probability accumulator pairs to the existing Volta +// A fragment layout. Exchange row ownership across lane bit 1, then concatenate +// the two eight-column halves across lane bit 3. No arithmetic in the exchange. +__device__ __forceinline__ void load_probability_fragment( + fi::AFragment& fragment, const unsigned* pairs) { const int lane = threadIdx.x & 31; - const int physical_row = - row + (lane & 3) + ((lane & 16) >> 2) + ((lane & 4) << 1); - const half* base = source + physical_row * (BK + 4) + col; - auto* values = reinterpret_cast(fragment.x); + const int row_bit = (lane >> 1) & 1; + const unsigned own0 = row_bit ? pairs[1] : pairs[0]; + const unsigned own1 = row_bit ? pairs[3] : pairs[2]; + const unsigned other0 = + __shfl_xor_sync(0xffffffff, row_bit ? pairs[0] : pairs[1], 2); + const unsigned other1 = + __shfl_xor_sync(0xffffffff, row_bit ? pairs[2] : pairs[3], 2); + const unsigned local[4] = {row_bit ? other0 : own0, row_bit ? own0 : other0, + row_bit ? other1 : own1, row_bit ? own1 : other1}; + auto* output = reinterpret_cast(fragment.x); #pragma unroll - for (int i = 0; i < 4; ++i) - values[i] = *reinterpret_cast(base + i * 4); + for (int part = 0; part < 4; ++part) { + const unsigned opposite = __shfl_xor_sync(0xffffffff, local[part], 8); + output[part] = (lane & 8) ? opposite : local[part]; + output[part + 4] = (lane & 8) ? local[part] : opposite; + } } __global__ __launch_bounds__(THREADS, @@ -69,15 +70,13 @@ __global__ __launch_bounds__(THREADS, const half* v, half* output, int length, int heads, float scale) { - constexpr int QLD = D + 8, PLD = BK + 4; + constexpr int QLD = D + 8; extern __shared__ __align__(32) unsigned char raw[]; half* qs = reinterpret_cast(raw); half* ks = qs + BQ * D; half* vs = ks + BK * QLD; - half* probabilities = vs + D * VLD; - float* scores = reinterpret_cast(probabilities + BQ * PLD); - float* maximum = scores + BQ * KEY_WARPS * 2; - float* denominator = maximum + BQ; + float running_max[2] = {-INFINITY, -INFINITY}; + float running_sum[2] = {0.f, 0.f}; const int tid = threadIdx.x, warp = tid / 32; const int warp_q = warp / KEY_WARPS, warp_k = warp % KEY_WARPS; const int lane = tid % 32; @@ -101,10 +100,6 @@ __global__ __launch_bounds__(THREADS, row < length ? q[base + int64_t(row) * heads * D + i % D] : __float2half(0.f); } - if (tid < BQ) { - maximum[tid] = -INFINITY; - denominator[tid] = 0.f; - } __syncthreads(); for (int start = 0; start < length; start += BK) { if (start == 0) { @@ -135,76 +130,73 @@ __global__ __launch_bounds__(THREADS, fi::mma_sync_m16n16k16_row_col_f16f16f32(qk[n], qa, kb); } } + unsigned probability_pairs[KEY_FRAGMENTS][4]; { - // Volta distributes each accumulator row across lanes differing in - // bits 1 and 3. Reduce its 16 columns in registers, then combine only - // the two warp partials through shared memory. - float row_max[2] = {-INFINITY, -INFINITY}; + // Preserve the original two logical K32 partials and their FP32 sum + // order even though one warp now owns both halves of this K64 tile. + float row_max[2][2] = {{-INFINITY, -INFINITY}, {-INFINITY, -INFINITY}}; #pragma unroll for (int n = 0; n < KEY_FRAGMENTS; ++n) { #pragma unroll for (int i = 0; i < qk[n].num_elements; ++i) { - const int col = warp_k * (BK / KEY_WARPS) + n * 16 + fragment_col + - (i & 1) + ((i >> 2) & 1) * 4; + const int col = n * 16 + fragment_col + (i & 1) + ((i >> 2) & 1) * 4; + const int row = (i >> 1) & 1; qk[n].x[i] = start + col < length ? qk[n].x[i] * scale : -INFINITY; - row_max[(i >> 1) & 1] = fmaxf(row_max[(i >> 1) & 1], qk[n].x[i]); + row_max[n / 2][row] = fmaxf(row_max[n / 2][row], qk[n].x[i]); } } #pragma unroll - for (int r = 0; r < 2; ++r) { - row_max[r] = - fmaxf(row_max[r], __shfl_xor_sync(0xffffffff, row_max[r], 2)); - row_max[r] = - fmaxf(row_max[r], __shfl_xor_sync(0xffffffff, row_max[r], 8)); - const int row = warp_q * 16 + fragment_row + r * 2; - if ((lane & 10) == 0) scores[row * KEY_WARPS + warp_k] = row_max[r]; - } - sync_query_pair(warp_q); - float new_max[2], row_sum[2] = {0.f, 0.f}; + for (int half = 0; half < 2; ++half) { #pragma unroll - for (int r = 0; r < 2; ++r) { - const int row = warp_q * 16 + fragment_row + r * 2; - new_max[r] = maximum[row]; + for (int row = 0; row < 2; ++row) { + auto& maximum = row_max[half][row]; + maximum = fmaxf(maximum, __shfl_xor_sync(0xffffffff, maximum, 2)); + maximum = fmaxf(maximum, __shfl_xor_sync(0xffffffff, maximum, 8)); + } + } + float new_max[2], row_sum[2][2] = {{0.f, 0.f}, {0.f, 0.f}}; #pragma unroll - for (int w = 0; w < KEY_WARPS; ++w) - new_max[r] = fmaxf(new_max[r], scores[row * KEY_WARPS + w]); - register_alpha[r] = __expf(maximum[row] - new_max[r]); + for (int row = 0; row < 2; ++row) { + new_max[row] = + fmaxf(fmaxf(running_max[row], row_max[0][row]), row_max[1][row]); + register_alpha[row] = __expf(running_max[row] - new_max[row]); } #pragma unroll for (int n = 0; n < KEY_FRAGMENTS; ++n) { #pragma unroll for (int i = 0; i < qk[n].num_elements; ++i) { - const int r = (i >> 1) & 1; - const int row = warp_q * 16 + fragment_row + r * 2; - const int col = warp_k * (BK / KEY_WARPS) + n * 16 + fragment_col + - (i & 1) + ((i >> 2) & 1) * 4; - const float p = __expf(qk[n].x[i] - new_max[r]); - probabilities[row * PLD + col] = __float2half_rn(p); - row_sum[r] += p; + const int row = (i >> 1) & 1; + const float p = __expf(qk[n].x[i] - new_max[row]); + qk[n].x[i] = p; + row_sum[n / 2][row] += p; } - } - float* partial_sums = scores + BQ * KEY_WARPS; #pragma unroll - for (int r = 0; r < 2; ++r) { - row_sum[r] += __shfl_xor_sync(0xffffffff, row_sum[r], 2); - row_sum[r] += __shfl_xor_sync(0xffffffff, row_sum[r], 8); - const int row = warp_q * 16 + fragment_row + r * 2; - if ((lane & 10) == 0) - partial_sums[row * KEY_WARPS + warp_k] = row_sum[r]; + for (int i = 0; i < 4; ++i) { + union PackedPair { + half2 value; + unsigned bits; + } pair; + pair.value = __floats2half2_rn(qk[n].x[2 * i], qk[n].x[2 * i + 1]); + probability_pairs[n][i] = pair.bits; + } } - sync_query_pair(warp_q); - if (warp_k == 0 && (lane & 10) == 0) { #pragma unroll - for (int r = 0; r < 2; ++r) { - const int row = warp_q * 16 + fragment_row + r * 2; - float sum = 0.f; + for (int half = 0; half < 2; ++half) { #pragma unroll - for (int w = 0; w < KEY_WARPS; ++w) - sum += partial_sums[row * KEY_WARPS + w]; - denominator[row] = denominator[row] * register_alpha[r] + sum; - maximum[row] = new_max[r]; + for (int row = 0; row < 2; ++row) { + auto& sum = row_sum[half][row]; + sum += __shfl_xor_sync(0xffffffff, sum, 2); + sum += __shfl_xor_sync(0xffffffff, sum, 8); } } +#pragma unroll + for (int row = 0; row < 2; ++row) { + float sum = 0.f; + sum += row_sum[0][row]; + sum += row_sum[1][row]; + running_sum[row] = running_sum[row] * register_alpha[row] + sum; + running_max[row] = new_max[row]; + } } // Issue the next K/V global loads while the current V tile is consumed. union StagedVector { @@ -219,7 +211,7 @@ __global__ __launch_bounds__(THREADS, const int tile_row = (tile / (D / 32)) * 8 + (lane >> 2); const int next_row = start + BK + tile_row; const int next_col = (tile % (D / 32)) * 32 + (lane & 3) * 8; - if (next_row < length) { + if (tile_row < BK && next_row < length) { const int64_t position = base + int64_t(next_row) * heads * D + next_col; if ((reinterpret_cast(k) % 16 == 0) && @@ -247,19 +239,19 @@ __global__ __launch_bounds__(THREADS, } #pragma unroll for (int part = 0; part < OUTPUT_FRAGMENTS; ++part) { - const int col = warp_k * (D / KEY_WARPS) + part * 16; - const int row = warp_q * 16; - auto& pv = accumulators[part]; #pragma unroll - for (int i = 0; i < pv.num_elements; ++i) - pv.x[i] *= register_alpha[(i >> 1) & 1]; + for (int i = 0; i < accumulators[part].num_elements; ++i) + accumulators[part].x[i] *= register_alpha[(i >> 1) & 1]; + } +#pragma unroll + for (int kv = 0; kv < KEY_FRAGMENTS; ++kv) { + fi::AFragment pa; + load_probability_fragment(pa, probability_pairs[kv]); #pragma unroll - for (int kv = 0; kv < BK; kv += 16) { - fi::AFragment pa; + for (int part = 0; part < OUTPUT_FRAGMENTS; ++part) { fi::QKBFragment vb; - load_p_fragment(pa, probabilities, row, kv); - fi::load_qk_b_fragment(vb, vs + col * VLD + kv, VLD); - fi::mma_sync_m16n16k16_row_col_f16f16f32(pv, pa, vb); + fi::load_qk_b_fragment(vb, vs + part * 16 * VLD + kv * 16, VLD); + fi::mma_sync_m16n16k16_row_col_f16f16f32(accumulators[part], pa, vb); } } __syncthreads(); @@ -269,6 +261,7 @@ __global__ __launch_bounds__(THREADS, const int tile = (tid + n * THREADS) / 32; const int tile_row = (tile / (D / 32)) * 8 + (lane >> 2); const int next_col = (tile % (D / 32)) * 32 + (lane & 3) * 8; + if (tile_row >= BK) continue; *reinterpret_cast(ks + tile_row * QLD + next_col) = next_k[n].packed; // Transpose four rows with exact 32-bit lane exchanges. Each lane @@ -292,7 +285,8 @@ __global__ __launch_bounds__(THREADS, __shfl_xor_sync(0xffffffff, transposed[2 * j], 8); const unsigned other1 = __shfl_xor_sync(0xffffffff, transposed[2 * j + 1], 8); - const unsigned local = transposed[2 * j + ((lane & 8) >> 3)]; + const unsigned local = + ((lane & 8) ? transposed[2 * j + 1] : transposed[2 * j]); const unsigned other = (lane & 8) ? other1 : other0; const uint2 vector = (lane & 8) ? make_uint2(other, local) : make_uint2(local, other); @@ -313,7 +307,7 @@ __global__ __launch_bounds__(THREADS, (i & 1) + ((i >> 2) & 1) * 4; if (q_start + row < length) output[base + int64_t(q_start + row) * heads * D + col] = - __float2half_rn(pv.x[i] / denominator[row]); + __float2half_rn(pv.x[i] / running_sum[(i >> 1) & 1]); } } } diff --git a/tests/video/test_h3_numerics.py b/tests/video/test_h3_numerics.py index 90ef61b9ac..1d39f8543d 100644 --- a/tests/video/test_h3_numerics.py +++ b/tests/video/test_h3_numerics.py @@ -178,7 +178,7 @@ def test_attention_padding_excludes_poisoned_suffix(used, padded): @pytest.mark.skipif(not torch.cuda.is_available(), reason="requires GPU") -@pytest.mark.parametrize("length", [127, 128, 129, 12323]) +@pytest.mark.parametrize("length", [127, 128, 129, 191, 192, 193, 12323]) def test_flashinfer_online_softmax_across_tiles_and_batches(length): from vllm.model_executor.models.minimax_h3.cuda_ops import flashinfer_extension @@ -197,7 +197,9 @@ def test_flashinfer_online_softmax_across_tiles_and_batches(length): @pytest.mark.skipif(not torch.cuda.is_available(), reason="requires GPU") -@pytest.mark.parametrize("length", [31, 32, 33, 63, 64, 65, 127, 128, 129]) +@pytest.mark.parametrize( + "length", [31, 32, 33, 63, 64, 65, 127, 128, 129, 191, 192, 193, 385] +) def test_flashinfer_prefetch_tail_and_unaligned_storage(length): from vllm.model_executor.models.minimax_h3.cuda_ops import flashinfer_extension @@ -218,7 +220,7 @@ def test_flashinfer_prefetch_tail_and_unaligned_storage(length): @pytest.mark.skipif(not torch.cuda.is_available(), reason="requires GPU") -@pytest.mark.parametrize("length", [129, 257]) +@pytest.mark.parametrize("length", [129, 193, 257, 385]) def test_flashinfer_query_groups_have_independent_softmax_state(length): from vllm.model_executor.models.minimax_h3.cuda_ops import flashinfer_extension From c69cfc7024460e314e79a0bba37a3b736340bc6e Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 19:22:52 +0800 Subject: [PATCH 11/25] [Doc] Record complete FI register-kernel media preservation Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- .../FLASHINFER_REGISTER_PROBABILITY.md | 27 ++++++++++++++++--- 1 file changed, 24 insertions(+), 3 deletions(-) diff --git a/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md b/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md index dd16c95348..c0a2040098 100644 --- a/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md +++ b/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md @@ -71,6 +71,27 @@ packaging, and includes only one measurement per implementation. It is not the required full-request warmup-plus-three acceptance. Existing FA query-128 remains faster in its separately recorded full-request measurements. -Full native media preservation, formal repeated request measurements, -independent official references and human review remain separate gates. The -campaign remains incomplete. No precision gate or sampler setting is relaxed. +## Complete native media control + +The formatted native build at `f8b85c681a402ad26aac13adc701687a342881d9` +completes a separate pair of full TP4 W8A16 LightX2V four-step v1.2 requests. +Both implementations use the same current Python source, shared pageable VAE +host storage, zero persistent FP16 cache and the five-second request's +1280x736/124-frame internal canvas. Only the immutable FI binary differs. + +`fi-register-native-quality.json` passes every numerical gate. Final video +and audio latents, all 124 unencoded RGB frames and PCM match bitwise. Video +PSNR is infinite and SSIM is 1; audio spectral cosine is +0.9999999999999756 and RMS ratio is 1. Both requests pass native media and +strict actual-work validation. Their peak allocation is unchanged at +19,501,498,880 bytes/card. Commands, binary hashes and clean source provenance +are in `fi-register-native-pair.json` and `fi-register-native-summary.json`. + +The single captured cold requests take 71.085578 / 63.649644 seconds denoise +and 107.358301 / 100.664110 seconds request (control/candidate). These captures +include different first-use setup costs and are not a formal speed comparison. +Use the matched warmed denoise control above for the isolated 7.179% result. + +Formal repeated request measurements, independent official references and +human review remain separate gates. The campaign remains incomplete. No +precision gate or sampler setting is relaxed. From 4ed70419e7f42c6e9f4625f7fa92cf8dea2126d8 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 19:50:16 +0800 Subject: [PATCH 12/25] [Doc] Record formal FI results and all eight H3 Turbo controls Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- .../FLASHINFER_REGISTER_PROBABILITY.md | 27 ++++++++++++-- docs/design/minimax_h3/SM70_EPILOGUES.md | 37 ++++++++++++++++++- 2 files changed, 60 insertions(+), 4 deletions(-) diff --git a/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md b/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md index c0a2040098..7c32cb8526 100644 --- a/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md +++ b/docs/design/minimax_h3/FLASHINFER_REGISTER_PROBABILITY.md @@ -92,6 +92,27 @@ and 107.358301 / 100.664110 seconds request (control/candidate). These captures include different first-use setup costs and are not a formal speed comparison. Use the matched warmed denoise control above for the isolated 7.179% result. -Formal repeated request measurements, independent official references and -human review remain separate gates. The campaign remains incomplete. No -precision gate or sampler setting is relaxed. +## Formal repeated requests + +Source `c69cfc7024460e314e79a0bba37a3b736340bc6e` completes one full native +warmup and three requests without profiler or captures, with the same +media-checked configuration and immutable binary. The warmup takes 63.412712 +seconds denoise and 98.849412 seconds request. + +| Measurement | Denoise seconds | Complete request seconds | +| --- | ---: | ---: | +| 1 | 62.321408 | 91.585107 | +| 2 | 62.339754 | 91.385871 | +| 3 | 62.257260 | 95.401157 | + +Minimum-to-maximum rank median throughput is **49.794813-49.794831 useful +TFLOP/s/card**, using the slowest rank's complete denoise time. Denoise CV is +0.056762%; peak allocation remains 19,501,498,880 bytes/card. The full contract, +per-rank stages/steps, loaded binary hashes, source hashes and NVML samples are +retained in `fi-register-720p-three-runs/`. + +**The >80 performance gate fails.** The earlier FA query-128 configuration's +51.939 TFLOP/s/card remains the campaign's best formal result. The complete +native FI numerical control passes, but independent official references and +human review remain pending. No AUTO selection, precision relaxation or +campaign completion is claimed. diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 10c6d89577..593d4dd4c9 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -4,6 +4,11 @@ This development branch is stacked on the common prepared execution (#571) and workflow accounting (#578) branches. No configuration has passed the campaign's >80 useful TFLOP/s/card and complete official quality gates. +The separate [FI register-probability update](FLASHINFER_REGISTER_PROBABILITY.md) +now also has full native media preservation and a formal warmup-plus-three +result of 49.795 useful TFLOP/s/card. It remains below the FA query-128 result +and the campaign target. Both use the same shared projection interface. + ## Measured problem and implementation The matching four-step FA denoise profile spends 6.890 seconds in miscellaneous @@ -238,9 +243,39 @@ These are captured cold quality controls, not warmed three-run performance. Together with the existing v1.2 four-step and v1.0_768p eight-step controls, all six official FL2V Turbo artifacts now have a complete W8A16 T2VA numerical preservation result. Original-weight and keyframe combinations remain separate -pending coverage; Ref2V four-step is also still awaiting its full control. +pending coverage. The subsequent Ref2V four-step control is recorded below. No independent official quality or human acceptance is inferred. Evidence: `remaining-turbo-pairs.json`, `remaining-turbo-summary.json`, `light4-v{10,11,01}-720p-quality.json`, `light8-v10-non768-720p-quality.json` and the corresponding captured runs under the campaign's artifact root. + +## Four-step mixed references and complete adapter inventory + +Source `c69cfc7024460e314e79a0bba37a3b736340bc6e` completes the official +Ref2V four-step v0.1 adapter with a W8A16 Ref2VA base, one image, one +2.5-second video and one standalone audio reference. The video starts at zero; +seed 42, five sigma points, video/audio shifts 12/3 and alpha 8 are retained. +The candidate uses the same general FA query-128 path as the other adapters. + +`ref4-mixed-quality.json` passes all declared numerical gates against frozen +native `4f19ef7`: final video/audio latents, all 124 RGB frames and PCM match +bitwise, PSNR is infinite, SSIM is 1 and RMS ratio is 1. Spectral cosine is +0.9999999999999695. Both native generations complete. The candidate's strict +actual-work checks pass; its complete denoise is 184.563686 seconds, request +269.210367 seconds, and peak allocation 19,613,711,360 bytes/card. Corrected +useful throughput is 52.585556-52.585562 TFLOP/s/card. + +The frozen control takes 214.108869 seconds denoise and 321.250580 seconds +request with the same peak allocation. These are captured cold requests with +different host VAE sharing policies; they are not formal speed acceptance. +The old control script did not embed Git metadata. The separate +`baseline-source-audit.json` verifies all 2,483 tracked `vllm` files in the +frozen archive against `4f19ef7`, without changing the historical contract. + +All eight official LightX2V artifacts now have complete native numerical +preservation evidence: six FL2V adapters on T2VA and both Ref2V adapters with +mixed references. This is not the full task/weight/reference cross-product, +independent official quality, human review or >80 acceptance. Exact paths, +source identities and timing scope are retained in `ref4-mixed-pair.json`, +`ref4-mixed-summary.json` and the campaign result index. From ef18c81783fa1bc75c8b73881aa59595d0bd6f98 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 20:12:23 +0800 Subject: [PATCH 13/25] [Doc] Record full H3 first and last frame parity controls Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/SM70_EPILOGUES.md | 30 ++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 593d4dd4c9..9418cb385c 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -279,3 +279,33 @@ mixed references. This is not the full task/weight/reference cross-product, independent official quality, human review or >80 acceptance. Exact paths, source identities and timing scope are retained in `ref4-mixed-pair.json`, `ref4-mixed-summary.json` and the campaign result index. + +## First, last and both-frame controls + +Source `4ed70419e7f42c6e9f4625f7fa92cf8dea2126d8` also completes all three +FL2VA keyframe modes using the official Light4 v1.2_768p adapter and W8A16, +with one immutable engine per implementation. The candidate uses the same +general FA query-128/prepared/exact-residual/epilogue path. First and last +images are the retained frames 0 and 123 of the campaign sample, selected +with indices `[0]`, `[-1]` and `[0,-1]` respectively. + +| Constraint | Frozen-control denoise seconds | Candidate denoise seconds | Candidate request seconds | Candidate useful TFLOP/s/card, minimum | +| --- | ---: | ---: | ---: | ---: | +| First frame | 77.275887 | 66.419303 | 106.123582 | 50.697486 | +| Last frame | 71.537050 | 64.852263 | 102.142091 | 51.922501 | +| First and last frames | 77.299085 | 69.615145 | 108.431893 | 52.304580 | + +Every pair passes complete numerical preservation: final video/audio latents, +all 124 decoded frames and PCM match bitwise; SSIM and RMS ratio are 1. +Peak allocation is unchanged within each pair: 19,512,957,952 bytes/card +for a single image and 19,522,796,544 bytes/card for both images. + +These are captured requests with different first-use state, reference lengths +and host VAE sharing policies. They establish full execution and native +preservation, not a formal performance comparison or independent verification +of reference fidelity. Original-weight/other-adapter keyframes and the full +legal Ref2VA combination matrix remain pending. + +Evidence: `keyframe-pairs.json`, `keyframe-summary.json`, and +`keyframe-{first,last,first-last}-quality.json`. The candidate contract retains +shared-operator, model and video-source hashes plus clean Git provenance. From 3280edbfcc00b39cc3c2ca3c08ea6487b6e85a5b Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 20:51:29 +0800 Subject: [PATCH 14/25] [Doc] Record H3 workflow controls and remaining performance gates Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 54 ++++++++++++++++++++++ docs/design/minimax_h3/SM70_EPILOGUES.md | 32 +++++++++++++ 2 files changed, 86 insertions(+) create mode 100644 docs/design/minimax_h3/CAMPAIGN_RESULTS.md diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md new file mode 100644 index 0000000000..d92ab3a598 --- /dev/null +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -0,0 +1,54 @@ +# H3 SM70 campaign results + +No configuration has completed the >80 useful TFLOP/s/card, independent official quality and human-review gates. + +All rows use TP4. Rows marked formal use a complete request warmup plus three unprofiled requests. Other timings are captured cold diagnostics. Original floating-weight controls with legacy FLOP accounting omit throughput. + +| Workflow | Weights | Backend | Denoise seconds | Minimum card TFLOP/s | Native numerical control | Formal performance | +| --- | --- | --- | ---: | ---: | --- | --- | +| FL2V Light4 v1.2_768p | W8A16 | FLASH_ATTN_V100 | 59.749 | 51.939 | passed | failed >80 | +| light4-v10 | W8A16 | FLASH_ATTN_V100 | 61.263 | 50.655 | passed | not measured | +| light4-v11 | W8A16 | FLASH_ATTN_V100 | 61.019 | 50.858 | passed | not measured | +| light4-v01 | W8A16 | FLASH_ATTN_V100 | 61.000 | 50.873 | passed | not measured | +| light8-v10-non768 | W8A16 | FLASH_ATTN_V100 | 120.478 | 51.516 | passed | not measured | +| FL2V Light8 v1.0_768p | W8A16 | FLASH_ATTN_V100 | 121.542 | 51.065 | passed | not measured | +| Ref2V Light8 v1.0_768p, image/video/audio | W8A16 | FLASH_ATTN_V100 | 366.800 | 52.919 | passed | not measured | +| Ref2V Light4 v0.1, image/video/audio | W8A16 | FLASH_ATTN_V100 | 184.564 | 52.586 | passed | not measured | +| FL2V Light4 v1.2_768p, register FI | W8A16 | FLASHINFER_SM70 | 62.321 | 49.795 | passed | failed >80 | +| FlashGen four-step | original floating | FLASH_ATTN_V100 | 59.488 | 51.621 | passed | not measured | +| FastH3 Dense data-free | original floating | FLASH_ATTN_V100 | 56.224 | 54.125 | passed | not measured | +| FastH3 VSA data-free | original floating | FASTVIDEO_VSA | 37.387 | 45.268 | failed | not measured | +| FL2V Light4 v1.2_768p, original floating | original floating | FLASH_ATTN_V100 | 66.366 | not measured | passed | not measured | +| FL2V Light4 v1.2_768p, first | W8A16 | FLASH_ATTN_V100 | 66.419 | 50.697 | passed | not measured | +| FL2V Light4 v1.2_768p, last | W8A16 | FLASH_ATTN_V100 | 64.852 | 51.923 | passed | not measured | +| FL2V Light4 v1.2_768p, first-last | W8A16 | FLASH_ATTN_V100 | 69.615 | 52.305 | passed | not measured | + +The VSA failure is against an explicitly labeled FP32 selected-key diagnostic, not the unmodified official GPU kernel. Generation alone does not establish numerical quality. + +- Only explicitly marked formal rows are full-request warmup-plus-three result; other timings are captured cold diagnostics. +- Useful FLOPs exclude padding, duplicate work and skipped sparse/cache work; denominator is the slowest complete-denoise rank. +- Native parity does not establish independent official-model or human audiovisual quality. +- The 720p-family request uses an internal 1280x736/124-frame canvas; Ref2VA has longer conditioning sequences. +- First/last/both keyframes pass native controls for W8A16 Light4 v1.2; remaining adapter/weight keyframes, legal reference combinations, and primary-shape/TP matrix remain incomplete. +- TeaCache and Cache-DiT/SCM currently have separate small-shape lifecycle evidence, not primary >80 or official quality acceptance. + +Exact source/run paths and the evidence index are retained in `campaign-results.json` and `campaign-results.csv`. + +## Additional bounded experiments + +These artifact-only experiments did not change production defaults: + +- Increasing NCCL CTA counts from 8 to 16/32 changes FP32 reduction bits. + The preliminary bitwise gate rejected them before timing or a full-model + comparison; this is not a measured full-model quality failure. +- Sequential FA warp-operand loading reduces one register count from 248 to + 235 but leaves the proposed occupancy budget unmet. Build evidence was + sufficient to reject the hypothesis; no GPU benchmark was run. +- K-only swizzling in the new FI kernel preserves operator bits but changes + the paired median by only 0.192%, within observed clock variation. No + complete-model run or production integration was justified. + +The artifact folders `nccl-cta-control`, `attention-single-warp-buffer` and +`attention-fi-key-only-swizzle` retain hypotheses, source hashes and results. +Do not repeat these experiments without a changed hypothesis. Communication +work continues separately; it has no accepted model-level speedup yet. diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 9418cb385c..9cd9f7be1c 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -309,3 +309,35 @@ legal Ref2VA combination matrix remain pending. Evidence: `keyframe-pairs.json`, `keyframe-summary.json`, and `keyframe-{first,last,first-last}-quality.json`. The candidate contract retains shared-operator, model and video-source hashes plus clean Git provenance. + +## Original floating FlashGen and FastH3 Dense controls + +Both four-interval T2VA variants now have complete native comparisons with +frozen mainline `4f19ef7a20db60bb0685e599bd3f4dd156202eed`. Each pair uses +original floating FL2VA weights, its matching official adapter, seed 42, +TP4, flow shifts 12/3 and the same 1280x736/124-frame internal canvas. +Candidate source and individual file hashes are retained in each contract; +the frozen source audit covers all 2,483 tracked package files with no mismatch. + +Both final video/audio latents, all 124 pre-encoding RGB frames and PCM match +bitwise for both variants. SSIM and RMS ratio are 1; spectral cosine exceeds +0.99999999999996. This establishes native numerical preservation for these +configurations, not independent official-model or human quality acceptance. + +| Variant | Baseline denoise / request seconds | Candidate denoise / request seconds | Candidate useful TFLOP/s/card | Candidate peak bytes/card | +| --- | ---: | ---: | ---: | ---: | +| FlashGen four-step | 73.290174 / 162.679376 | 59.488178 / 121.697365 | 51.620900 | 20,781,940,224 | +| FastH3 Dense data-free | 61.574164 / 121.787763 | 56.223567 / 114.632281 | 54.124949 | 20,023,082,496 | + +These are captured cold controls. Both arms use pageable host weights; the +candidate also shares host VAE weights and enables prepared column weights, +exact residual sharding, shared epilogues and explicit FA query tile 128. +The measurements combine these changes and do not isolate a kernel effect. +Neither variant has completed a full warmup plus three unprofiled requests. + +Evidence: `original-variant-pairs.json`, `original-variant-summary.json`, +`flashgen-original-quality.json`, `fasth3-dense-original-quality.json` and +`baseline-source-audit.json` in the campaign artifact root. Complete media and +contracts reside in the corresponding `*-original-{baseline,candidate}` runs. +The [campaign table](CAMPAIGN_RESULTS.md) separates these diagnostic timings +from formal acceptance measurements. Every configuration remains unqualified. From 6d2a44b8d022b8d87226a811d291e0196a9d6da3 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 21:33:26 +0800 Subject: [PATCH 15/25] [Kernel] Add explicit calibrated SM70 local-row reduction Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- .../benchmark_sm70_exact_row_reduce.py | 228 +++++++++++++++ csrc/sm70_turbomind/ops/exact_row_reduce.cu | 165 +++++++++++ docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 6 + docs/design/minimax_h3/EXACT_ROW_REDUCTION.md | 87 ++++++ tests/video/test_h3_provenance.py | 15 + tests/video/test_sm70_collectives.py | 55 ++++ .../layers/sm70_collective_calibration.py | 276 ++++++++++++++++++ .../model_executor/layers/sm70_collectives.py | 265 +++++++++++++++++ vllm/video/metrics.py | 2 +- 9 files changed, 1098 insertions(+), 1 deletion(-) create mode 100644 benchmarks/kernels/benchmark_sm70_exact_row_reduce.py create mode 100644 csrc/sm70_turbomind/ops/exact_row_reduce.cu create mode 100644 docs/design/minimax_h3/EXACT_ROW_REDUCTION.md create mode 100644 tests/video/test_sm70_collectives.py create mode 100644 vllm/model_executor/layers/sm70_collective_calibration.py create mode 100644 vllm/model_executor/layers/sm70_collectives.py diff --git a/benchmarks/kernels/benchmark_sm70_exact_row_reduce.py b/benchmarks/kernels/benchmark_sm70_exact_row_reduce.py new file mode 100644 index 0000000000..19736734c4 --- /dev/null +++ b/benchmarks/kernels/benchmark_sm70_exact_row_reduce.py @@ -0,0 +1,228 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""TP4 explicit row-plan correctness controls; launch under an owned GPU lease. + +Use torchrun --standalone --nproc_per_node=4 with this script. Measurements +are isolated communication diagnostics, never full-model acceptance. +""" + +import argparse +import hashlib +import importlib.util +import json +import os +import statistics +import sys +import time +from pathlib import Path + +import torch + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--output", type=Path, required=True) + parser.add_argument("--extension", type=Path) + parser.add_argument("--full-shape", action="store_true") + args = parser.parse_args() + from vllm.config import ParallelConfig, VllmConfig, set_current_vllm_config + from vllm.distributed import ( + cleanup_dist_env_and_memory, + get_tp_group, + init_distributed_environment, + initialize_model_parallel, + ) + from vllm.model_executor.layers import sm70_collectives as shared + from vllm.video.benchmark import source_provenance + + rank, local = int(os.environ["RANK"]), int(os.environ["LOCAL_RANK"]) + if int(os.environ["WORLD_SIZE"]) != 4: + raise ValueError("This control requires exactly four TP ranks") + torch.cuda.set_device(local) + torch.set_num_threads(4) + torch.manual_seed(5091 + rank) + record = dict( + rank=rank, + state="running", + cases=[], + guards=[], + source=source_provenance(), + scope="isolated operator control; no model acceptance", + ) + if args.extension: + spec = importlib.util.spec_from_file_location( + args.extension.stem, args.extension + ) + extension = importlib.util.module_from_spec(spec) + spec.loader.exec_module(extension) + shared._extension = lambda: extension + sys.modules["onecat_sm70_exact_reduce"] = extension + record["extension"] = { + "path": str(args.extension), + "sha256": hashlib.sha256(args.extension.read_bytes()).hexdigest(), + } + source = ( + Path(shared.__file__).resolve().parents[3] + / "csrc/sm70_turbomind/ops/exact_row_reduce.cu" + ) + record["cuda_source_sha256"] = hashlib.sha256(source.read_bytes()).hexdigest() + record["benchmark_source_sha256"] = hashlib.sha256( + Path(__file__).read_bytes() + ).hexdigest() + with set_current_vllm_config( + VllmConfig(parallel_config=ParallelConfig(tensor_parallel_size=4)) + ): + init_distributed_environment(4, rank, "env://", local, "nccl") + initialize_model_parallel(4) + try: + group = get_tp_group() + for label, shape, budget in ( + ("one-rank-invalid-shape", (3, 3) if rank == 0 else (4, 3), 2**30), + ("one-rank-small-budget", (4, 3), 1 if rank == 0 else 2**30), + ("different-valid-shapes", (4, 3) if rank == 0 else (8, 3), 2**30), + ): + try: + shared.SM70ExactRowReductionPlan( + group, shape, memory_budget_bytes=budget + ) + except ValueError: + record["guards"].append(label) + else: + raise AssertionError(f"Expected collective rejection: {label}") + shapes = [ + (4, 3), + (12, 65), + (68, 257), + (128, 768), + (260, 1024), + (1028, 3072), + ] + if args.full_shape: + shapes.append((34560, 5376)) + for shape in shapes: + started = time.perf_counter() + plan = shared.SM70ExactRowReductionPlan( + group, shape, memory_budget_bytes=4 * 2**30 + ) + calibration_seconds = time.perf_counter() - started + try: + with torch.inference_mode(): + for label, scale in ( + ("ordinary", 1.0), + ("wide", 1e30), + ("subnormal", 1e-40), + ): + storage = ( + torch.randn(shape[0] * shape[1] + 1, device="cuda") + * scale + ) + value = storage[1:].view(shape) + ref = group.all_reduce(value).chunk(4)[rank] + actual = plan.reduce(value) + mismatch = int( + torch.count_nonzero( + ref.view(torch.int32) != actual.view(torch.int32) + ) + ) + record["cases"].append( + dict( + shape=shape, + input=label, + storage_offset=1, + mismatch=mismatch, + calibration_seconds=calibration_seconds, + raw_ipc_bytes=plan.raw_ipc_bytes, + ) + ) + value = torch.full( + shape, + float("inf") if rank < 2 else -float("inf"), + device="cuda", + ) + ref = group.all_reduce(value).chunk(4)[rank] + actual = plan.reduce(value) + record["cases"].append( + dict( + shape=shape, + input="opposing-inf", + mismatch=int( + torch.count_nonzero( + ref.view(torch.int32) + != actual.view(torch.int32) + ) + ), + ) + ) + with torch.cuda.stream(torch.cuda.Stream()): + try: + plan.reduce(value) + except RuntimeError: + record["guards"].append("different-stream") + else: + raise AssertionError("Different stream was accepted") + vote = torch.tensor( + int(all(x["mismatch"] == 0 for x in record["cases"])), + device="cuda", + ) + torch.distributed.all_reduce( + vote, op=torch.distributed.ReduceOp.MIN + ) + if not vote.item(): + raise AssertionError("Native FP32 bits changed") + if shape == (34560, 5376): + value.normal_() + for _ in range(3): + group.all_reduce(value) + plan.reduce(value) + torch.cuda.synchronize() + times = {"native": [], "peer_rows": []} + for repeat in range(7): + for name in ( + ("native", "peer_rows") + if repeat % 2 == 0 + else ("peer_rows", "native") + ): + torch.distributed.barrier(group=group.cpu_group) + torch.cuda.synchronize() + start, end = ( + torch.cuda.Event(enable_timing=True), + torch.cuda.Event(enable_timing=True), + ) + start.record() + output = ( + group.all_reduce(value) + if name == "native" + else plan.reduce(value) + ) + end.record() + end.synchronize() + times[name].append(start.elapsed_time(end)) + del output + record["times_ms"] = times + record["median_ms"] = { + key: statistics.median(values) + for key, values in times.items() + } + finally: + plan.close() + try: + plan.reduce(value) + except RuntimeError: + record["guards"].append("closed-plan") + else: + raise AssertionError("Closed plan was accepted") + print( + json.dumps(dict(rank=rank, shape=shape, state="passed")), flush=True + ) + record["state"] = "passed_operator_control" + except BaseException as error: + record.update(state="failed", error=repr(error)) + raise + finally: + args.output.mkdir(parents=True, exist_ok=True) + (args.output / f"rank-{rank}.json").write_text(json.dumps(record, indent=2)) + cleanup_dist_env_and_memory() + + +if __name__ == "__main__": + main() diff --git a/csrc/sm70_turbomind/ops/exact_row_reduce.cu b/csrc/sm70_turbomind/ops/exact_row_reduce.cu new file mode 100644 index 0000000000..57e7472379 --- /dev/null +++ b/csrc/sm70_turbomind/ops/exact_row_reduce.cu @@ -0,0 +1,165 @@ +// SPDX-License-Identifier: Apache-2.0 +// SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +#include +#include +#include +#include +#include +#include +#include + +struct Peers { + float* data[4]; + unsigned* flags[4]; +}; +constexpr int GRID = 80; +__device__ __forceinline__ void publish(unsigned* p, unsigned value) { + asm volatile("st.release.sys.global.u32 [%0], %1;" ::"l"(p), "r"(value) + : "memory"); +} +__device__ __forceinline__ unsigned acquire(unsigned* p) { + unsigned value; + asm volatile("ld.acquire.sys.global.u32 %0, [%1];" + : "=r"(value) + : "l"(p) + : "memory"); + return value; +} +__device__ __forceinline__ void barrier(Peers peers, int rank, unsigned epoch, + int phase) { + __syncthreads(); + if (threadIdx.x < 4) { + int peer = threadIdx.x; + int location = phase * GRID * 4 + blockIdx.x * 4; + publish(peers.flags[peer] + location + rank, epoch); + while (acquire(peers.flags[rank] + location + peer) != epoch) { + } + } + __syncthreads(); +} +__device__ __forceinline__ float tree(unsigned code, float x0, float x1, + float x2, float x3) { + switch (code) { + case 0: + return __fadd_rn(x0, __fadd_rn(x1, __fadd_rn(x2, x3))); + case 1: + return __fadd_rn(x0, __fadd_rn(__fadd_rn(x1, x2), x3)); + case 2: + return __fadd_rn(x0, __fadd_rn(__fadd_rn(x1, x3), x2)); + case 3: + return __fadd_rn(__fadd_rn(x0, x1), __fadd_rn(x2, x3)); + case 4: + return __fadd_rn(__fadd_rn(x0, x2), __fadd_rn(x1, x3)); + case 5: + return __fadd_rn(__fadd_rn(x0, x3), __fadd_rn(x1, x2)); + case 6: + return __fadd_rn(__fadd_rn(x0, __fadd_rn(x1, x2)), x3); + case 7: + return __fadd_rn(__fadd_rn(__fadd_rn(x0, x1), x2), x3); + case 8: + return __fadd_rn(__fadd_rn(__fadd_rn(x0, x2), x1), x3); + case 9: + return __fadd_rn(__fadd_rn(x0, __fadd_rn(x1, x3)), x2); + case 10: + return __fadd_rn(__fadd_rn(__fadd_rn(x0, x1), x3), x2); + case 11: + return __fadd_rn(__fadd_rn(__fadd_rn(x0, x3), x1), x2); + case 12: + return __fadd_rn(__fadd_rn(x0, __fadd_rn(x2, x3)), x1); + case 13: + return __fadd_rn(__fadd_rn(__fadd_rn(x0, x2), x3), x1); + case 14: + return __fadd_rn(__fadd_rn(__fadd_rn(x0, x3), x2), x1); + default: + return __int_as_float(0x7fffffff); + } +} +__global__ void reduce_rows(Peers peers, const uint8_t* codes, float* output, + int64_t count, int64_t offset, int rank, + unsigned epoch) { + barrier(peers, rank, epoch, 0); + for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x; i < count; + i += int64_t(gridDim.x) * blockDim.x) { + int64_t at = offset + i; + float x0 = __ldcg(peers.data[0] + at); + float x1 = __ldcg(peers.data[1] + at); + float x2 = __ldcg(peers.data[2] + at); + float x3 = __ldcg(peers.data[3] + at); + output[i] = tree(codes[i], x0, x1, x2, x3); + } + barrier(peers, rank, epoch, 1); +} +std::tuple allocate(int64_t bytes) { + TORCH_CHECK(bytes > 0); + void* ptr = nullptr; + C10_CUDA_CHECK(cudaMalloc(&ptr, bytes)); + cudaIpcMemHandle_t handle; + try { + C10_CUDA_CHECK(cudaIpcGetMemHandle(&handle, ptr)); + C10_CUDA_CHECK(cudaMemset(ptr, 0, bytes)); + } catch (...) { + cudaFree(ptr); + throw; + } + return {reinterpret_cast(ptr), + pybind11::bytes(reinterpret_cast(&handle), sizeof(handle))}; +} +int64_t open_handle(pybind11::bytes bytes) { + std::string value = bytes; + TORCH_CHECK(value.size() == sizeof(cudaIpcMemHandle_t)); + cudaIpcMemHandle_t handle; + std::memcpy(&handle, value.data(), sizeof(handle)); + void* ptr = nullptr; + C10_CUDA_CHECK( + cudaIpcOpenMemHandle(&ptr, handle, cudaIpcMemLazyEnablePeerAccess)); + return reinterpret_cast(ptr); +} +void release(int64_t ptr, bool owner) { + if (owner) + C10_CUDA_CHECK(cudaFree(reinterpret_cast(ptr))); + else + C10_CUDA_CHECK(cudaIpcCloseMemHandle(reinterpret_cast(ptr))); +} +void run(torch::Tensor input, torch::Tensor codes, torch::Tensor output, + std::vector pointers, std::vector flags, int rank, + unsigned epoch) { + TORCH_CHECK(input.is_cuda() && input.scalar_type() == torch::kFloat32 && + input.is_contiguous()); + TORCH_CHECK(codes.is_cuda() && codes.scalar_type() == torch::kUInt8 && + codes.is_contiguous()); + TORCH_CHECK(output.is_cuda() && output.scalar_type() == torch::kFloat32 && + output.is_contiguous()); + TORCH_CHECK(input.device() == codes.device() && + input.device() == output.device()); + TORCH_CHECK(rank >= 0 && rank < 4 && pointers.size() == 4 && + flags.size() == 4 && epoch > 0); + TORCH_CHECK(input.numel() == output.numel() * 4 && + codes.numel() == output.numel()); + c10::cuda::CUDAGuard guard(input.device()); + auto stream = at::cuda::getCurrentCUDAStream(); + const auto* properties = at::cuda::getDeviceProperties(input.get_device()); + TORCH_CHECK( + properties->major == 7 && properties->minor == 0 && + properties->multiProcessorCount >= GRID, + "Exact peer reduction requires an SM70 device with at least 80 SMs"); + Peers peers; + for (int i = 0; i < 4; ++i) { + TORCH_CHECK(pointers[i] && flags[i]); + peers.data[i] = reinterpret_cast(pointers[i]); + peers.flags[i] = reinterpret_cast(flags[i]); + } + C10_CUDA_CHECK(cudaMemcpyAsync(peers.data[rank], input.data_ptr(), + input.nbytes(), cudaMemcpyDeviceToDevice, + stream)); + reduce_rows<<>>( + peers, codes.data_ptr(), output.data_ptr(), + output.numel(), output.numel() * rank, rank, epoch); + C10_CUDA_KERNEL_LAUNCH_CHECK(); +} +PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { + m.def("allocate", &allocate); + m.def("open_handle", &open_handle); + m.def("release", &release); + m.def("run", &run); +} diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index d92ab3a598..1e7b1e0916 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -52,3 +52,9 @@ The artifact folders `nccl-cta-control`, `attention-single-warp-buffer` and `attention-fi-key-only-swizzle` retain hypotheses, source hashes and results. Do not repeat these experiments without a changed hypothesis. Communication work continues separately; it has no accepted model-level speedup yet. + +The explicit [shared row-reduction interface](EXACT_ROW_REDUCTION.md) has +separate operator and prototype full-media controls. Its 2.49348% paired +denoise improvement is a development measurement, not a new formal campaign +result. H3 runtime integration and final-interface media validation remain +pending, and the ordinary reduction stays selected. diff --git a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md new file mode 100644 index 0000000000..94294c7765 --- /dev/null +++ b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md @@ -0,0 +1,87 @@ +# Explicit SM70 local-row reduction + +The shared `SM70ExactRowReductionPlan` interface is experimental and has no +automatic dispatch or H3 runtime selection yet. The ordinary residual path +continues to use FP32 all-reduce followed by a local-row slice. No configuration +has passed the campaign's >80 useful TFLOP/s/card and complete quality gates. + +## Arithmetic and ownership + +A conventional reduce-scatter changes FP32 addition order relative to the +existing all-reduce, and the earlier full H3 control failed numerical gates. +This interface instead classifies the native communicator's addition order +for the actual two-dimensional FP32 shape. Ten fixed finite probes distinguish +all 15 four-input binary addition trees. Unmatched or ambiguous elements reject +setup collectively. Calibration never reads model weights or model activations. + +The CUDA implementation copies each rank's partial input into owned IPC +storage, reads only its destination rows from peers and evaluates the calibrated +tree with rounded FP32 additions. System release/acquire flags establish input +visibility and completion. The grid has 80 blocks; the interface requires +SM70 devices with at least 80 SMs, four distinct peer-accessible devices on +one host, and consistent CUDA visibility. All row-parallel adapter contributions +must already be included in the input. + +The caller supplies an explicit budget covering persistent IPC buffers, +local output, one-byte arithmetic codes and calibration scratch. GPU buffers +belong to the plan, not a global shape cache. Returned tensors alias the plan's +output; consume them before the next call. Call `close()` collectively before +tearing down the TP group. Rank-dependent setup errors are exchanged over the +CPU group, and peer handles close before owners free their allocations. + +The plan rejects another stream/device, autograd inputs, incompatible layouts, +CUDA Graph execution and epoch exhaustion. It does not silently change a +backend or precision. Callers retain their ordinary collective when a plan +is unsuitable. Only the shared operator is provided in this change; a model +must explicitly own its lifecycle before integrating it into requests. + +## Validation and measured limits + +Environment: Torch 2.10.0+cu128, CUDA toolkit 12.8.93, NCCL 2.27.5, four leased +V100 SXM2 32GB cards. Evidence root: +`/data/minimax-h3/sm70-general-20260909/exact-peer-reduction/`. + +- The first arithmetic classification covers every element of the real + 34560x5376 projection. Three independent wide-range inputs and signed-zero + controls match the native all-reduce bitwise. +- The source implementation's TP4 control covers seven shapes from 4x3 through + 34560x5376, including non-H3 DiT widths, tails and non-aligned storage offsets. + All four ranks preserve bits for ordinary, wide, subnormal and opposing + infinity inputs: 112 numerical cases. Mismatched shapes, insufficient + per-rank budgets, different streams and closed plans are rejected. +- The prototype's independent four-rank memcheck fixture reports zero errors + on every rank. An earlier NCCL-bearing fixture reported only initialization + `cudaFuncGetAttributes` probes for unsupported kernels; NCCL explicitly + skips that return code in its [corresponding source](https://github.com/NVIDIA/nccl/blob/v2.27.5-1/src/enqueue.cc#L37-L38). + The isolated fixture uses Gloo for coordination and an independent FP32 + arithmetic reference; no CUDA API error suppression was applied. +- Source operator medians, including the full input copy and device barriers: + native 14.561–14.641 ms, peer rows 10.939–10.991 ms. These seven alternating + measurements are communication diagnostics only. +- A separate artifact override at source `3280edbfcc` completes a full denoise + warmup per implementation, then one measurement each: 59.717282 seconds + native versus 58.228244 seconds peer rows, a 2.49348% reduction. Candidate + useful throughput is 53.295148–53.295167 TFLOP/s/card. Every final video/audio + latent bit matches across all four passes, and the baseline also matches + the previously frozen FA query-128 control. Each candidate pass uses 400 + peer reductions. This is not the full-request warmup-plus-three protocol. +- The prototype's full native media control also preserves both final latents, + all 124 RGB frames and PCM bitwise. SSIM and RMS ratio are 1; spectral cosine + exceeds 0.99999999999998. Its captured cold request takes 94.269191 seconds, + including 61.688596 seconds denoise. This is native preservation, not an + independent official-model or human audiovisual review. +- That full request peaks at 19,732,554,240 PyTorch-allocated bytes/card plus + 743,180,800 persistent raw IPC bytes/card. Their sum is 20,475,735,040 bytes; + driver/library overhead is additional. Do not report only the PyTorch number. + +The prototype full-model evidence precedes the packaged plan's dynamic +calibration and setup guards. It is not substituted for full-model validation +of this final interface. Native integration, finalized-interface media controls, +formal repeated requests, TP/shape breadth and official/human quality gates +remain incomplete. No AUTO promotion is made. + +Reproduce the operator control with an owned native GPU lease and +`torchrun --standalone --nproc_per_node=4 +benchmarks/kernels/benchmark_sm70_exact_row_reduce.py --output +--full-shape`. The optional `--extension` pins an already-built library; the +report records its SHA256 plus benchmark, CUDA and shared Python source hashes. diff --git a/tests/video/test_h3_provenance.py b/tests/video/test_h3_provenance.py index 1048b546e9..63ee157cc1 100644 --- a/tests/video/test_h3_provenance.py +++ b/tests/video/test_h3_provenance.py @@ -34,3 +34,18 @@ def test_provenance_tracks_shared_operators_outside_model_directory( assert {k: v for k, v in before.items() if k != changed} == { k: v for k, v in after.items() if k != changed } + + +def test_provenance_tracks_loaded_generic_sm70_binary(tmp_path, monkeypatch): + import sys + from types import SimpleNamespace + + from vllm.video.metrics import loaded_kernel_provenance + + binary = tmp_path / "exact_reduce.so" + binary.write_bytes(b"generic collective binary") + monkeypatch.setitem( + sys.modules, "onecat_sm70_exact_reduce", SimpleNamespace(__file__=str(binary)) + ) + result = loaded_kernel_provenance() + assert result[str(binary)] == hashlib.sha256(binary.read_bytes()).hexdigest() diff --git a/tests/video/test_sm70_collectives.py b/tests/video/test_sm70_collectives.py new file mode 100644 index 0000000000..6d63a8bca6 --- /dev/null +++ b/tests/video/test_sm70_collectives.py @@ -0,0 +1,55 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Host checks for the explicit, non-automatic SM70 collective interface.""" + +import numpy as np +import pytest + +from vllm.model_executor.layers.sm70_collective_calibration import PROBES +from vllm.model_executor.layers.sm70_collectives import _layout + + +@pytest.mark.parametrize( + "shape", [(0, 8), (3, 8), (4, 0), (-4, 8), (4,), (True, 8), (4, 2**63)] +) +def test_reject_invalid_layout(shape): + with pytest.raises((TypeError, ValueError)): + _layout(shape) + + +def test_budget_covers_ipc_output_codes_and_calibration(): + shape, raw, resident, calibration = _layout((34560, 5376)) + assert shape == (34560, 5376) + assert raw == 743_180_800 + assert resident > raw + assert calibration > resident + 2 * 34560 * 5376 * 4 + + +def _trees(indices): + if len(indices) == 1: + return [indices[0]] + result: list[tuple] = [] + # Anchor the first leaf on the left to remove commutative duplicates. + for mask in range(1, (1 << len(indices)) - 1, 2): + left = tuple(x for i, x in enumerate(indices) if mask & (1 << i)) + right = tuple(x for i, x in enumerate(indices) if not mask & (1 << i)) + result.extend((a, b) for a in _trees(left) for b in _trees(right)) + return result + + +def _evaluate(tree, values): + if isinstance(tree, int): + return np.float32(values[tree]) + return np.float32(_evaluate(tree[0], values) + _evaluate(tree[1], values)) + + +def test_fixed_probes_cover_all_fp32_addition_trees(): + trees = _trees((0, 1, 2, 3)) + assert len(trees) == 15 + actual = { + tuple(int(_evaluate(tree, values).view(np.uint32)) for values, _ in PROBES) + for tree in trees + } + stored = {tuple(bits[i] for _, bits in PROBES) for i in range(15)} + assert len(stored) == 15 + assert actual == stored diff --git a/vllm/model_executor/layers/sm70_collective_calibration.py b/vllm/model_executor/layers/sm70_collective_calibration.py new file mode 100644 index 0000000000..119840ac36 --- /dev/null +++ b/vllm/model_executor/layers/sm70_collective_calibration.py @@ -0,0 +1,276 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""GPU classification of FP32 addition trees for explicit SM70 row plans.""" + +from vllm.triton_utils import tl, triton + + +@triton.jit +def update_mask( + reference, + masks, + expected, + N: tl.constexpr, + OFFSET: tl.constexpr, + BLOCK: tl.constexpr, +): + i = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + bits = tl.load(reference + OFFSET + i, i < N, other=0).to(tl.uint32, bitcast=True) + mask = tl.load(masks + i, i < N, other=0).to(tl.uint32) + for tree in tl.static_range(15): + correct = bits == tl.load(expected + tree) + mask = mask & tl.where(correct, 0x7FFF, 0x7FFF ^ (1 << tree)).to(tl.uint32) + tl.store(masks + i, mask.to(tl.int16), i < N) + + +@triton.jit +def decode_mask(masks, codes, stats, N: tl.constexpr, BLOCK: tl.constexpr): + i = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + mask = tl.load(masks + i, i < N, other=0).to(tl.uint32) + unique = (mask != 0) & ((mask & (mask - 1)) == 0) + code = tl.full((BLOCK,), 255, tl.uint8) + for tree in tl.static_range(15): + code = tl.where(mask == (1 << tree), tree, code).to(tl.uint8) + tl.store(codes + i, code, i < N) + tl.atomic_add(stats, tl.sum(((i < N) & (mask == 0)).to(tl.int32))) + tl.atomic_add(stats + 1, tl.sum(((i < N) & (mask != 0) & (~unique)).to(tl.int32))) + + +# Fixed finite probes distinguish all 15 four-input binary addition trees. +PROBES = ( + ( + (-55.133113861083984, -124364.703125, -361660.78125, 0.006582669448107481), + ( + 3370997780, + 3370997780, + 3370997779, + 3370997780, + 3370997779, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + 3370997780, + ), + ), + ( + ( + -1.0654968036760692e-06, + 7.75692081451416, + -3.822844155365601e-05, + 0.0002807896235026419, + ), + ( + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009773, + 1090009772, + 1090009772, + 1090009772, + ), + ), + ( + ( + 2.3551223193862825e-07, + 142.72190856933594, + -489.4043884277344, + 77.70187377929688, + ), + ( + 3280371076, + 3280371077, + 3280371076, + 3280371076, + 3280371076, + 3280371077, + 3280371077, + 3280371077, + 3280371077, + 3280371076, + 3280371076, + 3280371076, + 3280371076, + 3280371076, + 3280371076, + ), + ), + ( + ( + 204.39108276367188, + -0.6318132877349854, + -46.683990478515625, + 0.0033716242760419846, + ), + ( + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979171, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + 1125979170, + ), + ), + ( + (-1198.590576171875, -0.01978362910449505, 12.744208335876465, -956837.625), + ( + 3379160183, + 3379160183, + 3379160183, + 3379160184, + 3379160184, + 3379160183, + 3379160184, + 3379160184, + 3379160184, + 3379160183, + 3379160184, + 3379160183, + 3379160183, + 3379160184, + 3379160183, + ), + ), + ( + ( + 4.6832619204906223e-07, + -0.45981907844543457, + -60.8045768737793, + -15.982765197753906, + ), + ( + 3264904843, + 3264904844, + 3264904844, + 3264904843, + 3264904844, + 3264904844, + 3264904844, + 3264904844, + 3264904844, + 3264904844, + 3264904844, + 3264904844, + 3264904843, + 3264904843, + 3264904843, + ), + ), + ( + (-185.40750122070312, -7355930.0, 2.246054172515869, 54945.8671875), + ( + 3403599967, + 3403599967, + 3403599967, + 3403599967, + 3403599966, + 3403599967, + 3403599967, + 3403599967, + 3403599966, + 3403599967, + 3403599967, + 3403599967, + 3403599967, + 3403599967, + 3403599967, + ), + ), + ( + ( + -2.810704131661623e-07, + -1.8851414651521736e-08, + 1.4957002303361833e-10, + 1.1888499784618034e-07, + ), + ( + 3024239082, + 3024239082, + 3024239082, + 3024239081, + 3024239080, + 3024239081, + 3024239082, + 3024239080, + 3024239080, + 3024239081, + 3024239081, + 3024239081, + 3024239081, + 3024239080, + 3024239081, + ), + ), + ( + ( + 125.35652160644531, + -16.128854751586914, + -30.964479446411133, + -1.538109358989459e-06, + ), + ( + 1117554368, + 1117554368, + 1117554368, + 1117554368, + 1117554368, + 1117554368, + 1117554368, + 1117554369, + 1117554369, + 1117554368, + 1117554369, + 1117554369, + 1117554369, + 1117554369, + 1117554369, + ), + ), + ( + ( + -10.238574981689453, + 7.449639797210693, + 0.18453934788703918, + 0.0034961116034537554, + ), + ( + 3223745828, + 3223745828, + 3223745828, + 3223745828, + 3223745826, + 3223745828, + 3223745828, + 3223745828, + 3223745826, + 3223745828, + 3223745828, + 3223745828, + 3223745830, + 3223745826, + 3223745826, + ), + ), +) diff --git a/vllm/model_executor/layers/sm70_collectives.py b/vllm/model_executor/layers/sm70_collectives.py new file mode 100644 index 0000000000..ee4f0b7961 --- /dev/null +++ b/vllm/model_executor/layers/sm70_collectives.py @@ -0,0 +1,265 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Explicit, calibrated FP32 local-row reduction for SM70 TP4 callers. + +This experimental interface has no automatic dispatch. Prepare collectively +outside CUDA graphs, use one bound stream, consume each returned view before +calling again, and close collectively before destroying the process group. +Calibration depends on the communicator and shape, never on model values. +""" + +import sys +from functools import lru_cache +from importlib import import_module +from operator import index +from pathlib import Path +from typing import Any + +import torch +import torch.distributed as dist + + +def _layout(shape): + if len(shape) != 2 or any(isinstance(value, bool) for value in shape): + raise ValueError("Exact row reduction requires two integer dimensions") + rows, columns = map(index, shape) + if rows <= 0 or columns <= 0 or rows % 4: + raise ValueError("Exact row reduction requires positive TP4-aligned rows") + count = rows * columns + if count > (2**63 - 1) // 4: + raise ValueError("Exact row reduction allocation size overflows int64") + local = count // 4 + raw = count * 4 + 2 * 80 * 4 * 4 + resident = raw + local * 5 + calibration_peak = resident + count * 8 + local * 2 + 256 + return (rows, columns), raw, resident, calibration_peak + + +@lru_cache(maxsize=1) +def _extension(): + try: + return import_module("vllm._sm70_exact_reduce_C") + except ImportError: + from torch.utils.cpp_extension import load + + root = Path(__file__).resolve().parents[3] + source = root / "csrc/sm70_turbomind/ops/exact_row_reduce.cu" + if not source.is_file(): + raise RuntimeError( + "SM70 exact row reduction requires the source build" + ) from None + extension = load( + name="onecat_sm70_exact_reduce", + sources=[str(source)], + extra_cuda_cflags=[ + "-O3", + "--fmad=false", + "-gencode=arch=compute_70,code=sm_70", + ], + verbose=False, + ) + sys.modules["onecat_sm70_exact_reduce"] = extension + return extension + + +class SM70ExactRowReductionPlan: + """Own one shape's IPC buffers and calibrated native FP32 addition order. + + ``group`` is a vLLM TP group with four ranks and a CPU process group. + ``memory_budget_bytes`` must cover persistent buffers and calibration + scratch. Callers retain their ordinary collective when this explicit plan + is unsuitable. Returned tensors alias plan storage. CUDA graphs and use + from another stream/device are rejected rather than silently changing + synchronization semantics. + """ + + def __init__(self, group, shape, *, memory_budget_bytes): + error = None + try: + self.shape, self.raw_ipc_bytes, self.resident_bytes, peak = _layout(shape) + budget = index(memory_budget_bytes) + if isinstance(memory_budget_bytes, bool) or budget < peak: + error = "Exact row reduction exceeds the explicit memory budget" + except (TypeError, ValueError, OverflowError) as exc: + self.shape, peak = None, 0 + error = str(exc) + self.calibration_peak_bytes = peak + self.group = group + self.rank = group.rank_in_group + self.device = torch.accelerator.current_device_index() + self.stream = torch.cuda.current_stream(self.device).cuda_stream + self._closed = False + self._buffers = [] + self._epoch = 0 + self.calls = 0 + properties = torch.cuda.get_device_properties(self.device) + if group.world_size != 4 or not 0 <= self.rank < 4: + error = "Exact row reduction requires TP4" + elif (properties.major, properties.minor) != (7, 0): + error = "Exact row reduction requires SM70" + elif properties.multi_processor_count < 80: + error = "Exact row reduction requires at least 80 SMs" + elif torch.cuda.is_current_stream_capturing(): + error = "Prepare exact row reduction outside CUDA graphs" + metadata: list[Any] = [None] * group.world_size + with torch.inference_mode(False): + dist.all_gather_object( + metadata, + (self.shape, self.device, error, str(properties.uuid)), + group=group.cpu_group, + ) + errors = [item[2] for item in metadata if item[2]] + if errors or any(item[0] != self.shape for item in metadata): + raise ValueError(errors or "Ranks requested different reduction shapes") + if len({item[3] for item in metadata}) != 4 or any( + not 0 <= item[1] < torch.accelerator.device_count() + or str(torch.cuda.get_device_properties(item[1]).uuid) != item[3] + for item in metadata + ): + error = ( + "Exact row reduction requires one host with " + "consistent CUDA device visibility" + ) + elif any( + peer != self.rank + and not torch.cuda.can_device_access_peer(self.device, item[1]) + for peer, item in enumerate(metadata) + ): + error = "Exact row reduction requires peer access to every TP rank" + self._agree(error) + self.ops = _extension() + self.output = None + self.codes = None + try: + with torch.inference_mode(False): + count = self.shape[0] * self.shape[1] + self.pointers = self._shared(count * 4) + self.flags = self._shared(2 * 80 * 4 * 4) + self.output = torch.empty( + (self.shape[0] // 4, self.shape[1]), + device=self.device, + dtype=torch.float32, + ) + self.codes = self._calibrate(count) + dist.barrier(group=group.cpu_group) + except BaseException: + self.close() + raise + + def _agree(self, error): + errors = [None] * self.group.world_size + with torch.inference_mode(False): + dist.all_gather_object(errors, error, group=self.group.cpu_group) + if any(errors): + raise RuntimeError(f"Exact row reduction setup failed: {errors}") + + def _shared(self, size): + pointer, handle, error = 0, None, None + try: + pointer, handle = self.ops.allocate(size) + torch.accelerator.synchronize() + except RuntimeError as exc: + error = str(exc) + handles: list[Any] = [None] * 4 + dist.all_gather_object(handles, (handle, error), group=self.group.cpu_group) + if any(item[1] for item in handles): + if pointer: + self.ops.release(pointer, True) + raise RuntimeError(f"Exact row reduction IPC allocation failed: {handles}") + pointers = [0] * 4 + pointers[self.rank] = pointer + self._buffers.append(pointers) + for peer, (handle, _) in enumerate(handles): + if peer != self.rank: + try: + pointers[peer] = self.ops.open_handle(handle) + except RuntimeError as exc: + error = str(exc) + break + self._agree(error) + return pointers + + def _calibrate(self, count): + from vllm.triton_utils import triton + + from .sm70_collective_calibration import PROBES, decode_mask, update_mask + + n = count // 4 + masks = torch.full((n,), 0x7FFF, device=self.device, dtype=torch.int16) + codes = torch.empty(n, device=self.device, dtype=torch.uint8) + stats = torch.zeros(2, device=self.device, dtype=torch.int64) + value = torch.empty(self.shape, device=self.device, dtype=torch.float32) + for inputs, expected_bits in PROBES: + value.fill_(inputs[self.rank]) + reference = self.group.all_reduce(value) + expected = torch.tensor( + expected_bits, device=self.device, dtype=torch.uint32 + ) + update_mask[(triton.cdiv(n, 256),)]( + reference, masks, expected, n, n * self.rank, 256 + ) + del reference, expected + decode_mask[(triton.cdiv(n, 256),)](masks, codes, stats, n, 256) + self._agree( + None + if stats.tolist() == [0, 0] + else "Native FP32 addition order cannot be classified uniquely" + ) + return codes + + def reduce(self, value): + """Return this rank's local rows; adapters must already be included.""" + if self._closed: + raise RuntimeError("Exact row reduction plan is closed") + if ( + not value.is_cuda + or value.device.index != self.device + or value.dtype != torch.float32 + or tuple(value.shape) != self.shape + or not value.is_contiguous() + or value.requires_grad + ): + raise ValueError("Exact row reduction requires the prepared FP32 layout") + if ( + torch.accelerator.current_device_index() != self.device + or torch.cuda.current_stream(self.device).cuda_stream != self.stream + or torch.cuda.is_current_stream_capturing() + ): + raise RuntimeError( + "Exact row reduction requires its original uncaptured stream" + ) + if self._epoch == 2**32 - 1: + raise RuntimeError( + "Exact row reduction epoch exhausted; prepare a new plan" + ) + self._epoch += 1 + self.calls += 1 + self.ops.run( + value, + self.codes, + self.output, + self.pointers, + self.flags, + self.rank, + self._epoch, + ) + return self.output + + def close(self): + """Collectively release peer handles before freeing their owners.""" + if self._closed: + return + with torch.accelerator.device_index(self.device): + torch.accelerator.synchronize() + dist.barrier(group=self.group.cpu_group) + for pointers in self._buffers: + for peer, pointer in enumerate(pointers): + if pointer and peer != self.rank: + self.ops.release(pointer, False) + dist.barrier(group=self.group.cpu_group) + for pointers in self._buffers: + if pointers[self.rank]: + self.ops.release(pointers[self.rank], True) + self._buffers.clear() + self.output = self.codes = None + self._closed = True diff --git a/vllm/video/metrics.py b/vllm/video/metrics.py index 18eb769e26..ea4ee0650c 100644 --- a/vllm/video/metrics.py +++ b/vllm/video/metrics.py @@ -21,7 +21,7 @@ def loaded_kernel_provenance(): paths = { str(Path(filename).resolve()) for name, module in list(sys.modules.items()) - if name.startswith(("vllm._h3_", "onecat_h3_")) + if name.startswith(("vllm._h3_", "onecat_h3_", "vllm._sm70_", "onecat_sm70_")) and (filename := getattr(module, "__file__", None)) } return { From 6951ff5260e8538495c8cfd5a9fe53426e3f2c9f Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 21:44:17 +0800 Subject: [PATCH 16/25] [Doc] Record native media preservation for shared row reduction Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 5 +++-- docs/design/minimax_h3/EXACT_ROW_REDUCTION.md | 18 +++++++++++++----- 2 files changed, 16 insertions(+), 7 deletions(-) diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index 1e7b1e0916..0b0154d9e1 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -56,5 +56,6 @@ work continues separately; it has no accepted model-level speedup yet. The explicit [shared row-reduction interface](EXACT_ROW_REDUCTION.md) has separate operator and prototype full-media controls. Its 2.49348% paired denoise improvement is a development measurement, not a new formal campaign -result. H3 runtime integration and final-interface media validation remain -pending, and the ordinary reduction stays selected. +result. The final shared interface also passes a complete native media control via +an explicit forward override. Native H3 runtime integration remains pending, +and the ordinary reduction stays selected. diff --git a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md index 94294c7765..399d64572a 100644 --- a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md +++ b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md @@ -74,11 +74,19 @@ V100 SXM2 32GB cards. Evidence root: 743,180,800 persistent raw IPC bytes/card. Their sum is 20,475,735,040 bytes; driver/library overhead is additional. Do not report only the PyTorch number. -The prototype full-model evidence precedes the packaged plan's dynamic -calibration and setup guards. It is not substituted for full-model validation -of this final interface. Native integration, finalized-interface media controls, -formal repeated requests, TP/shape breadth and official/human quality gates -remain incomplete. No AUTO promotion is made. +The committed shared interface at `6d2a44b8d0` now also passes a complete +native H3 control using an explicit forward override. Dynamic calibration uses +the current native group and the actual shape, with no saved arithmetic-code +map. Final video/audio latents, all 124 RGB frames and PCM match the frozen +FA query-128 control bitwise; SSIM and RMS ratio are 1. The captured request +records 58.310740 seconds denoise and 99.562953 seconds total. Its contract, +source/binary manifests, `review-native-quality.json` and +`review-native-summary.json` are retained separately from prototype evidence. + +This validates the final shared operator in one H3 request, but does not add +native model routing or establish full-request warmup-plus-three performance. +API integration, formal repeated requests, TP/shape breadth and official/human +quality gates remain incomplete. No AUTO promotion is made. Reproduce the operator control with an owned native GPU lease and `torchrun --standalone --nproc_per_node=4 From ca82c279bca00664fcb10881bcceaa913f539846 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 21:54:05 +0800 Subject: [PATCH 17/25] [Core] Expose budgeted native H3 residual reduction Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 4 +- docs/design/minimax_h3/EXACT_ROW_REDUCTION.md | 40 +++++++-- tests/video/test_h3_residual_collectives.py | 87 +++++++++++++++++++ vllm/entrypoints/cli/video.py | 14 +++ .../model_executor/layers/sm70_collectives.py | 5 ++ .../models/minimax_h3/collectives.py | 66 ++++++++++++++ .../models/minimax_h3/config.py | 13 +++ .../models/minimax_h3/pipeline.py | 27 +++++- .../models/minimax_h3/transformer.py | 21 +++-- vllm/video/engine.py | 18 +++- 10 files changed, 277 insertions(+), 18 deletions(-) create mode 100644 tests/video/test_h3_residual_collectives.py create mode 100644 vllm/model_executor/models/minimax_h3/collectives.py diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index 0b0154d9e1..7a1d048312 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -57,5 +57,5 @@ The explicit [shared row-reduction interface](EXACT_ROW_REDUCTION.md) has separate operator and prototype full-media controls. Its 2.49348% paired denoise improvement is a development measurement, not a new formal campaign result. The final shared interface also passes a complete native media control via -an explicit forward override. Native H3 runtime integration remains pending, -and the ordinary reduction stays selected. +an explicit forward override. Native H3 selection is now explicit and still requires its final GPU +validation; the ordinary reduction remains the default. diff --git a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md index 399d64572a..9b99ae8bc8 100644 --- a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md +++ b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md @@ -1,7 +1,8 @@ # Explicit SM70 local-row reduction The shared `SM70ExactRowReductionPlan` interface is experimental and has no -automatic dispatch or H3 runtime selection yet. The ordinary residual path +automatic dispatch. H3 exposes an explicit runtime selection; its final +integration is still undergoing GPU and full-request validation. The ordinary residual path continues to use FP32 all-reduce followed by a local-row slice. No configuration has passed the campaign's >80 useful TFLOP/s/card and complete quality gates. @@ -32,8 +33,9 @@ CPU group, and peer handles close before owners free their allocations. The plan rejects another stream/device, autograd inputs, incompatible layouts, CUDA Graph execution and epoch exhaustion. It does not silently change a backend or precision. Callers retain their ordinary collective when a plan -is unsuitable. Only the shared operator is provided in this change; a model -must explicitly own its lifecycle before integrating it into requests. +is unsuitable. H3 owns one plan per pipeline, reuses it for identical shapes and closes it +collectively on shape changes or worker shutdown. Other models must likewise +own the plan lifecycle explicitly. ## Validation and measured limits @@ -83,9 +85,8 @@ records 58.310740 seconds denoise and 99.562953 seconds total. Its contract, source/binary manifests, `review-native-quality.json` and `review-native-summary.json` are retained separately from prototype evidence. -This validates the final shared operator in one H3 request, but does not add -native model routing or establish full-request warmup-plus-three performance. -API integration, formal repeated requests, TP/shape breadth and official/human +This validates the final shared operator in one H3 request, but does not establish full-request warmup-plus-three performance. +Final native integration validation, formal repeated requests, TP/shape breadth and official/human quality gates remain incomplete. No AUTO promotion is made. Reproduce the operator control with an owned native GPU lease and @@ -93,3 +94,30 @@ Reproduce the operator control with an owned native GPU lease and benchmarks/kernels/benchmark_sm70_exact_row_reduce.py --output --full-shape`. The optional `--extension` pins an already-built library; the report records its SHA256 plus benchmark, CUDA and shared Python source hashes. + +## Native API selection and memory records + +Start `vllm video serve` or `vllm video generate` with +`--residual-sequence-parallel --residual-reduction peer +--residual-reduction-memory-gib 4` to select the explicit candidate. The default +remains `native`; HTTP clients continue to use the existing video API. Selection +does not depend on floating versus W8A16 weights, adapters or task labels. + +TP1 retains its ordinary path. TP2 uses the original all-reduce and local slice. +TP4 creates a shared plan only when its complete calibration/storage requirement +fits the explicit budget; larger shapes use the ordinary collective. A plan +remains valid only for its original communicator and shape. Setup happens on +the first actual projection; its time is included in complete denoise and is +also reported separately in `residual_communication.setup_seconds`. Repeated +requests of the same shape reuse calibration without reading model state. + +Each result records peer/native call counts and the fallback reason. The +`torch_peak_allocated_bytes` and `raw_ipc_peak_bytes` fields remain separate; +`peak_allocated_bytes` is their conservative sum and is marked as an upper +bound when raw IPC storage is present. The performance validator therefore +includes external communication allocations in its memory gate. CUDA driver +and library overhead still require the retained NVML measurements. + +CPU request ownership, budget fallback, config, API and existing residual +regressions pass. GPU generation through the final native selection and formal +measurements remain required; preceding forward overrides do not satisfy them. diff --git a/tests/video/test_h3_residual_collectives.py b/tests/video/test_h3_residual_collectives.py new file mode 100644 index 0000000000..33f1527839 --- /dev/null +++ b/tests/video/test_h3_residual_collectives.py @@ -0,0 +1,87 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +from types import SimpleNamespace + +import pytest +import torch + +from vllm.model_executor.models.minimax_h3 import collectives +from vllm.model_executor.models.minimax_h3.config import H3Config, H3InputError + + +@pytest.mark.parametrize("tp", [1, 2, 4]) +@pytest.mark.parametrize("partition", ["fl2va", "ref2va"]) +def test_explicit_peer_option_keeps_tp_and_partition_compatibility(tp, partition): + config = H3Config( + tensor_parallel_size=tp, + partition=partition, + residual_sequence_parallel=True, + residual_reduction="peer", + ) + assert config.residual_reduction == "peer" + assert H3Config().residual_reduction == "native" + + +@pytest.mark.parametrize("budget", [0, -1, True, float("nan"), float("inf"), 1e308]) +def test_reject_invalid_communication_budget(budget): + with pytest.raises(H3InputError): + H3Config(residual_reduction_memory_gib=budget) + + +def test_peer_option_requires_residual_rows(): + with pytest.raises(H3InputError, match="residual sequence"): + H3Config(residual_reduction="peer") + + +def test_reuse_request_accounting_shape_eviction_and_budget_fallback(monkeypatch): + plans = [] + + class Plan: + raw_ipc_bytes = 64 + + @staticmethod + def required_memory_bytes(shape): + return shape[0] * shape[1] * 4 + + def __init__(self, group, shape, *, memory_budget_bytes): + self.shape = shape + self.closed = False + plans.append(self) + + def reduce(self, value): + return (value * 4).chunk(4)[1] + + def close(self): + self.closed = True + + monkeypatch.setattr(collectives, "SM70ExactRowReductionPlan", Plan) + group = SimpleNamespace(world_size=4, rank_in_group=1, all_reduce=lambda x: x * 4) + owner = collectives.H3ResidualReduction(group, memory_budget_bytes=100) + value = torch.arange(12, dtype=torch.float32).view(4, 3) + expected = (value * 4).chunk(4)[1] + assert torch.equal(owner.reduce(value), expected) + owner.begin_request() + assert owner.snapshot()["peer_calls"] == 0 + assert owner.snapshot()["raw_ipc_peak_bytes"] == 64 + assert torch.equal(owner.reduce(value), expected) + assert len(plans) == 1 + assert owner.snapshot()["setup_seconds"] == 0 + large = torch.ones(16, 3) + assert torch.equal(owner.reduce(large), torch.full((4, 3), 4.0)) + assert plans[0].closed + assert owner.snapshot()["native_calls"] == 1 + assert owner.snapshot()["fallback_reason"] + owner.begin_request() + assert owner.snapshot()["raw_ipc_peak_bytes"] == 0 + owner.close() + + +def test_tp2_uses_ordinary_reduction_without_plan(monkeypatch): + def unexpected(*args, **kwargs): + raise AssertionError("TP2 must not construct a TP4 peer plan") + + monkeypatch.setattr(collectives, "SM70ExactRowReductionPlan", unexpected) + group = SimpleNamespace(world_size=2, rank_in_group=1, all_reduce=lambda x: x * 2) + owner = collectives.H3ResidualReduction(group, memory_budget_bytes=100) + assert torch.equal(owner.reduce(torch.ones(6, 3)), torch.full((3, 3), 2.0)) + assert owner.snapshot()["native_calls"] == 1 diff --git a/vllm/entrypoints/cli/video.py b/vllm/entrypoints/cli/video.py index 4253330577..6408b53452 100644 --- a/vllm/entrypoints/cli/video.py +++ b/vllm/entrypoints/cli/video.py @@ -72,6 +72,18 @@ def subparser_init(self, subparsers): action="store_true", help="Experimental FP32 residual sharding for TP2/TP4; TP1 is a no-op", ) + mode.add_argument( + "--residual-reduction", + choices=("native", "peer"), + default="native", + help="Explicit TP4 SM70 row reduction; requires residual sharding", + ) + mode.add_argument( + "--residual-reduction-memory-gib", + type=float, + default=4.0, + help="Communication setup and buffer budget; larger shapes use native", + ) mode.add_argument("--output-dir", type=Path, default=Path("h3-output")) mode.add_argument( "--video-encoder", @@ -135,6 +147,8 @@ def cmd(args): int8_weight_layout=args.int8_weight_layout, fp16_weight_layout=args.fp16_weight_layout, residual_sequence_parallel=args.residual_sequence_parallel, + residual_reduction=args.residual_reduction, + residual_reduction_memory_gib=args.residual_reduction_memory_gib, video_encoder=args.video_encoder, host_weight_pin_memory=args.host_weight_pin_memory, share_host_vae_weights=args.share_host_vae_weights, diff --git a/vllm/model_executor/layers/sm70_collectives.py b/vllm/model_executor/layers/sm70_collectives.py index ee4f0b7961..ba47473030 100644 --- a/vllm/model_executor/layers/sm70_collectives.py +++ b/vllm/model_executor/layers/sm70_collectives.py @@ -73,6 +73,11 @@ class SM70ExactRowReductionPlan: synchronization semantics. """ + @staticmethod + def required_memory_bytes(shape): + """Conservative explicit budget for setup scratch and persistent buffers.""" + return _layout(shape)[3] + def __init__(self, group, shape, *, memory_budget_bytes): error = None try: diff --git a/vllm/model_executor/models/minimax_h3/collectives.py b/vllm/model_executor/models/minimax_h3/collectives.py new file mode 100644 index 0000000000..9ecb0d9d63 --- /dev/null +++ b/vllm/model_executor/models/minimax_h3/collectives.py @@ -0,0 +1,66 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Pipeline ownership and request accounting for explicit residual collectives.""" + +import time + +from vllm.model_executor.layers.sm70_collectives import SM70ExactRowReductionPlan + + +class H3ResidualReduction: + def __init__(self, group, *, memory_budget_bytes): + self.group = group + self.memory_budget_bytes = memory_budget_bytes + self.plan: SM70ExactRowReductionPlan | None = None + self.begin_request() + + def begin_request(self): + self.peer_calls = 0 + self.native_calls = 0 + self.setup_seconds = 0.0 + self.raw_peak_bytes = self.plan.raw_ipc_bytes if self.plan is not None else 0 + self.fallback_reason = None + + def reduce(self, value): + shape = tuple(value.shape) + if self.plan is not None and self.plan.shape != shape: + self.plan.close() + self.plan = None + if self.group.world_size != 4: + self.fallback_reason = "peer execution requires TP4" + elif ( + SM70ExactRowReductionPlan.required_memory_bytes(shape) + > self.memory_budget_bytes + ): + self.fallback_reason = "shape exceeds residual communication budget" + else: + if self.plan is None: + started = time.perf_counter() + self.plan = SM70ExactRowReductionPlan( + self.group, shape, memory_budget_bytes=self.memory_budget_bytes + ) + self.setup_seconds += time.perf_counter() - started + self.raw_peak_bytes = max(self.raw_peak_bytes, self.plan.raw_ipc_bytes) + self.peer_calls += 1 + return self.plan.reduce(value) + self.native_calls += 1 + rows = value.shape[0] // self.group.world_size + return self.group.all_reduce(value).narrow( + 0, self.group.rank_in_group * rows, rows + ) + + def snapshot(self): + return { + "configured_backend": "peer", + "peer_calls": self.peer_calls, + "native_calls": self.native_calls, + "fallback_reason": self.fallback_reason, + "setup_seconds": self.setup_seconds, + "raw_ipc_peak_bytes": self.raw_peak_bytes, + "memory_budget_bytes": self.memory_budget_bytes, + } + + def close(self): + if self.plan is not None: + self.plan.close() + self.plan = None diff --git a/vllm/model_executor/models/minimax_h3/config.py b/vllm/model_executor/models/minimax_h3/config.py index 3d28b16f13..d3d17a480a 100644 --- a/vllm/model_executor/models/minimax_h3/config.py +++ b/vllm/model_executor/models/minimax_h3/config.py @@ -45,12 +45,25 @@ class H3Config: int8_weight_layout: str = "column" fp16_weight_layout: Literal["row", "column"] = "row" residual_sequence_parallel: bool = False + residual_reduction: Literal["native", "peer"] = "native" + residual_reduction_memory_gib: float = 4.0 host_weight_pin_memory: bool = True share_host_vae_weights: bool = False weight_offload: Literal["component", "layer"] = "component" video_encoder: Literal["libx264", "h264_nvenc"] = "libx264" def __post_init__(self) -> None: + if self.residual_reduction not in ("native", "peer"): + raise H3InputError("residual reduction must be native or peer") + if self.residual_reduction == "peer" and not self.residual_sequence_parallel: + raise H3InputError("peer reduction requires residual sequence parallelism") + if ( + isinstance(self.residual_reduction_memory_gib, bool) + or not math.isfinite(self.residual_reduction_memory_gib) + or not math.isfinite(self.residual_reduction_memory_gib * 2**30) + or self.residual_reduction_memory_gib <= 0 + ): + raise H3InputError("residual communication budget must be finite and > 0") if self.attention_query_tile not in (64, 128): raise H3InputError("Attention query tile must be 64 or 128") if ( diff --git a/vllm/model_executor/models/minimax_h3/pipeline.py b/vllm/model_executor/models/minimax_h3/pipeline.py index fa0da0270b..31b12a7f28 100644 --- a/vllm/model_executor/models/minimax_h3/pipeline.py +++ b/vllm/model_executor/models/minimax_h3/pipeline.py @@ -88,7 +88,7 @@ minimax_h3_align_frame_count, minimax_h3_time_shift_sigmas, ) -from .transformer import MiniMaxH3DiTModel +from .transformer import MiniMaxH3DiTBlock, MiniMaxH3DiTModel from .vae import MiniMaxH3AudioVAE, MiniMaxH3VideoVAE from .weight_cache import FP16WeightCache from .weights import iter_checkpoint_weights, resolve_model_root @@ -508,6 +508,17 @@ def __init__(self, config: H3Config, *, shared_weights_dir: str | None = None): for module in self.transformer.modules(): if isinstance(module, Attention): module.query_tile = config.attention_query_tile + self._residual_reduction = None + if config.residual_reduction == "peer": + from .collectives import H3ResidualReduction + + self._residual_reduction = H3ResidualReduction( + get_tp_group(), + memory_budget_bytes=int(config.residual_reduction_memory_gib * 2**30), + ) + for module in self.transformer.modules(): + if isinstance(module, MiniMaxH3DiTBlock): + module.residual_reducer = self._residual_reduction weights = iter_checkpoint_weights(transformer_path) if restore_adaln: weights = restore_dense_adaln_weights(weights, path / "transformer") @@ -688,9 +699,23 @@ def _resident_dit_layers_on_device(self, *, enabled=True): def progress_bar(self, *, total): return tqdm(total=total, desc="H3 denoise", disable=self._dit_rank != 0) + def residual_reduction_stats(self): + reducer = getattr(self, "_residual_reduction", None) + if reducer is None: + return {"configured_backend": "native", "raw_ipc_peak_bytes": 0} + return reducer.snapshot() + + def close(self): + reducer = getattr(self, "_residual_reduction", None) + if reducer is not None: + reducer.close() + @torch.inference_mode() def forward(self, request: H3Request): self.stage_durations = {} + reducer = getattr(self, "_residual_reduction", None) + if reducer is not None: + reducer.begin_request() self.actual_dit_calls = 0 started = time.perf_counter() context = self._prepare_request_inputs( diff --git a/vllm/model_executor/models/minimax_h3/transformer.py b/vllm/model_executor/models/minimax_h3/transformer.py index e1d7961265..f716ef3e3c 100644 --- a/vllm/model_executor/models/minimax_h3/transformer.py +++ b/vllm/model_executor/models/minimax_h3/transformer.py @@ -53,6 +53,8 @@ QuantizationConfig, ) + from .collectives import H3ResidualReduction + logger = init_logger(__name__) @@ -824,6 +826,7 @@ def __init__( quant_config, prefix=f"{prefix}.mlp", ) + self.residual_reducer: H3ResidualReduction | None = None self.residual_group = ( get_tp_group() if residual_sequence_parallel and get_tensor_model_parallel_world_size() > 1 @@ -921,9 +924,12 @@ def forward( input_is_rotated=input_is_rotated, ) if group is not None: - h = group.all_reduce(h).narrow( - 0, group.rank_in_group * residual.shape[0], residual.shape[0] - ) + if self.residual_reducer is not None: + h = self.residual_reducer.reduce(h) + else: + h = group.all_reduce(h).narrow( + 0, group.rank_in_group * residual.shape[0], residual.shape[0] + ) x, h = indexed_gate_rms_norm_scale_shift( residual, gate_msa, @@ -942,9 +948,12 @@ def forward( h = group.all_gather(h, dim=0) h = self.mlp(h, input_is_rotated=input_is_rotated) if group is not None: - h = group.all_reduce(h).narrow( - 0, group.rank_in_group * residual.shape[0], residual.shape[0] - ) + if self.residual_reducer is not None: + h = self.residual_reducer.reduce(h) + else: + h = group.all_reduce(h).narrow( + 0, group.rank_in_group * residual.shape[0], residual.shape[0] + ) return indexed_gate(residual, gate_mlp, h, combined_indices) diff --git a/vllm/video/engine.py b/vllm/video/engine.py index 881ac8be7e..abd7658391 100644 --- a/vllm/video/engine.py +++ b/vllm/video/engine.py @@ -45,6 +45,7 @@ def _worker(rank, config, gpu_ids, endpoint, connection, shared_weights_dir=None current = VllmConfig( parallel_config=ParallelConfig(tensor_parallel_size=config.tensor_parallel_size) ) + pipeline = None try: with set_current_vllm_config(current): init_distributed_environment( @@ -77,8 +78,15 @@ def _worker(rank, config, gpu_ids, endpoint, connection, shared_weights_dir=None from .metrics import loaded_kernel_provenance kernel_provenance = loaded_kernel_provenance() + communication = pipeline.residual_reduction_stats() + torch_peak = torch.accelerator.max_memory_allocated() + raw_peak = communication["raw_ipc_peak_bytes"] result = { "rank": rank, + "residual_communication": communication, + "torch_peak_allocated_bytes": torch_peak, + "raw_ipc_peak_bytes": raw_peak, + "peak_allocation_is_upper_bound": bool(raw_peak), "stage_seconds": pipeline.stage_durations, "dit_calls": pipeline.actual_dit_calls, "useful_denoise_flops": pipeline.useful_denoise_flops, @@ -89,7 +97,7 @@ def _worker(rank, config, gpu_ids, endpoint, connection, shared_weights_dir=None "denoise_steps": pipeline.denoise_steps, "denoise_executed_blocks": pipeline.denoise_executed_blocks, "kernel_provenance": kernel_provenance, - "peak_allocated_bytes": torch.accelerator.max_memory_allocated(), + "peak_allocated_bytes": torch_peak + raw_peak, } if rank == 0: from .media import export_video @@ -134,8 +142,12 @@ def _worker(rank, config, gpu_ids, endpoint, connection, shared_weights_dir=None except BaseException: connection.send({"error": traceback.format_exc(), "rank": rank}) finally: - cleanup_dist_env_and_memory() - connection.close() + try: + if pipeline is not None: + pipeline.close() + finally: + cleanup_dist_env_and_memory() + connection.close() class H3Engine: From 58cafcd05e840b28ec164a3f1eac468d80a1d17c Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 22:22:16 +0800 Subject: [PATCH 18/25] [Doc] Record formal native H3 peer reduction results Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 30 ++---------- docs/design/minimax_h3/EXACT_ROW_REDUCTION.md | 49 +++++++++++++++++-- 2 files changed, 50 insertions(+), 29 deletions(-) diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index 7a1d048312..9e653c64ad 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -22,6 +22,7 @@ All rows use TP4. Rows marked formal use a complete request warmup plus three un | FL2V Light4 v1.2_768p, first | W8A16 | FLASH_ATTN_V100 | 66.419 | 50.697 | passed | not measured | | FL2V Light4 v1.2_768p, last | W8A16 | FLASH_ATTN_V100 | 64.852 | 51.923 | passed | not measured | | FL2V Light4 v1.2_768p, first-last | W8A16 | FLASH_ATTN_V100 | 69.615 | 52.305 | passed | not measured | +| FL2V Light4 v1.2_768p, native peer rows | W8A16 | FLASH_ATTN_V100 | 58.293 | 53.236 | passed | failed >80 | The VSA failure is against an explicitly labeled FP32 selected-key diagnostic, not the unmodified official GPU kernel. Generation alone does not establish numerical quality. @@ -34,28 +35,7 @@ The VSA failure is against an explicitly labeled FP32 selected-key diagnostic, n Exact source/run paths and the evidence index are retained in `campaign-results.json` and `campaign-results.csv`. -## Additional bounded experiments - -These artifact-only experiments did not change production defaults: - -- Increasing NCCL CTA counts from 8 to 16/32 changes FP32 reduction bits. - The preliminary bitwise gate rejected them before timing or a full-model - comparison; this is not a measured full-model quality failure. -- Sequential FA warp-operand loading reduces one register count from 248 to - 235 but leaves the proposed occupancy budget unmet. Build evidence was - sufficient to reject the hypothesis; no GPU benchmark was run. -- K-only swizzling in the new FI kernel preserves operator bits but changes - the paired median by only 0.192%, within observed clock variation. No - complete-model run or production integration was justified. - -The artifact folders `nccl-cta-control`, `attention-single-warp-buffer` and -`attention-fi-key-only-swizzle` retain hypotheses, source hashes and results. -Do not repeat these experiments without a changed hypothesis. Communication -work continues separately; it has no accepted model-level speedup yet. - -The explicit [shared row-reduction interface](EXACT_ROW_REDUCTION.md) has -separate operator and prototype full-media controls. Its 2.49348% paired -denoise improvement is a development measurement, not a new formal campaign -result. The final shared interface also passes a complete native media control via -an explicit forward override. Native H3 selection is now explicit and still requires its final GPU -validation; the ordinary reduction remains the default. +The [native row-reduction record](EXACT_ROW_REDUCTION.md) includes explicit +API selection, native bitwise media preservation and the 53.236 TFLOP/s/card +formal result. Its host-memory policy differs from the older query-128 run, +so complete-request times do not isolate the communication change. diff --git a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md index 9b99ae8bc8..2a855a3c11 100644 --- a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md +++ b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md @@ -1,8 +1,9 @@ # Explicit SM70 local-row reduction The shared `SM70ExactRowReductionPlan` interface is experimental and has no -automatic dispatch. H3 exposes an explicit runtime selection; its final -integration is still undergoing GPU and full-request validation. The ordinary residual path +automatic dispatch. H3 exposes an explicit runtime selection; its native +four-step integration passes media preservation and remains below the +formal >80 throughput gate. The ordinary residual path continues to use FP32 all-reduce followed by a local-row slice. No configuration has passed the campaign's >80 useful TFLOP/s/card and complete quality gates. @@ -119,5 +120,45 @@ includes external communication allocations in its memory gate. CUDA driver and library overhead still require the retained NVML measurements. CPU request ownership, budget fallback, config, API and existing residual -regressions pass. GPU generation through the final native selection and formal -measurements remain required; preceding forward overrides do not satisfy them. +regressions pass. The final native selection now has the separate controls and measurements +below. Wider workflow validation remains incomplete. + +## Final native four-step measurements + +At source `ca82c279bc`, the ordinary engine selects peer rows through H3Config, +with no forward replacement. The separate captured native request preserves +both final latents, all 124 RGB frames and PCM bitwise. All four ranks report +400 peer calls, zero native fallbacks and about 0.374 seconds initial plan +setup. `peer-api-native-quality.json` and `peer-api-native-summary.json` retain +this control and its precise configuration. + +`peer-api-720p-three-runs/performance.json` records one complete request warmup +and three unprofiled, uncaptured requests of that same configuration: + +| Measurement | Value | +| --- | ---: | +| Warmup denoise | 59.698783 s | +| Measured denoise | 58.344130 / 58.293218 / 58.234342 s | +| Median useful TFLOP/s/card | 53.235745–53.235764 | +| Denoise coefficient of variation | 0.076959% | +| Complete request | 91.071940 / 90.560870 / 91.905029 s | +| Live allocation upper bound, including IPC | 20,475,227,136 bytes/card | +| Memory gate | Pass | +| >80 throughput gate | Fail | + +The companion `peer-api-formal-telemetry.json` retains 1,995 NVML samples at a +0.25-second interval. Across startup, warmup and measurement, the maximum +sampled device usage is 24,387,256,320 bytes, including allocator caches and +runtime overhead. High-utilization samples have median power of about +278–280 W/card. Telemetry is independent of the CUDA work counters. + +This formal request uses pageable host masters and shared VAE weights; the +older FA query-128 formal control used pinned masters. The latter's shorter +complete-request time must not be presented as a matched comparison of the +communication kernels. The isolated 2.49348% denoise comparison above used +matching host and compute settings. No overall request speedup is claimed +across the different host-memory policies. + +The interface remains explicit. The >80 target, official/human quality, wider +workflow and shape/TP matrix are still incomplete. Initial setup, skipped work, +raw IPC storage and slower end-to-end outcomes are retained in the records. From f12803d4579588001fc3ea029674a253313f97ee Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 22:58:51 +0800 Subject: [PATCH 19/25] [Doc] Record native peer backend and workload controls Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 8 +++--- docs/design/minimax_h3/EXACT_ROW_REDUCTION.md | 25 +++++++++++++++++++ 2 files changed, 28 insertions(+), 5 deletions(-) diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index 9e653c64ad..7113877ace 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -18,11 +18,14 @@ All rows use TP4. Rows marked formal use a complete request warmup plus three un | FlashGen four-step | original floating | FLASH_ATTN_V100 | 59.488 | 51.621 | passed | not measured | | FastH3 Dense data-free | original floating | FLASH_ATTN_V100 | 56.224 | 54.125 | passed | not measured | | FastH3 VSA data-free | original floating | FASTVIDEO_VSA | 37.387 | 45.268 | failed | not measured | +| FL2V Light4 v1.2_768p, original native peer rows | original floating | FLASH_ATTN_V100 | 59.324 | 52.312 | passed | not measured | +| Ref2V Light4 v0.1, mixed native peer rows | W8A16 | FLASH_ATTN_V100 | 181.261 | 53.544 | passed | not measured | | FL2V Light4 v1.2_768p, original floating | original floating | FLASH_ATTN_V100 | 66.366 | not measured | passed | not measured | | FL2V Light4 v1.2_768p, first | W8A16 | FLASH_ATTN_V100 | 66.419 | 50.697 | passed | not measured | | FL2V Light4 v1.2_768p, last | W8A16 | FLASH_ATTN_V100 | 64.852 | 51.923 | passed | not measured | | FL2V Light4 v1.2_768p, first-last | W8A16 | FLASH_ATTN_V100 | 69.615 | 52.305 | passed | not measured | | FL2V Light4 v1.2_768p, native peer rows | W8A16 | FLASH_ATTN_V100 | 58.293 | 53.236 | passed | failed >80 | +| FL2V Light4 v1.2_768p, register FI native peer rows | W8A16 | FLASHINFER_SM70 | 60.970 | 50.899 | passed | failed >80 | The VSA failure is against an explicitly labeled FP32 selected-key diagnostic, not the unmodified official GPU kernel. Generation alone does not establish numerical quality. @@ -34,8 +37,3 @@ The VSA failure is against an explicitly labeled FP32 selected-key diagnostic, n - TeaCache and Cache-DiT/SCM currently have separate small-shape lifecycle evidence, not primary >80 or official quality acceptance. Exact source/run paths and the evidence index are retained in `campaign-results.json` and `campaign-results.csv`. - -The [native row-reduction record](EXACT_ROW_REDUCTION.md) includes explicit -API selection, native bitwise media preservation and the 53.236 TFLOP/s/card -formal result. Its host-memory policy differs from the older query-128 run, -so complete-request times do not isolate the communication change. diff --git a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md index 2a855a3c11..2456f3c2db 100644 --- a/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md +++ b/docs/design/minimax_h3/EXACT_ROW_REDUCTION.md @@ -162,3 +162,28 @@ across the different host-memory policies. The interface remains explicit. The >80 target, official/human quality, wider workflow and shape/TP matrix are still incomplete. Initial setup, skipped work, raw IPC storage and slower end-to-end outcomes are retained in the records. + +## Native backend and workload breadth + +The same native API path with register-probability FI also passes a complete +latent/RGB/PCM bitwise control against its frozen FI baseline. Full-request +warmup plus three unprofiled measurements record denoise +60.986616 / 60.889535 / 60.969633 seconds, median +50.898828-50.898847 useful TFLOP/s/card and CV 0.069457%. Complete requests take +89.859034 / 89.086770 / 89.525699 seconds. The allocation upper bound including +raw IPC is 20,475,227,136 bytes/card. FA and FI peer runs share the pageable +host/shared VAE policy. Both fail the >80 gate. Evidence: +`peer-api-fi-720p-three-runs/performance.json`, +`peer-api-fi-native-quality.json` and `peer-api-fi-formal-telemetry.json`. + +Additional complete native controls preserve video/audio latents, all 124 RGB +frames and PCM bitwise with original floating Light4 weights and W8A16 Ref4 +mixed image/video/audio conditioning. Original Light4 records 59.323955 seconds +denoise and 22,022,771,200 bytes/card allocation upper bound. Ref4 records +181.260984 seconds and 21,572,765,184 bytes/card; its longer reference sequence +uses an explicit 8 GiB reduction budget and 1,491,864,064 raw IPC bytes/card. +Both record 400 peer calls with zero native fallbacks. These are captured cold +quality controls, not formal repeated performance measurements. Evidence: +`peer-api-breadth-summary.json` and both corresponding `*-quality.json` files. +Independent official quality, human review and the full task/weight/shape +matrix remain incomplete. From 44679e8ae855973e6ee41c69b739a1e57cedfad4 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:11:07 +0800 Subject: [PATCH 20/25] [Doc] Focus H3 attention optimization on FlashAttention Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index 7113877ace..736b9ed87c 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -2,6 +2,13 @@ No configuration has completed the >80 useful TFLOP/s/card, independent official quality and human-review gates. +Development now concentrates on FlashAttention-V100, as requested by the user. +With matched native peer-reduction and pageable host/shared VAE settings, FA +records 53.235745 useful TFLOP/s/card versus FI's 50.898828. FI retains its +validated implementation; additional FI optimization and exhaustive acceptance +are paused. Workflow coverage and quality gates continue on the FA/shared path. +Historical FI evidence remains below for reproducibility. + All rows use TP4. Rows marked formal use a complete request warmup plus three unprofiled requests. Other timings are captured cold diagnostics. Original floating-weight controls with legacy FLOP accounting omit throughput. | Workflow | Weights | Backend | Denoise seconds | Minimum card TFLOP/s | Native numerical control | Formal performance | From dcaad0c8d3b930d021fade8ee110f5f9a98ae7db Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:25:46 +0800 Subject: [PATCH 21/25] [Doc] Record FA bottlenecks and rejected operator candidates Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/FA_DEVELOPMENT.md | 62 ++++++++++++++++++++++++ 1 file changed, 62 insertions(+) create mode 100644 docs/design/minimax_h3/FA_DEVELOPMENT.md diff --git a/docs/design/minimax_h3/FA_DEVELOPMENT.md b/docs/design/minimax_h3/FA_DEVELOPMENT.md new file mode 100644 index 0000000000..c10406ff32 --- /dev/null +++ b/docs/design/minimax_h3/FA_DEVELOPMENT.md @@ -0,0 +1,62 @@ +# FlashAttention development focus + +Further Attention development and exhaustive workflow qualification concentrate +on FlashAttention-V100 at the user's request. Validated FlashInfer remains +available; new FI optimization and full FI acceptance are paused. + +With native TP4 peer rows, shared pageable VAE weights and a complete request +warmup plus three unprofiled requests, FA records a minimum-card median +53.235745 useful TFLOP/s and 58.293218 seconds denoise. The corresponding FI +measurements are 50.898828 TFLOP/s and 60.969633 seconds. Both pass native +latent/RGB/PCM preservation and fail the >80 throughput gate. Independent +official and human quality gates remain pending. + +## Evidence guiding the next operator change + +A development-only FA timer samples thread zero in head zero and every 32nd +query CTA. Nine boundary lengths and the actual [1,34551,14,128] captured +input remain bitwise equal to the retained FA implementation. Instrumented +operator median is 150.045700 ms versus 149.676025 ms without instrumentation +(0.247% overhead). Main-shape compilation uses 250 registers and no spills. +The query-64/key-128 diagnostic specialization spills and is not measured. + +| Sampled phase | Fraction of sampled warp spans | +| --- | ---: | +| QK including operand loads and barrier | 38.85% | +| V prefetch and softmax | 18.76% | +| Probability stores and publication | 7.94% | +| PV including operand loads and barrier | 33.20% | +| Iteration join | 1.25% | + +These phase spans include scheduling and waits; they are not whole-kernel +critical-path percentages or formal model performance. The measurement guides +operand and register-lifetime work, without qualifying a faster configuration. + +The primary 34,560-row GEMM diagnostic also compares all 18 returned eligible +zero-workspace, no-split, FP32-accumulation Lt choices across QKV, FC1, output +projection and FC2. Timed alternatives preserve outputs bitwise, and the +existing algorithm 21/tile 24 remains the fastest choice in each shape. +No GEMM plan change is justified by this measurement. + +## Rejected candidates + +- Normal-range hardware exp2 retains the library outside [-126,0]. Boundary + and actual-input outputs are bitwise, but paired median regresses from + 149.407745 to 158.803970 ms. No full-model run or source promotion follows. +- A new explicit WMMA Q16/K64-owner prototype retains FA K128 softmax panels, + FP32 accumulation and output scaling. It lowers register use from 248 to + 128 without spilling and doubles a Q128 CTA to 512 threads. Nine boundary + lengths and actual input are bitwise, but operator median regresses from + 147.507202 to 323.808258 ms. Register count alone is insufficient evidence + of a speedup. This is distinct from the old invalid CUTLASS warp16 shape. +- Adding vector Q/K loads, V lane-exchange transpose and swizzled probability + storage to that prototype spills 160 bytes per thread at the 128-register + limit. The resource gate rejects it before GPU timing. A separate scoped + staging experiment checks whether shortening vector lifetime can avoid + that spill; it remains an unqualified development artifact. + +Exact code, binary hashes, clocks, numerical results and paired measurements +are retained under `/data/minimax-h3/sm70-general-20260909/` in +`attention-fa-phase-clock`, `gemm-primary-heuristics`, +`attention-fa-normal-exp2`, `attention-fa-warp16-native` and +`attention-fa-warp16-staging`. None is installed as a production replacement. From 570be8d407644732f2b2ed66fc9fff78fc86de8f Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:26:51 +0800 Subject: [PATCH 22/25] [Doc] Close FA staging resource experiment Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/FA_DEVELOPMENT.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/design/minimax_h3/FA_DEVELOPMENT.md b/docs/design/minimax_h3/FA_DEVELOPMENT.md index c10406ff32..2a3a7a3971 100644 --- a/docs/design/minimax_h3/FA_DEVELOPMENT.md +++ b/docs/design/minimax_h3/FA_DEVELOPMENT.md @@ -52,8 +52,8 @@ No GEMM plan change is justified by this measurement. - Adding vector Q/K loads, V lane-exchange transpose and swizzled probability storage to that prototype spills 160 bytes per thread at the 128-register limit. The resource gate rejects it before GPU timing. A separate scoped - staging experiment checks whether shortening vector lifetime can avoid - that spill; it remains an unqualified development artifact. + staging experiment reduces the spill to 116 bytes with a 120-byte stack, + which still fails the resource gate. Neither staging variant is GPU timed. Exact code, binary hashes, clocks, numerical results and paired measurements are retained under `/data/minimax-h3/sm70-general-20260909/` in From e9dea1272195c39234a7c166844eff523931069e Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:58:38 +0800 Subject: [PATCH 23/25] [Doc] Align H3 control record with FA development focus Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CONTROL.md | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/docs/design/minimax_h3/CONTROL.md b/docs/design/minimax_h3/CONTROL.md index b875f18e60..c20483d863 100644 --- a/docs/design/minimax_h3/CONTROL.md +++ b/docs/design/minimax_h3/CONTROL.md @@ -1,5 +1,26 @@ # Native MiniMax H3 migration control +## Current campaign direction + +New Attention development and exhaustive workflow acceptance concentrate on +FlashAttention-V100 at the user's request. The already validated FlashInfer +implementation remains available, but new FI optimization and exhaustive FI +qualification are paused. The historical parallel-backend decisions below +are retained as evidence, not current work allocation. + +The full-request TP4 four-step native peer-row FA measurement records a +minimum-card median of 53.235745 useful TFLOP/s, 58.293218 seconds denoise, +and CV 0.076959%. It does not meet the >80 gate. Original floating weights, +W8A16, legal adapters/reference inputs, explicit VSA/cache algorithms and +independent official/human quality remain part of the campaign scope. +Only configurations satisfying the unchanged quality and full performance +gates may enter automatic selection. + +See [FA_DEVELOPMENT.md](FA_DEVELOPMENT.md) for current bottleneck evidence and +rejected kernel candidates, and [CAMPAIGN_RESULTS.md](CAMPAIGN_RESULTS.md) +for the per-workflow qualification table. The owned stack remains Draft; +local operator results are not end-to-end model acceptance. + Latest FlashInfer change: [FLASHINFER_LOCAL_ROTATION.md](FLASHINFER_LOCAL_ROTATION.md). Rotate FP16 rows on their owner before all-gather, avoiding duplicated QKV/MLP From 692dd344c21bdd035687c94bb29b4830d4f247b4 Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Thu, 10 Sep 2026 00:18:17 +0800 Subject: [PATCH 24/25] [Test] Record complete original H3 sampling preservation Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CAMPAIGN_RESULTS.md | 8 ++---- docs/design/minimax_h3/FA_DEVELOPMENT.md | 18 ++++++++++++++ docs/design/minimax_h3/SM70_EPILOGUES.md | 29 ++++++++++++++++++++++ 3 files changed, 49 insertions(+), 6 deletions(-) diff --git a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md index 736b9ed87c..64c8f4de7d 100644 --- a/docs/design/minimax_h3/CAMPAIGN_RESULTS.md +++ b/docs/design/minimax_h3/CAMPAIGN_RESULTS.md @@ -2,12 +2,7 @@ No configuration has completed the >80 useful TFLOP/s/card, independent official quality and human-review gates. -Development now concentrates on FlashAttention-V100, as requested by the user. -With matched native peer-reduction and pageable host/shared VAE settings, FA -records 53.235745 useful TFLOP/s/card versus FI's 50.898828. FI retains its -validated implementation; additional FI optimization and exhaustive acceptance -are paused. Workflow coverage and quality gates continue on the FA/shared path. -Historical FI evidence remains below for reproducibility. +Development concentrates on FlashAttention-V100. FI retains its validated implementation; further FI optimization and exhaustive acceptance are paused at the user request. Quality and workflow gates remain unchanged. All rows use TP4. Rows marked formal use a complete request warmup plus three unprofiled requests. Other timings are captured cold diagnostics. Original floating-weight controls with legacy FLOP accounting omit throughput. @@ -27,6 +22,7 @@ All rows use TP4. Rows marked formal use a complete request warmup plus three un | FastH3 VSA data-free | original floating | FASTVIDEO_VSA | 37.387 | 45.268 | failed | not measured | | FL2V Light4 v1.2_768p, original native peer rows | original floating | FLASH_ATTN_V100 | 59.324 | 52.312 | passed | not measured | | Ref2V Light4 v0.1, mixed native peer rows | W8A16 | FLASH_ATTN_V100 | 181.261 | 53.544 | passed | not measured | +| Base H3, no LoRA, 49 updates, native peer rows | original floating | FLASH_ATTN_V100 | 649.973 | 57.353 | passed | not measured | | FL2V Light4 v1.2_768p, original floating | original floating | FLASH_ATTN_V100 | 66.366 | not measured | passed | not measured | | FL2V Light4 v1.2_768p, first | W8A16 | FLASH_ATTN_V100 | 66.419 | 50.697 | passed | not measured | | FL2V Light4 v1.2_768p, last | W8A16 | FLASH_ATTN_V100 | 64.852 | 51.923 | passed | not measured | diff --git a/docs/design/minimax_h3/FA_DEVELOPMENT.md b/docs/design/minimax_h3/FA_DEVELOPMENT.md index 2a3a7a3971..31bd620366 100644 --- a/docs/design/minimax_h3/FA_DEVELOPMENT.md +++ b/docs/design/minimax_h3/FA_DEVELOPMENT.md @@ -60,3 +60,21 @@ are retained under `/data/minimax-h3/sm70-general-20260909/` in `attention-fa-phase-clock`, `gemm-primary-heuristics`, `attention-fa-normal-exp2`, `attention-fa-warp16-native` and `attention-fa-warp16-staging`. None is installed as a production replacement. + +The vector-staging Q96 follow-up raises the per-thread register budget to 168. +Its only 4-byte local spill is stored before the key loop and reloaded after +its final back edge, as verified in SASS. This admits a bounded numerical and +performance probe without claiming zero spills. Nine boundaries and the +actual 34,551-token input are bitwise, but median regresses from 148.462585 to +281.825287 ms. This closes the split-key warp16 staging route. + +A separate one-owner FA K128 prototype keeps rounded probability fragments +in registers, preserves the FA K64-half sum order and retains Q across key +panels. It uses 199 registers, no spills and 96 KiB shared memory. Nine +boundaries and the actual input are bitwise, but median is 195.892136 ms +versus 147.519104 ms. Its V-prefetch follow-up compiles to 240 registers with +no spills, but is not GPU timed: the parent exceeds the predeclared 10% +slowdown limit. No full-model run or production promotion follows. + +Evidence: `attention-fa-warp16-staging-q96`, `attention-fa-register-k128`, +`attention-fa-register-k128-vprefetch`, and `fa-vprefetch-after-register.json`. diff --git a/docs/design/minimax_h3/SM70_EPILOGUES.md b/docs/design/minimax_h3/SM70_EPILOGUES.md index 9cd9f7be1c..59c4f05b2b 100644 --- a/docs/design/minimax_h3/SM70_EPILOGUES.md +++ b/docs/design/minimax_h3/SM70_EPILOGUES.md @@ -341,3 +341,32 @@ Evidence: `original-variant-pairs.json`, `original-variant-summary.json`, contracts reside in the corresponding `*-original-{baseline,candidate}` runs. The [campaign table](CAMPAIGN_RESULTS.md) separates these diagnostic timings from formal acceptance measurements. Every configuration remains unqualified. + +## Original floating weights without an adapter + +The complete default-sampling T2VA native control uses original floating +weights, no LoRA, 50 sigma points and 49 actual updates. Both requests retain +seed 42, the same prompt and 1280x736/124-frame internal canvas. The frozen +`4f19ef7` control uses ordinary residuals and row-major floating weights; +the candidate at runtime source `570be8d407` uses prepared column-major +projections, FA query128 and explicit native peer rows. Shared host VAE +masters and pageable staging are recorded separately in the run contracts. + +All declared numerical gates pass: video/audio latents, all 124 pre-encoding +RGB frames and PCM are bitwise equal; SSIM and RMS ratio are 1. Denoise falls +from 725.819030 to 649.973214 seconds (10.449687%). Complete captured requests +are 835.355602 and 700.121402 seconds. Candidate actual-work validation passes +on all ranks and records 57.353041 useful TFLOP/s/card. Tracked allocation upper bounds +are 21,633,302,016 bytes/card for the baseline and 20,998,705,664 for the +candidate, including the candidate's raw IPC memory. + +These are single captured cold requests, not warmup-plus-three performance +acceptance. The combined result covers shared operators, residual layout and +host residency; it does not isolate an Attention-only or host-policy-only +speedup. Independent official reference, continuous human audiovisual review +and >80 acceptance remain incomplete. Four sampled baseline frames were +inspected, which does not replace those quality gates. + +Evidence: `original-base-pair.json`, `original-base-summary.json`, +`original-base-no-lora-original-quality.json` and the corresponding original +base run directories under the campaign artifact roots. From 9f220dc2fc27c90e561bb887ae7b0d7ca866910a Mon Sep 17 00:00:00 2001 From: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> Date: Thu, 10 Sep 2026 00:36:10 +0800 Subject: [PATCH 25/25] [Doc] Consolidate retained FA workflow delivery Assisted-by: OpenAI Codex Signed-off-by: yangzhuxinyzx <153831768+yangzhuxinyzx@users.noreply.github.com> --- docs/design/minimax_h3/CONTROL.md | 6 ++ docs/design/minimax_h3/CURRENT_STATUS.md | 93 ++++++++++++++++++++++++ docs/design/minimax_h3/FA_DEVELOPMENT.md | 30 +++++++- 3 files changed, 126 insertions(+), 3 deletions(-) create mode 100644 docs/design/minimax_h3/CURRENT_STATUS.md diff --git a/docs/design/minimax_h3/CONTROL.md b/docs/design/minimax_h3/CONTROL.md index c20483d863..bd90c0de78 100644 --- a/docs/design/minimax_h3/CONTROL.md +++ b/docs/design/minimax_h3/CONTROL.md @@ -2,6 +2,12 @@ ## Current campaign direction +The latest user instruction narrows current delivery to already working H3 +dense workflows using retained FA. Further slower Attention prototypes and +new workflow expansion are paused. No experimental replacement is installed. +See [CURRENT_STATUS.md](CURRENT_STATUS.md) for the current report. Historical +experiments and the unfinished full campaign remain below for reference. + New Attention development and exhaustive workflow acceptance concentrate on FlashAttention-V100 at the user's request. The already validated FlashInfer implementation remains available, but new FI optimization and exhaustive FI diff --git a/docs/design/minimax_h3/CURRENT_STATUS.md b/docs/design/minimax_h3/CURRENT_STATUS.md new file mode 100644 index 0000000000..f1db0d7ce3 --- /dev/null +++ b/docs/design/minimax_h3/CURRENT_STATUS.md @@ -0,0 +1,93 @@ +# H3 retained FA delivery status + +Current delivery concentrates on already working H3 dense workflows and the +retained `FLASH_ATTN_V100` implementation. This is already the native H3 +default backend. Slower experimental replacements are excluded. Further +Attention prototype work and new workflow expansion are paused at the user's +request; no GPU experiments remain queued. Native CLI/HTTP APIs remain the +frontend entry point. + +## Formal four-step result + +TP4, V100 SXM2 32GB, single request, 1280x736 internal canvas, 120 requested +frames aligned to 124, 24 FPS, LightX2V v1.2 four-step W8A16. Both backend +runs use the same pageable host/shared VAE policy and native peer reduction. +Each has one complete request warmup and three unprofiled measurements. + +| Metric | Retained FA | Retained FI | +| --- | ---: | ---: | +| Complete denoise median | 58.293218 s | 60.969633 s | +| Denoise divided by four updates | 14.573304 s | 15.242408 s | +| Minimum-card median useful TFLOP/s | 53.235745 | 50.898828 | +| Denoise CV | 0.076959% | 0.069457% | +| Complete request median | 91.071940 s | 89.525699 s | + +FA reduces denoise by 4.389750%. The complete request includes other stages +and is slower in this measurement; an overall FA request speedup is not +established. The tracked live-allocation upper bound including IPC is +20,475,227,136 bytes/card. All >80 gates remain incomplete. + +## Retained coverage + +These are complete native output-preservation controls. Except for the +explicit formal result above, the times below are captured cold diagnostics. +They are not repeated performance acceptance or independent official-model +quality acceptance. + +| Workflow | Weights | Denoise | Numerical control | +| --- | --- | ---: | --- | +| Base H3, no adapter, 49 updates | Original floating | 649.973 s | Passed | +| LightX2V four-step v1.2, peer reduction | Original floating | 59.324 s | Passed | +| Six official FL2V four/eight-step adapters | W8A16 | Four: 58.293-61.263 s; eight: 120.478-121.542 s | Passed | +| Ref2VA four-step, image/video/audio, peer reduction | W8A16 | 181.261 s | Passed | +| Ref2VA eight-step, image/video/audio | W8A16 | 366.800 s | Passed | +| Light4 v1.2 first / last / both keyframes | W8A16 | 66.419 / 64.852 / 69.615 s | Passed | +| FlashGen four-step | Original floating | 59.488 s | Passed | +| FastH3 Dense data-free | Original floating | 56.224 s | Passed | + +The no-adapter original-weight control decreases denoise from 725.819 to +649.973 seconds (10.449687%) and complete request from 835.356 to 700.121 +seconds. Final video/audio latents, all pre-encoding RGB frames and PCM are +bitwise equal. Combined host, projection and residual changes contribute; +this is not an isolated Attention comparison. + +Original floating and W8A16 bases share FP16 projection, scale restoration, +LoRA addition and Attention interfaces. Adapter increments retain unrotated +inputs and enter row reduction before the collective. Residual sharding and +native peer reduction are explicit options, with memory-budget fallback. +No quantization-only or adapter-free restriction selects shared Attention. + +## Reproducing the measured configuration + +Use the native `H3Config` with the model and matching adapter identifiers +from the recorded request contract. The measured TP4 primary configuration +sets `attention_backend="FLASH_ATTN_V100"`, `attention_query_tile=128`, +`fp16_weight_layout="column"`, `residual_sequence_parallel=True`, +`residual_reduction="peer"`, `residual_reduction_memory_gib=4`, +`host_weight_pin_memory=False`, and `share_host_vae_weights=True`. +The mixed Ref4 control uses an explicit 8 GiB communication budget. +Keep each adapter's official sigma, flow-shift and task contract. + +This records explicit measured options, not a universal default or AUTO +promotion. Peer reduction has complete controls for the primary four-step, +original four-step/no-adapter and mixed Ref4 cases. Other rows retain their +measured configurations; the latest peer option is not claimed validated for +every shape/adapter combination. TP1/TP2 capacity checks use layer offload +and ordinary reduction, with no primary-shape >80 claim. + +FastH3 VSA remains outside qualified delivery because its full numerical +diagnostic fails. TeaCache and Cache-DiT/SCM have request-lifecycle and +small-shape GPU evidence; primary-shape official quality and performance are +incomplete. The 243-frame and 15-second runs establish generation/memory +compatibility only. Human audiovisual review and independent official-model +controls remain pending. The implementation stack remains Draft, not merged +to main. + +Evidence: [CAMPAIGN_RESULTS.md](CAMPAIGN_RESULTS.md), +[FA_DEVELOPMENT.md](FA_DEVELOPMENT.md), and artifact root +`/data/minimax-h3/sm70-general-20260909/` with +`peer-api-720p-three-runs/performance.json`, +`peer-api-fi-720p-three-runs/performance.json`, +`peer-api-breadth-summary.json`, `original-base-summary.json`, and full +request/source/binary contracts. No new timing is inferred from the scope +change or documentation update. diff --git a/docs/design/minimax_h3/FA_DEVELOPMENT.md b/docs/design/minimax_h3/FA_DEVELOPMENT.md index 31bd620366..e128af6798 100644 --- a/docs/design/minimax_h3/FA_DEVELOPMENT.md +++ b/docs/design/minimax_h3/FA_DEVELOPMENT.md @@ -1,8 +1,14 @@ # FlashAttention development focus -Further Attention development and exhaustive workflow qualification concentrate -on FlashAttention-V100 at the user's request. Validated FlashInfer remains -available; new FI optimization and full FI acceptance are paused. +Current delivery uses the retained FlashAttention-V100 implementation on +already working H3 dense workflows. At the user's latest request, further +investigation of slower Attention prototypes and new workflow expansion are +paused. Validated FlashInfer remains available, with further FI optimization +and exhaustive FI acceptance paused. The unfinished matrix and unchanged +quality/performance gates remain recorded as incomplete. + +See [CURRENT_STATUS.md](CURRENT_STATUS.md) for the retained configuration, +measured workflow coverage and delivery limits. With native TP4 peer rows, shared pageable VAE weights and a complete request warmup plus three unprofiled requests, FA records a minimum-card median @@ -78,3 +84,21 @@ slowdown limit. No full-model run or production promotion follows. Evidence: `attention-fa-warp16-staging-q96`, `attention-fa-register-k128`, `attention-fa-register-k128-vprefetch`, and `fa-vprefetch-after-register.json`. + +The artifact-only D128 tiled-GEMM port also closes without promotion. Padding +only PV's reduction width to 32 fixes its nonaligned-key vector read; the +original failing boundary then passes Compute Sanitizer with zero errors. +Eleven boundary/stress cases and the actual input pass the independent FP32 +operator gate. Actual-input relative L2 is 0.000231 against that oracle and +0.000335 against FA; output is not bitwise equal to FA. Its paired median is +212.961273 ms versus retained FA's 149.299194 ms, so no complete sampling +quality run is justified. + +A bounded Nsight Systems attribution records 209.361943 ms operator wall +time: QK plus tile softmax 142.181160 ms, PV plus probability rescaling +61.613049 ms, and GPU idle/host gaps 0.582682 ms. API durations overlap GPU +execution and cannot be added to those kernel durations. Packing or CPU +scheduling does not explain this regression. A row-maximum epilogue follow-up +completed its CPU build but was stopped at the user's scope change before +any GPU numerical check or timing. Neither artifact replaces production FA. +Evidence: `attention-fa-tiled-gemm-d128` and `attention-fa-tiled-gemm-rowmax`.