Repository navigation
Fix test helper discovery in pytest and sanitizer runner - #446
lucifer1004 wants to merge 2 commits into
Conversation
| return decorator | ||
|
|
||
|
|
||
| test_filter.__test__ = False |
There was a problem hiding this comment.
🔵 suggestion: Consider adding a one-line comment explaining why test_filter.__test__ = False is set (pytest / test_sanitizer.py opt-out marker for a test-prefixed helper). Future readers may otherwise be tempted to remove what looks like a stray attribute assignment.
🤖 v5
| for module_name in [os.path.splitext(f)[0] for f in files if f not in exclude_files] | ||
| for name, obj in inspect.getmembers(importlib.import_module(module_name)) | ||
| if inspect.isfunction(obj) and name.startswith('test') and 'test_filter' not in name | ||
| if inspect.isfunction(obj) and name.startswith('test') and getattr(obj, '__test__', True) |
There was a problem hiding this comment.
🔵 suggestion: With default discovery, the new tests/test_test_discovery.py module (test_pytest_discovery, test_sanitizer_discovery) will now itself be picked up and run under compute-sanitizer for every tool. These are pure-Python CPU tests that spawn pytest subprocesses and mock subprocess.run, so running them under the sanitizer is wasted time (and test_sanitizer_discovery re-executes this runner via runpy inside the sanitized process). Setting __test__ = False on them is not appropriate since pytest should collect them, so adding 'test_test_discovery.py' to exclude_files is the right lever. Not blocking.
🤖 v5
There was a problem hiding this comment.
The whole test has been removed.
|
|
||
|
|
||
| def test_pytest_discovery(): | ||
| with tempfile.TemporaryDirectory() as directory: |
There was a problem hiding this comment.
🔵 suggestion: The 60s timeout for the pytest subprocess may be tight: the sample module imports torch (via deep_gemm/testing/utils.py) in a fresh interpreter twice (collect-only + run). On cold caches or slow CI runners torch import alone can take 10–30s, so this could be flaky. Consider raising the timeout (e.g. 300s); a hang is still caught, just later.
🤖 v5
There was a problem hiding this comment.
The whole test has been removed.
| spec = importlib.util.spec_from_file_location(path.stem, path) | ||
| module = importlib.util.module_from_spec(spec) | ||
| spec.loader.exec_module(module) | ||
| package = ModuleType('deep_gemm') |
There was a problem hiding this comment.
🔵 suggestion: patch.object(os, 'listdir', ...), patch.object(importlib, 'import_module', ...) and patch.object(subprocess, 'run', ...) are global patches held while runpy.run_path executes the runner. This works today because runpy/argparse don't call these, but if the runner later grows other imports or filesystem calls this test could break in confusing ways. A short comment noting the assumption (and that import_module returns the same sample module for any name) would help maintainers.
🤖 v5
There was a problem hiding this comment.
The whole test has been removed.
🤖 ds-review-bot Code Reviewv6v5This MR fixes accidental collection of the v4本 MR 修复了测试助手被 pytest / sanitizer runner 误收集的问题。 标记只作用于装饰器工厂,被装饰的测试函数仍可被发现,条件过滤行为不变;实现简洁、与现有模式一致。未发现正确性、安全性、性能或行为回归问题,建议合并。 Files reviewed: 3 |
FP8 paged MQA logits gains PAGE_KV=32 (BLOCK_KV derived as min(PAGE_KV, 64), mirroring the FP4 sibling). Device-only update: this fork's host launcher still restricts paged FP8 to page 64, so page32 stays inert until host glue opts in (vllm-project#14). Validated on sm_120a: test_sm120_mqa.py + test_sm120_fp8_fp4.py 23/23 passed from a fresh JIT cache (only the pre-existing test_filter collection quirk remains, deepseek-ai#446). Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
* feat(sm120): vendor DeepGEMM-sm120 v0.1.3 device layer Replace the in-tree SM120 device headers with byte-identical vendored copies from lucifer1004/DeepGEMM-sm120 v0.1.3 (merged superset of both lineages). Marker-free vendoring: no manifests, tooling, or header comments enter this fork; provenance and drift control live on the canonical repo's side (per-tag manifest, check_vendor.py --fork, scheduled downstream watcher). Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> * feat(sm120): adapt host launchers to vendored kernel signatures The vendored DeepGEMM-sm120 v0.1.0 device layer widens several kernel ABIs; adapt the SM120 host glue with neutral mappings: - sm120_bf16_gemm: pass the epilogue operator as a runtime argument (EpilogueArgs, same marshalling as SM100) and split the old int64_t stride_cd_m/stride_cd_batch into u32 stride_d_m/stride_c_m/ stride_d_batch; stride_c_m=0 keeps C sharing D's row stride. - sm120_fp8_fp4_gemm_1d1d: pass shape_cd_m (= shape_m), the runtime epilogue argument, and stride_c_m (= 0) after stride_cd_batch. - sm120_split_k_reduce: the kWithAccumulation template flag is gone; instantiate the 3-arg template and map accumulation onto D in place to the new runtime C operand (gmem_c = D with D's own strides, with_alpha = false). - Masked m-grouped GEMM (bf16 + fp8/fp4): disable the TMA-store epilogue (swizzle_cd_mode = 0) when m % BLOCK_M != 0. The vendored kernels dropped the old masked-boundary scalar fallback, so a boundary tile's full-tile TMA store would spill into the next group's rows; the scalar store path carries the masked row bounds. This mirrors the proven nv_dev-lineage launchers. Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> * fix(sm120): guard k-grouped GEMM TMA store against unaligned M The vendored v0.1.0 GEMM kernels have no boundary-tile scalar fallback in the TMA-store epilogue, and the k-grouped D descriptor is one flat 2D map over all groups, so TMA can only clamp at the outermost M dim: with m % BLOCK_M != 0, a group's partial tail M tile TMA-store spills into the next group's slab. Mirror the masked-launcher guard (and the nv_dev-lineage reference glue) in both k-grouped launchers: force swizzle_cd_mode = 0 (group-bounded scalar store path) for unaligned m. Repro evidence (num_sms=2, groups=2, m=48, n=128, ks=[8192, 128]; the asymmetric K makes group 0's tile store last, so the spill overwrites group 1's head rows deterministically): - before: bf16/bf16-out 5/5 iterations contaminated (d[1][0:16] = 0.0, expected 768); fp8/fp32-out 5/5 contaminated (d[1][0:16] = 0.0, expected 771) - bf16/fp32-out was already clean (0/5): the bf16 kernel gates its TMA-store epilogue on sizeof(cd_dtype_t) <= 2, so fp32 outputs take the scalar path; the fp8 kernel has no such dtype gate - after: all three configurations 0/5 contaminated Adds test_sm120_kgroup_unaligned_m_tail_tile_isolation to tests/test_sm120_bf16.py (fp32 and bf16 outputs) and tests/test_sm120_fp8_fp4.py (NT and TN layouts), asserting per-group correctness including slab head rows plus flat-storage guard regions. Full SM120 suite: 48 passed (46 baseline nodes + 2 new). Memcheck on the two affected files: 20 passed, 0 sanitizer errors. Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> * feat(sm120): vendor DeepGEMM-sm120 v0.1.4 device layer FP8 paged MQA logits gains PAGE_KV=32 (BLOCK_KV derived as min(PAGE_KV, 64), mirroring the FP4 sibling). Device-only update: this fork's host launcher still restricts paged FP8 to page 64, so page32 stays inert until host glue opts in (#14). Validated on sm_120a: test_sm120_mqa.py + test_sm120_fp8_fp4.py 23/23 passed from a fresh JIT cache (only the pre-existing test_filter collection quirk remains, deepseek-ai#446). Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> * fix(sm120): vendor DeepGEMM-sm120 v0.1.5 generic-proxy WAR fences Vendor lucifer1004/DeepGEMM-sm120 v0.1.5 (b31a688): 14 fence_view_async_shared() sites across 8 SM120 kernel headers, ordering generic-proxy smem reads (ldmatrix/ld_shared) before the empty-barrier arrive that hands each stage back to the TMA producer. Ports deepseek-ai#453 (bf16 GEMM) and fixes the same hazard found by audit in bmk_bnk_mn, tf32_hc_prenorm, fp8_fp4_gemm_1d1d, and the fp8/fp4 dense+paged MQA logits kernels. Sparse MQA is exempt (legacy cp.async producer). Perf A/B on RTX PRO 6000: no regression beyond run-to-run noise. Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> --------- Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Summary
test_filterdecorator factory with pytest's existing__test__ = Falseopt-out convention.test_filter.Why
Running an entire test module through pytest can collect the imported
test_filter(condition)helper and fail withfixture 'condition' not found. Direct script execution avoids this, while the sanitizer runner currently has a name-specific exclusion. The shared marker handles aliases and explicit opt-outs without accidentally excluding real tests such astest_filter_behavior.The marker applies only to the decorator factory. Decorated test functions remain discoverable, and existing condition-based execution/filtering behavior is unchanged.
Validation
Used a temporary CPU-only verification harness (not included in this PR):
78b6900.getattr(obj, "__test__", True)directly.git diff --check: passed.No GPU kernels, actual Compute Sanitizer execution, or native rebuild were required. This small PR is independent of the SM120/API migration work.