mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-11-08 09:39:58 +08:00
qutaoandClaude Opus 4.8
c48e37333c
Address maintainer review follow-ups: - Move the four loose SM120 NVFP4 files under maint/gemm/ into a dedicated maint/gemm/gemm_sm120/ folder to distinguish them from other gemm types, updating the repo-root path resolution, the benchmark usage string, the cutlass_nvf4_ref.cu source path, and the quantizer loader in the scale-layout test. - Guard the SM120-hardware test points with tilelang.testing.requires_cuda_compute_version_eq(12, 0): replace the manual capability skip in the example CLI test and tighten the block-scale kernel tests from _ge(12, 0) to _eq(12, 0). - Reuse the native tilelang.cuda.target.target_is_sm120 helper in _is_explicit_non_sm120_cuda instead of hand-parsing the arch string, keeping the explicit-arch guard so targets without an arch attr never trip the C++ ICHECK on hosts that have no SM120 GPU. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
102 lines
3.5 KiB
Python
102 lines
3.5 KiB
Python
import importlib.util
|
|
import sys
|
|
from pathlib import Path
|
|
|
|
import tilelang.testing
|
|
|
|
|
|
def _load_sm120_example(monkeypatch):
|
|
repo_root = Path(__file__).resolve().parents[3]
|
|
example = repo_root / "examples/gemm_sm120/sm120_nvfp4_blockscaled_gemm.py"
|
|
monkeypatch.setattr(sys, "argv", [str(example)])
|
|
spec = importlib.util.spec_from_file_location("sm120_nvfp4_blockscaled_gemm_example", example)
|
|
module = importlib.util.module_from_spec(spec)
|
|
assert spec.loader is not None
|
|
spec.loader.exec_module(module)
|
|
return module
|
|
|
|
|
|
def test_sm120_nvfp4_example_cli_keeps_internal_strategy_fixed(monkeypatch):
|
|
module = _load_sm120_example(monkeypatch)
|
|
args = module.parse_args()
|
|
|
|
assert not hasattr(args, "micro_pipeline")
|
|
assert not hasattr(args, "sf_layout")
|
|
assert not hasattr(args, "scale_storage_layout")
|
|
assert not hasattr(args, "manual_ws2")
|
|
|
|
assert module._SM120_SCALE_LAYOUT == "blockscaled_chunk_kmajor"
|
|
|
|
|
|
def test_sm120_nvfp4_example_source_has_no_legacy_strategy_flags():
|
|
repo_root = Path(__file__).resolve().parents[3]
|
|
source = (repo_root / "examples/gemm_sm120/sm120_nvfp4_blockscaled_gemm.py").read_text()
|
|
|
|
forbidden_flags = [
|
|
"--micro-pipeline",
|
|
"--manual-ws2",
|
|
"--manual-ws2-split",
|
|
"--manual-ws2-sf-load",
|
|
"--manual-ws2-sf-layout",
|
|
"--manual-ws2-ab-shared-storage",
|
|
"--manual-ws2-ab-copy-view",
|
|
"--scale-storage-layout",
|
|
"--sf-layout",
|
|
"--path",
|
|
"--run-cutlass",
|
|
"--cutlass-build-dir",
|
|
"--cutlass-binary",
|
|
"--rebuild-cutlass",
|
|
"--cmake",
|
|
"--nvcc",
|
|
]
|
|
for flag in forbidden_flags:
|
|
assert flag not in source
|
|
|
|
forbidden_harness_terms = [
|
|
"build_cutlass",
|
|
"subprocess",
|
|
"CMAKE_CUDA_COMPILER",
|
|
]
|
|
for term in forbidden_harness_terms:
|
|
assert term not in source
|
|
|
|
|
|
@tilelang.testing.requires_cuda_compute_version_eq(12, 0)
|
|
def test_sm120_nvfp4_example_kernel_handles_mn_tail_tiles(monkeypatch):
|
|
import pytest
|
|
|
|
torch = pytest.importorskip("torch")
|
|
|
|
from tilelang.quantize import swizzle_blockscaled_chunk_kmajor_scale_words
|
|
|
|
module = _load_sm120_example(monkeypatch)
|
|
for M, N, K in [(257, 384, 512), (130, 128, 256), (128, 136, 256)]:
|
|
kernel = module.sm120_nvfp4_blockscaled_gemm(M, N, K)
|
|
|
|
A = module._make_packed_fp4(M, K, seed=3)
|
|
B = module._make_packed_fp4(N, K, seed=4)
|
|
SFA_semantic = module._make_binary_scale_words(M, K, seed=5)
|
|
SFB_semantic = module._make_binary_scale_words(N, K, seed=6)
|
|
SFA = swizzle_blockscaled_chunk_kmajor_scale_words(SFA_semantic)
|
|
SFB = swizzle_blockscaled_chunk_kmajor_scale_words(SFB_semantic)
|
|
assert SFA.shape[0] % 128 == 0
|
|
assert SFB.shape[0] % 128 == 0
|
|
|
|
C = torch.empty((M, N), device="cuda", dtype=torch.bfloat16)
|
|
kernel(A, B, SFA, SFB, C)
|
|
torch.cuda.synchronize()
|
|
module._verify(A, B, SFA_semantic, SFB_semantic, C, torch.bfloat16)
|
|
|
|
|
|
def test_sm120_nvfp4_example_kernel_rejects_unsupported_tails(monkeypatch):
|
|
import pytest
|
|
|
|
module = _load_sm120_example(monkeypatch)
|
|
# simultaneous M and N tails hit a known copy-lowering boundary bug
|
|
with pytest.raises(ValueError, match="simultaneous M and N tail"):
|
|
module.sm120_nvfp4_blockscaled_gemm(257, 136, 512)
|
|
# bf16 output rows must stay 16-byte aligned
|
|
with pytest.raises(AssertionError, match="multiple of 8"):
|
|
module.sm120_nvfp4_blockscaled_gemm(128, 130, 256)
|