forked from ccf-ai-infra/TileOPs-Metax
17 Commits
| Author | SHA1 | Message | Date |
|---|---|---|---|
wawahejun
|
ce1b15c28f |
docs: add MetaX C500 summer-camp guide on top of upstream dev
Source tree is identical to MetaX-MACA/TileOPs-Metax dev at f02d3d8; this commit carries only the summer-camp documentation and PR templates. Content verified by running everything on a real MetaX C500 (MACA 3.7.1.5, torch 2.8.0+metax3.7.1.3, tilelang 0.1.10+cuda.gitf549117c, sGPU slice 16000 MiB). Installation (high severity). The documented `make install`, `pip install tileops`, and bare `python3 -m venv .venv` steps destroy a working MACA environment. The container's TileLang is an in-place source build imported via PYTHONPATH, so pip reports it as absent and resolves the official CUDA wheel over it; a venv without --system-site-packages cuts off the MetaX PyTorch build and the ABI-coupled apache-tvm-ffi. Replace those steps with the PYTHONPATH setup, document that tileops needs no install at all, and note that --no-deps is the only safe install form (as scripts/ci/install_tileops.sh already does). Flag -c constraints.txt as CUDA-CI-only for the same ABI reason. Add TileLang provenance and backend checks to the verification list, which previously covered mx-smi, torch and einops but not the component most likely to be wrong. Quick start. GemmOp(M, N, K, dtype=...) does not match the implementation -- GemmOp is input-inferred and takes only trans_a/trans_b. Fix the signature and document the trans_b default, keeping the original M,N,K of 1024,1024,512, which passes on C500 via the MACA dispatch path. New sections 1.2 and 1.3 in the migration guide. Document the is_maca() dispatch to *_maca.py kernels, and that a gated kernel does not imply an unusable Op: GemmKernel declares [89, 90] and is gated on C500, yet GemmOp works because it dispatches to gemm_maca.py ([80, 86, 89, 90]). Availability must be judged from what the Op layer dispatches to, not from one kernel's supported_archs. List the 20 declarations that exclude 80 as unsuitable migration targets, and note that adding a *_maca.py kernel plus dispatch is a good target instead. Record that get_sm_version() reuses NVIDIA's encoding, so C500 reports 80 while sharing nothing with Ampere, and that the raw "architecture 80" message names no device. Document that a usable Op still has shape limits: SoftmaxFwdOp fails above a 1024-wide reduction dimension (mcErrorInvalidValue), independent of row count. Document that a parent process which has imported tilelang will see any subprocess that imports it again SIGKILLed with no output, which aborts tests/test_validate_manifest.py at exit 137, and give the deselect workaround. Roofline. Record the sGPU slice quota and state whether peaks are whole-card or slice-scaled; dividing a slice measurement by a whole-card peak yields an unexplainable efficiency. Verified on C500 against this tree: validate_manifest.py exit 0; 29 passed across test_ops_manifest.py, test_kernel_map_install.py and benchmarks/tests; GemmOp passes at 1024x1024x512, 1024^3 and 4096^3; the documented quick-start snippet and every self-check command run as written. pre-commit and ruff are unavailable in this container (installing them would invoke pip dependency resolution), so formatting was checked via git diff --check and end-of-file newlines instead. Squashed documentation commits by Beckylu <648245013@qq.com> and FrRay <1077376663@qq.com> covering the summer-camp guide, PR templates, and README translations. Co-Authored-By: Beckylu <648245013@qq.com> Co-Authored-By: FrRay <1077376663@qq.com> |
|
|
Cao Ying
|
0a9bf1e40a
|
[Chore][Cleanup] repo slimming: dead code, duplicated tests/benches, compat shims, file defragmentation (#1764)
Closes #1763 ## Summary - Remove dead helpers, redundant elementwise benchmark drivers, decorative banners, and stale compatibility aliases. - Consolidate duplicated tests and private benchmark wrappers onto shared parametrized fixtures and `ManifestBenchmark`. - Merge fragmented workload, reduction-op, MHC, and normalization benchmark modules; retarget imports and manifest source paths. - Preserve canonical runtime behavior while reducing the repository by roughly 4.3k net lines. ## Test plan - [x] pre-commit passed - [x] `python scripts/validate_manifest.py --strict` passed - [x] Repository-wide pytest collection completed: 5,178 tests, 0 collection errors - [x] Touched test and benchmark modules collect cleanly --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
ChongLi
|
d3b58074f9
|
[Perf][Conv2d] Optimize symmetric conv2d layout transform overhead (#1645)
Closes #1644 ## Summary This PR addresses the poor `conv2d` performance reported in #1644 (CI run `28398916051`) by introducing `Conv2dSymmetricKernel`, an NHWC-implicit-GEMM path for symmetric conv2d cases (`kernel_h == kernel_w`, `groups == 1`, `c_in % 32 == 0`). Key changes: - Added `Conv2dSymmetricKernel` and dispatched symmetric conv2d cases to it. - Unified the three layout transpose macros (`nchw_to_nhwc_input`, `kcrs_to_krsc_weight`, `nhwc_to_nchw_output`) into a single `transpose_spatial_channel` macro. - Parameterized tile size and fastest-varying dimension so the transpose kernels can be tuned per shape in follow-up work. - Reduced the weight transpose `spatial_block` from `32` to `16`, which increases the spatial grid for small kernels (`3x3`, `5x5`, `7x7`). - Added compile-time asserts for tile/thread constraints. No public API is modified. ## Test plan - [x] `python -m py_compile tileops/kernels/convolution.py` passed - [x] `tests/ops/test_convolution.py -k "test_conv2d"` — 20 passed - [x] `tests/ops/test_convolution.py -k "dispatches"` — 6 passed ## Benchmark Comparison of this branch (`Conv2dSymmetricKernel`) against `upstream/main` (`Conv2dKernel`) on H200, clock-locked 1500 MHz: | case | main Conv2dKernel (ms) | this branch Conv2dSymmetricKernel (ms) | speedup | |------|------------------------|----------------------------------------|---------| | deeplabv3-aspp-3x3-rate12-fp16 | 0.4278 | 0.1152 | **3.71x** | | stage-transition-5x5-s2-fp16 | 0.0748 | 0.0242 | **3.09x** | | midres-5x5-s1-fp16 | 0.0392 | 0.0183 | **2.14x** | | stride2-bf16 | 0.0279 | 0.0136 | **2.05x** | | stage-transition-3x3-s2-fp16 | 0.0279 | 0.0163 | **1.71x** | | resnet-3x3-fp16 | 0.0169 | 0.0150 | **1.13x** | `Conv2dSymmetricKernel` is faster on all representative symmetric cases, with larger gains on bigger/dilated kernels. Further improvements are expected once the new transpose tile parameters are exposed to autotune. ## Additional context Addressed Gemini Code Assist review feedback: - Renamed local `h`/`w` to `h_idx`/`w_idx` in `transpose_spatial_channel` to avoid shadowing outer parameters. - Switched to interleaved channel mapping for coalesced NHWC writes in the `channel_fastest` branch. - Added `c_in % block_k == 0` filtering in `Conv2dSymmetricKernel.autotune_configs`. - Made bias addition conditional on `has_bias` in `conv_nhwc_implicit_gemm_bias`. This change is the first step toward a larger performance fix for #1644. The unified macro makes it straightforward to add `input_spatial_block`, `output_spatial_block`, etc. to `Conv2dSymmetricKernel.autotune_configs` in a follow-up PR. |
|
|
ChongLi
|
1c41f43d8b
|
[Feat][Convolution] Support grouped conv2d and conv3d (#1568)
Closes #1521 ## Summary - Add native grouped Conv2d/Conv3d TileLang kernels for bias and no-bias variants. - Wire Conv2d/Conv3d op dispatch to grouped kernels while preserving existing dilation support from `main`. - Update the convolution manifest contract for grouped Conv2d/Conv3d status, kernel maps, and per-group roofline accounting. - Add grouped correctness coverage and model-derived grouped convolution benchmark cases. ## Benchmark ```bash CUDA_VISIBLE_DEVICES=1 TILELANG_CLEANUP_TEMP_FILES=1 python -m pytest \ 'benchmarks/ops/bench_convolution.py::test_conv2d_bench[mobilenetv2-depthwise-fp16]' \ 'benchmarks/ops/bench_convolution.py::test_conv2d_bench[resnext-grouped-3x3-fp16]' \ 'benchmarks/ops/bench_convolution.py::test_conv3d_bench[3d-resnext-grouped-k3-fp16]' \ 'benchmarks/ops/bench_convolution.py::test_conv3d_bench[3d-resnext-grouped-k3-b8-fp16]' \ -vvs ``` Result: `4 passed in 138.39s`. | Case | TileOps latency | TileOps TFLOP/s | TileOps bandwidth | Torch latency | Torch TFLOP/s | Torch bandwidth | | --- | ---: | ---: | ---: | ---: | ---: | ---: | | `conv2d mobilenetv2-depthwise-fp16` | 0.0064 ms | 0.2812 | 0.0626 TB/s | 0.0065 ms | 0.2772 | 0.0617 TB/s | | `conv2d resnext-grouped-3x3-fp16` | 0.0041 ms | 3.4882 | 0.1498 TB/s | 0.0200 ms | 0.7207 | 0.0309 TB/s | | `conv3d 3d-resnext-grouped-k3-fp16` | 0.0147 ms | 5.8872 | 0.1645 TB/s | 0.7442 ms | 0.1165 | 0.0033 TB/s | | `conv3d 3d-resnext-grouped-k3-b8-fp16` | 0.1269 ms | 5.4643 | 0.1519 TB/s | 5.7221 ms | 0.1212 | 0.0034 TB/s | |
|
|
Ang Gao
|
4c1ddeb827
|
[Fix] gpu-smoke on new tvm-ffi stack (#1609)
## Summary Fixes #1607. This PR restores `tests/ops -m smoke` on the new image-baked TileLang/tvm-ffi stack by addressing the import, binder, lowering, and timeout issues uncovered by the GPU smoke run. Changes include: - remove remaining `tvm.tir` smoke gates/import assumptions in GQA FP8 and topk selector paths - handle new TileLang binder argument errors in binary autotune fallback - avoid Conv1d local scalar binding that new TileLang rejects - inline GatedDeltaNet recurrence scalars that were treated as missing API args - materialize side-effecting MoE extern atomic results before indexed stores so lowering does not duplicate atomic calls - fix FP8LightingIndexer `block_Q > 1` GEMM output shape and use batched queries by default to avoid smoke timeout - make `gpu-smoke` use a run-local `TRITON_CACHE_DIR` under `/ci-cache/triton/gpu-smoke-${GITHUB_RUN_ID}-${GITHUB_RUN_ATTEMPT}` so TorchInductor/Triton writes do not collide with stale shared-cache hash directories owned by another runner user ## Validation Image: `ghcr.io/tile-ai/tileops-runner:65dbc98-torch2.10` GPU: local GPU1 (H200 visible as device 0 in container) ```text python -m pytest tests/ops/attention/test_gqa_fp8.py \ tests/ops/test_topk_selector.py \ tests/ops/test_binary_arith.py \ tests/ops/test_convolution.py \ tests/ops/test_gated_deltanet_fwd.py \ tests/ops/test_gated_deltanet_chunkwise_bwd.py \ tests/ops/test_moe_permute_nopad.py \ tests/ops/test_moe_permute_align.py \ tests/ops/test_moe_fused_moe.py \ tests/ops/test_fp8_lighting_indexer.py \ -m smoke -q --tb=short --disable-warnings 162 passed, 102 deselected python -m pytest tests/ops -m smoke -q --tb=short --disable-warnings 2180 passed, 24 skipped, 955 deselected, 283 warnings in 148.30s TRITON_CACHE_DIR=/data7/shared/ci-cache/triton/gpu-smoke-local-codex \ python -m pytest tests/ops/test_mamba.py tests/ops/test_moe_fused_moe.py \ -m smoke -q --tb=short --disable-warnings 42 passed, 21 deselected ``` |
|
|
ChongLi
|
d10b6b18d3
|
[Maintain][Convolution] Promote conv1d ops to implemented (#1556)
Closes #1555 ## Summary - Add Conv1d output-shape inference, dtype validation, and roofline methods for no-bias and bias variants. - Validate Conv1d dtypes in `forward()` before dispatching kernels. - Promote `Conv1dFwdOp` and `Conv1dBiasFwdOp` manifest entries from `spec-only` to `implemented`. ## Test plan - [x] `python scripts/validate_manifest.py --strict --check-op Conv1dFwdOp` - [x] `python scripts/validate_manifest.py --strict --check-op Conv1dBiasFwdOp` - [x] `PYTHONPATH="$PWD" python -m pytest -q tests/ops/test_convolution.py -k "conv1d"` - [x] `python scripts/validate_manifest.py --strict` - [x] `PYTHONPATH="$PWD" python -m pytest -q tests/ops/test_convolution.py` - [x] `git diff --check` - [x] `python -m compileall -q tileops/ops/convolution.py` - [x] Commit/push hooks passed ## Regression - Conv1d targeted tests: `19 passed, 24 deselected` - Full convolution test file: `43 passed` - Manifest strict validation passes for both Conv1d entries; remaining Conv1d L4 benchmark manifest-driven messages are advisory because `source.bench_manifest_driven` remains false. ## Additional context - `scripts/validate.sh` is not present in this checkout, so the PR lifecycle pre/post gates from the skill could not be run directly. The repo-available manifest, pytest, diff, compile, and hook checks above were used instead. |
|
|
ChongLi
|
4338516fa7
|
[Fix][Convolution] Add dilation support (#1554)
Closes #1519 ## Summary - Add dilation support to Conv2d/Conv2dBias and Conv3d/Conv3dBias kernel dispatch and indexing. - Extend convolution tests with dilation cases in the existing Conv2dFixture and Conv3dFixture, without adding standalone dilation-only tests. - Add model-derived dilation benchmark cases for DeepLabV3 ASPP Conv2d and 3D U-Net ASPP Conv3d. ## Test plan - [x] `CUDA_VISIBLE_DEVICES=1 TMPDIR=/home/lyc/Project/TileOPs/.tmp/tvm python -m pytest tests/ops/test_convolution.py -q` (`45 passed in 8.75s`) - [x] `CUDA_VISIBLE_DEVICES=1 TMPDIR=/home/lyc/Project/TileOPs/.tmp/tvm python -m pytest benchmarks/ops/bench_convolution.py -k "rate12 or rate6" -vvs` (`2 passed, 22 deselected in 8.05s`) - [x] `python -m ruff check tests/ops/test_convolution.py benchmarks/ops/bench_convolution.py tileops/ops/convolution.py tileops/kernels/convolution.py` - [x] `git diff --check -- tests/ops/test_convolution.py benchmarks/ops/bench_convolution.py tileops/ops/convolution.py tileops/kernels/convolution.py tileops/manifest/convolution.yaml` - [x] `python scripts/validate_manifest.py --check-op Conv2dFwdOp` (passed with existing advisory warnings) - [x] `python scripts/validate_manifest.py --check-op Conv3dFwdOp` (passed with existing advisory warnings) ## Benchmark Run on GPU 1 with `CUDA_VISIBLE_DEVICES=1`; GPU reported by benchmark harness: NVIDIA H200. Ratio uses the repository benchmark convention: `torch_latency / tileops_latency`. | Case | Source | TileOps latency | PyTorch baseline latency | Baseline / TileOps | | --- | --- | ---: | ---: | ---: | | `deeplabv3-aspp-3x3-rate12-fp16` | DeepLabV3/DeepLabV3+ ASPP 3x3 atrous conv on stride-16 encoder features | 0.2963 ms | 0.1273 ms | 0.43x | | `3d-unet-aspp-3x3x3-rate6-fp16` | 3D U-Net + 3D ASPP medical segmentation branch | 0.0970 ms | 0.5258 ms | 5.42x | ## Regression - Conv2d no-bias regression now covers `dilation=2` against `torch.nn.functional.conv2d`. - Conv3d no-bias regression now covers `dilation=2` against `torch.nn.functional.conv3d`. |
|
|
ChongLi
|
20397575ca
|
[Perf][Convolution] Align convolution layout kernels (#1527)
Close #1516 ## Summary - Move Conv2d and Conv3d layout handling into TileLang kernels so benchmark execution no longer depends on the skipped external layout path. - Align Conv2d pointwise with the Conv1d pointwise data/weight ordering pattern. - Re-enable Conv2d/Conv3d benchmark cases under TileLang 0.1.9 by clearing the convolution-specific skip prefixes. ## Test plan - `TMPDIR=/home/lyc/tvm_tmp conda run --no-capture-output -n tileops-dev python -m pytest benchmarks/ops/bench_convolution.py -vvs --tb=short` - Current branch: `22 passed in 643.83s` - main comparison worktree: `22 passed in 317.00s` ## Benchmark Compared current branch `e04aeb1` plus benchmark enable commit against main `63f3022` on NVIDIA H200, Driver 575.57.08, Torch 2.9.0+cu128, CUDA 12.8. Benchmarks used the default benchmark parameters; no tune-disabling override was passed, so cases without special handling ran with `tune=True`. Summary by TileOps latency, where speed ratio is `main_latency / current_latency`: | Op | Cases | Avg speed ratio | Faster | Similar | Slower | Notes | | --- | ---: | ---: | ---: | ---: | ---: | --- | | Conv1d | 6 | 1.01x | 0 | 6 | 0 | Essentially unchanged. | | Conv2d | 13 | 0.71x | 1 | 3 | 9 | Layout-integrated kernels regress most non-1x1 large-spatial and 5x5 cases; one small stem case improves. | | Conv3d | 3 | 0.89x | 1 | 0 | 2 | Stem case improves, but larger bf16 case regresses significantly. | Notable Conv2d changes: | Case | main ms | current ms | ratio | change | | --- | ---: | ---: | ---: | ---: | | N=1, Cin=3, HxW=112x112, Cout=64, K=3x3, S=2, fp16 | 0.0067 | 0.0035 | 1.91x | -47.8% latency | | N=1, Cin=256, HxW=112x112, Cout=512, K=3x3, S=1, fp16 | 0.0686 | 0.4381 | 0.16x | +538.6% latency | | N=1, Cin=128, HxW=56x56, Cout=256, K=5x5, S=2, fp16 | 0.0221 | 0.0882 | 0.25x | +299.1% latency | | N=2, Cin=64, HxW=56x56, Cout=256, K=1x1, fp16 | 0.0041 | 0.0043 | 0.95x | +4.9% latency | | N=1, Cin=512, HxW=7x7, Cout=2048, K=1x1, fp16 | 0.0050 | 0.0108 | 0.46x | +116.0% latency | Notable Conv3d changes: | Case | main ms | current ms | ratio | change | | --- | ---: | ---: | ---: | ---: | | N=1, Cin=3, DxHxW=16x112x112, Cout=64, K=3x3x3, fp16 | 0.0445 | 0.0253 | 1.76x | -43.1% latency | | N=1, Cin=64, DxHxW=8x56x56, Cout=128, K=3x3x3, S=2, fp16 | 0.0248 | 0.0333 | 0.74x | +34.3% latency | | N=1, Cin=32, DxHxW=32x64x64, Cout=64, K=3x3x3, bf16 | 0.0845 | 0.4726 | 0.18x | +459.3% latency | Full local report: `output_mid/conv_benchmark_full_tune_20260528/report.md`. |
|
|
ChongLi
|
7b9aca2e21
|
[Feat][Convolution] Add native groups support and string padding coverage for conv1d (#1516)
## Summary Closes #1515. Adds native `groups` support for `Conv1dFwdOp` / `Conv1dBiasFwdOp` with a dedicated `GroupConv1dKernel` path: - Adds `_conv1d_group_kernel`, which treats `groups` as an explicit grid dimension so output-channel tiles stay within one group. - Keeps `_conv1d_kernel` and `_conv1d_pointwise_kernel` as `groups=1` paths only. - Uses the grouped kernel for `groups > 1`, including pointwise-shaped grouped convs. - Keeps a low-level `_conv1d_direct_kernel` fast path for depthwise cases where `c_in_g == 1 && c_out_g == 1`; no public direct kernel class is exposed. - Supports non-16-multiple `c_out_g` by masking within the grouped kernel while keeping TensorCore block choices 16-aligned. ## Test plan - Extended `Conv1dFixture` so the main `test_conv1d` covers: - normal `groups=1` conv1d - string padding: `padding='valid'` and `padding='same'` - grouped conv: `groups=2` - non-16 output channels per group: `groups=3, c_out_g=24` - Conformer-style depthwise conv: `groups=channels, kernel_size=31` - Tightened fp16 conv1d comparison tolerance to `1e-3 / 1e-3`; bf16 keeps the existing tolerance. - Verified locally: - `python -m ruff check tests/ops/test_convolution.py tileops/kernels/convolution.py tileops/ops/convolution.py tileops/kernels/__init__.py` - `TMPDIR=/home/lyc/Project/TileOps-workspace2/.tmp python -m pytest tests/ops/test_convolution.py -q` - `python scripts/validate_manifest.py --check-op Conv1dFwdOp` - `git diff --check HEAD` ## Regression - No public API changes for `groups=1` users. - `groups=1` dispatch remains on the existing generic or pointwise conv1d kernels. - `Conv1dPointwiseKernel` remains `groups=1` only; grouped pointwise-shaped cases dispatch to `GroupConv1dKernel`. - Dilation behavior is unchanged. ## Additional context - Earlier grouped indexing could allow a tile to cross group boundaries unless `c_out_g` aligned with `block_m`. The new grouped kernel avoids this by launching tiles within each group instead of deriving `group_id` from a global output-channel tile. - The grouped kernel masks partial `block_m` and `block_k` tiles, so legal grouped shapes are covered for correctness even when per-group channels are not multiples of 16. - The depthwise direct kernel casts operands to the accumulator dtype before multiplication to keep fp16 numerical differences within the standard conv1d tolerance. |
|
|
ChongLi
|
acfea41b69
|
[BugFix][CONV1D] Align tensor layout with manifest (#1469)
## Summary This PR aligns Conv1d runtime tensor layout with the manifest-declared NCL semantics and adds a dedicated pointwise Conv1d fast path. - Updates Conv1d kernel inputs/outputs to use `x: (N, C_in, L)` and `out: (N, C_out, L_out)`. - Reworks the generic Conv1d tile grid to use `(L tile, C_out tile, N)` so the L dimension remains contiguous instead of forcing NCL data into the previous NLC-optimized flattened structure. - Flattens generic Conv1d weight as `(C_out, K * C_in)` before the TileLang kernel so the kernel can use a contiguous shared-memory copy for weights. - Adds an input tile fast path for fully in-bounds generic Conv1d tiles to reduce boundary checks on common interior tiles. - Adds `Conv1dPointwiseKernel` as a peer kernel to `Conv1dKernel`, similar to `Conv2d1x1Kernel`. - Moves pointwise dispatch to `Conv1dFwdOp`, matching the `Conv2dOp` pattern for selecting `Conv2d1x1Kernel` vs `Conv2dKernel`. ## Test plan - `python -m py_compile tileops/ops/convolution.py tileops/kernels/convolution.py tileops/kernels/__init__.py` - `python -m ruff check tileops/ops/convolution.py tileops/kernels/convolution.py tileops/kernels/__init__.py` - `git diff --check` - `python scripts/validate_manifest.py --check-op Conv1dFwdOp --levels schema,signature` - `python scripts/validate_manifest.py --check-op Conv1dBiasFwdOp --levels schema,signature` - `TMPDIR=/home/lyc/Project/TileOps-workspace2/.tmp TILELANG_CLEANUP_TEMP_FILES=1 python -m pytest tests/ops/test_convolution.py -k 'pointwise' -vvs` - `TMPDIR=/home/lyc/Project/TileOps-workspace2/.tmp TILELANG_CLEANUP_TEMP_FILES=1 python -m pytest tests/ops/test_convolution.py -k 'conv1d and not pointwise' -vvs` - `CUDA_VISIBLE_DEVICES=1 TMPDIR=/home/lyc/Project/TileOps-workspace2/.tmp TILELANG_CLEANUP_TEMP_FILES=1 python -m pytest benchmarks/ops/bench_convolution.py -k 'convtasnet-pointwise-k1-s1-fp16' -vvs` ## Benchmark Full Conv1d benchmark on NVIDIA H200, comparing against the parent before the layout change: | case | before TileOps ms | current TileOps ms | change | | --- | ---: | ---: | ---: | | convtasnet-pointwise-k1-s1-fp16 | 0.1898 | 0.2448 | +29.0% | | seanet-k3-s1-fp16 | 0.0175 | 0.0213 | +21.7% | | audio-downsample-k5-s2-fp16 | 0.0160 | 0.0199 | +24.4% | | seanet-stem-k7-s1-fp16 | 0.0431 | 0.0618 | +43.4% | | sequence-downsample-k3-s2-bf16 | 0.0097 | 0.0105 | +8.2% | | seanet-k3-s1-d2-fp16 | 0.0174 | 0.0192 | +10.3% | Pointwise rerun after adding `Conv1dPointwiseKernel`, on GPU1 with SM clock checked at 1830 MHz: | case | TileOps ms | Torch ms | TileOps TFLOPS | bandwidth TB/s | | --- | ---: | ---: | ---: | ---: | | convtasnet-pointwise-k1-s1-fp16 | 0.2448 | 0.5180 | 137.0594 | 0.8042 | Raw local reports: - `output_mid/conv1d_manifest/profile_before_24a036c.log` - `output_mid/conv1d_manifest/profile_current_full_gpu0.log` - `output_mid/conv1d_manifest/full_benchmark_comparison.md` - `output_mid/conv1d_manifest/profile_conv1d_pointwise_op_dispatch_gpu1.log` ## Regression This PR intentionally prioritizes manifest layout correctness over preserving the previous NLC-oriented Conv1d optimization. Current generic Conv1d performance is slower than the previous implementation because: - The corrected NCL layout no longer matches the original flattened NLC-optimized traversal. - The generic Conv1d path now carries additional index mapping and boundary handling for stride, padding, and dilation. - `weight.permute(0, 2, 1).contiguous()` is currently done in the Python forward path for generic Conv1d, which adds end-to-end overhead. - Reverting to a 2D flattened `(N * L)` grid reintroduces per-element `m_idx // out_l` and `m_idx % out_l` mapping and severely regresses K=7/S=1, so this PR keeps the 3D grid. The largest previous regression was K=1/S=1 pointwise Conv1d. This PR now routes that case through `Conv1dPointwiseKernel`, which removes generic stride/padding/dilation mapping from the TileLang kernel and lowers the pointwise result from the earlier corrected-layout measurement of `0.3015 ms` to `0.2448 ms`. Follow-up directions: - Further tune pointwise Conv1d to close the remaining gap against the old NLC-oriented baseline. - Add or cache a weight-packed representation to reduce repeated host-side weight layout conversion. - Continue tuning K=3/S=1 and K=7/S=1 after the correctness-aligned generic kernel is merged. - Use NCU on pointwise and K=3/K=7 cases to separate index overhead, memory load efficiency, and weight packing cost. |
|
|
Cao Ying
|
2de052e19a
|
[Refactor][Hygiene] strip development-process metadata from shipped source (#1290)
## Summary
Apply `.claude/rules/code-style.md` "no development-process metadata in
shipped source" across `tileops/` `tests/` `benchmarks/` `scripts/`.
Rewrite `pytest.skip`/`xfail` reasons, docstrings, comments to describe
invariants instead of referencing issue/PR numbers and AC labels.
Prose-only — no behavior change.
## Scope (24 files)
- 16 `pytest.skip`/`xfail` reason rewrites (tilelang 0.1.9 batch, rope,
attention, recurrent linear-attention) + shared
`TILELANG_019_SKIP_REASON` in two conftests
- AC-label strips in 9 test files / 1 bench /
`scripts/validate_manifest.py`
- `TODO(#NNNN)` strips in `tileops/kernels/convolution.py` (×2),
`reduction/logsumexp.py`, `reduction/softmax.py`
- Historical PR refs in `tests/test_validate_manifest.py` rewritten to
describe the contract
Excluded: `tileops/manifest/`, `docs/`, `.claude/` (out of rule scope).
`FIXME(staged-rollout)` blocks already compliant. Single-digit "Finding
#N" labels not matched by the rule's `#[0-9]{3,}` regex.
## Test plan
- [x] discovery scan returns 0 matches
- [x] `pytest tests/ --co -q` — 2752 collected
- [x] touched test files: 186 + 53 passed
- [x] pre-commit clean
Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
ChongLi
|
f905cd3ba6
|
[BugFix][CONV] Disable async copy for conv data loads (#1119)
Closes #1105 ## Summary - Disable TileLang automatic async-copy injection for the generic Conv2d and Conv3d kernels with `pass_configs={"tl.enable_async_copy": False}`. - Keep convolution weight tile loads as `T.copy(...)` so they remain eligible for TileLang TMA lowering. - Add `TODO(#1105)` comments at both JIT sites explaining that this is a temporary workaround until TileLang fixes scalar `cp.async` widening for vectorized manual data loads. - Re-enable the previously skipped Conv2d and Conv3d correctness cases now that the affected paths compile and pass. ## Test plan - [x] pre-commit passed during commit and push hooks - [x] `CUDA_VISIBLE_DEVICES=<target_gpu> TMPDIR=<repo>/.tmp/tvm_tmp conda run -n tileops-dev python -m pytest tests/ops/test_convolution.py -vvs` Result: ```text 36 passed in 73.41s (0:01:13) ``` ## Structural Readiness All checks passed. ## Benchmark **Environment**: NVIDIA H200 fixed at 1830 MHz, CUDA 12.8, PyTorch 2.9.0+cu128, TileLang 0.1.9 ### Conv2d All 13 official Conv2d benchmark cases completed successfully with autotuning enabled. Under the same test conditions, compared with the old-code TileLang 0.1.8 baseline: | Op | Shape group | dtype | Cases | Average latency change | |----|-------------|-------|------:|-----------------------:| | Conv2d | All official benchmark shapes | fp16/bf16 | 13 | -3.6% | | Conv2d | Non-1x1 convolution shapes | fp16/bf16 | 7 | -3.8% | | Conv2d | 1x1 convolution shapes | fp16/bf16 | 6 | -3.5% | **Takeaways:** - Keeping weight loads as `T.copy(...)` preserves the TMA-friendly path and avoids the large performance regression seen with guarded scalar weight loads. - Disabling automatic async-copy injection for generic Conv2d avoids the TileLang 0.1.9 scalar `cp.async` widening failure without reducing benchmark throughput under the measured conditions. - PyTorch 2.10.0+cu128 **Benchmark command:** ```bash CUDA_VISIBLE_DEVICES=<target_gpu> TMPDIR=<repo>/.tmp/tvm_tmp \ conda run -n tileops-dev python -m pytest \ benchmarks/ops/bench_convolution.py::test_conv2d_bench -vvs ``` ## Regression ### Root Cause Analysis Issue #1105 is caused by a TileLang 0.1.9 lowering/codegen path, not by NVCC. In the generic convolution kernels, manual `data_shared` loads are scalar fp16/bf16 global-to-shared stores inside a vectorized `T.Parallel` loop. When the weight `T.copy(...)` is recognized as a TMA producer, TileLang's warp-specialized producer scheduling can also mark the manual SIMT data producer for async-copy handling. The relevant TileLang 0.1.9 source locations are: - `src/transform/producer_consumer_ws.cc`: classifies generic `T.copy` as a TMA producer when bulk-load checks pass, and annotates SIMT producers for pipeline-managed async-copy scheduling in mixed producer groups. - `src/transform/lower_ptx_async_copy.cc:80-89`: tracks `ForKind::kVectorized` loop extents in `current_vectorized_lanes_`. - `src/transform/lower_ptx_async_copy.cc:401-412`: validates the expected final `cp.async` byte width using `effective_lanes * dtype_bits * current_vectorized_lanes_`. - `src/transform/lower_ptx_async_copy.cc:489-508`: creates `tl::ptx_cp_async` with the per-access logical element count. - `src/transform/vectorize_loop.cc:702-767`: is expected to widen vectorized `tl::ptx_cp_async` calls by multiplying the logical element count by the vector size. - `src/target/codegen_cuda.cc:59-90`: performs the final legality check and rejects `tl::ptx_cp_async` widths outside `{4, 8, 16}` bytes. For the Conv2d `smoke-fp16-3x3` reproducer, the scalar fp16 load has `effective_lanes=1`, while the surrounding vectorized loop contributes `current_vectorized_lanes_=8`. The pass therefore predicts: ```text 1 * 16 bits * 8 = 128 bits = 16 bytes ``` That predicted width is legal for PTX `cp.async`, so TileLang injects `tl::ptx_cp_async`. However, the injected intrinsic still carries the scalar logical element count (`num_elems=1`) and relies on a later vectorization pass to widen it to `num_elems=8`. In the failing path that widening does not survive, so CUDA codegen eventually sees: ```text num_elems=1, dtype=fp16 => 2 bytes ``` and correctly rejects it: ```text tl::ptx_cp_async requires a final PTX byte width in {4, 8, 16}, but got 2 ``` This patch avoids the faulty automatic `cp.async` rewrite for the generic Conv2d/Conv3d manual data-load paths. It deliberately does not disable TMA or replace weight `T.copy(...)`, because doing so regresses larger convolution cases where the weight tile load benefits from TMA. ## Additional context The ideal long-term fix should be in TileLang: - `InjectPTXAsyncCopy` should not leave a scalar-width `tl::ptx_cp_async` in IR when its legality depends on a later vectorization pass. - If later vectorization cannot widen the intrinsic to a legal 4/8/16-byte transfer, TileLang should fall back to normal load/store before CUDA codegen. - CUDA codegen should remain the final validation layer, not the first place where this cross-pass contract failure is discovered. Once TileLang guarantees that failed widening falls back safely, the `tl.enable_async_copy=False` workaround in Conv2d and Conv3d can be removed. |
|
|
ChongLi
|
38e3b0e708
|
[Refactor][CONV1D] Add dilation support to forward kernel (#1038)
## Summary Closes #1031. Related #853. Related #1037. This PR adds PyTorch-compatible `dilation` support to the Conv1d forward kernel path and keeps the change scoped away from grouped convolution and manifest status promotion. It also aligns the two Conv1d forward variants with their manifest semantics: - `Conv1dFwdOp` is now the no-bias variant and exposes `forward(input, weight)`. - `Conv1dBiasFwdOp` is now the bias-required variant and exposes `forward(input, weight, bias)`. Implementation details: - Adds `dilation` / `dilation_l` plumbing through `Conv1dFwdOp`, `Conv1dBiasFwdOp`, `Conv1dKernel`, `_conv1d_wrapped_kernel`, and `_conv1d_kernel`. - Updates Conv1d output length to `(L_in + 2 * padding - dilation * (kW - 1) - 1) // stride + 1`. - Updates kernel input indexing to `il = ol * stride_l + kw * dilation_l - pad_l`. - Keeps `groups=1` as the only implemented path while validating the manifest-declared groups constraints and rejecting unsupported grouped Conv1d explicitly. - Keeps Conv1d spatial params aligned with the manifest by accepting `int | tuple[int]`. ## Test plan - `python -m py_compile tileops/ops/convolution.py tileops/kernels/convolution.py tests/ops/test_convolution.py benchmarks/ops/bench_convolution.py` - `python -m ruff check tileops/ops/convolution.py tileops/kernels/convolution.py tests/ops/test_convolution.py benchmarks/ops/bench_convolution.py` - `python -m pytest tests/ops/test_convolution.py --collect-only -q` - `python -m pytest tests/ops/test_convolution.py -vvs -k conv1d` - `python -m pytest benchmarks/ops/bench_convolution.py --collect-only -q -k conv1d_bench` - `python -m pytest tests --collect-only -q` - `python scripts/validate_manifest.py --check-op Conv1dFwdOp` - `python scripts/validate_manifest.py --check-op Conv1dBiasFwdOp` ## Benchmark Adds a focused Conv1d benchmark case with `dilation=2` while preserving existing `dilation=1` benchmark coverage. The benchmark still uses the existing hand-written benchmark harness; #1037 tracks making Conv1d benchmarks fully manifest-driven so manifest validation no longer emits L4 benchmark warnings. ## Regression Existing Conv1d `dilation=1` behavior remains covered by the original Conv1d test and benchmark cases. New tests cover `dilation=2`, no-bias `Conv1dFwdOp`, and bias-required `Conv1dBiasFwdOp`. ## Additional context `python scripts/validate_manifest.py --check-op Conv1dFwdOp` and `python scripts/validate_manifest.py --check-op Conv1dBiasFwdOp` now pass, but still report warnings for missing generated shape/dtype parity methods and the non-manifest-driven benchmark path. The benchmark warning is intentionally tracked separately in #1037. |
|
|
ChongLi
|
3ac048f43f
|
[BugFix][Conv] Include output shared buffer in memory estimate (#1019)
Closes #932 ## Summary - Include the convolution output shared buffer in `conv_shared_memory_bytes`. - Keep kernel logic unchanged; only the autotune shared memory estimator changes. ## Test plan - [x] pre-commit hooks passed during commit - [ ] `scripts/validate.sh --pre` not run because `scripts/validate.sh` is not present in this checkout - [ ] pytest not run; no related unit test was added per request ## Regression - This tightens autotune filtering for convolution configs that would exceed shared memory once `out_shared` is included. |
|
|
ChongLi
|
3c88b48df1
|
[BugFix][Conv] thread has_bias through conv2d 1x1 wrapper (#1018)
Closes #933 ## Summary - Thread `has_bias` through the Conv2d 1x1 torch custom op wrapper and fake registration. - Pass `self.has_bias` from `Conv2d1x1Kernel.forward` so the compiled kernel can skip the bias branch when disabled. ## Test plan - [x] `python -m pytest tests/ops/test_convolution.py` (24 passed, 8 skipped) - [x] pre-commit hooks passed during commit ## Benchmark - Not run. This only fixes wrapper plumbing so the existing no-bias kernel path is reachable under `torch.compile`. ## Regression - Verified existing convolution op tests still pass. - `scripts/validate.sh --pre` was unavailable in this checkout (`No such file or directory`), so title and branch patterns were checked directly against `.claude/conventions/types.sh`. |
|
|
Cao Ying
|
bdbd7a8324
|
[Refactor] Rename base modules to <module>_base.py convention (#950)
## Summary Unify base module file naming to `<module>_base.py` convention across four modules: - `workloads/base.py` -> `workloads/workload_base.py` - `tileops/ops/op.py` -> `tileops/ops/op_base.py` - `tileops/kernels/kernel.py` -> `tileops/kernels/kernel_base.py` - `benchmarks/benchmark.py` -> `benchmarks/benchmark_base.py` All imports, string references, and doc paths updated. No logic changes — pure rename + import propagation. Closes #947 ## Test plan - [x] **AC-1**: All four files renamed to <module>_base.py -- pass - [x] **AC-2**: All imports updated, no broken references -- pass - [x] **AC-3**: docs/testing.md path references updated -- pass - [x] **AC-4**: Pre-commit and existing tests pass -- pass **Test results**: 2368/2368 tests passed, 0 failed --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
a83d6bfa46
|
[Refactor] Flatten kernels/convolution/ into single module (#930)
## Summary - Merge `tileops/kernels/convolution/` directory (5 files: `__init__.py`, `common.py`, `conv1d.py`, `conv2d.py`, `conv3d.py`) into a single flat `tileops/kernels/convolution.py` - No code logic changes — pure file consolidation - Follows the same flattening pattern applied to other single-purpose kernel directories in #928 ## Test plan - [x] `python -c "from tileops.kernels.convolution import Conv1dKernel, Conv2d1x1Kernel, Conv2dKernel, Conv3dKernel"` resolves - [x] `pytest tests/ops/test_convolution.py` — 29 passed, 0 failed - [x] Pre-commit lint clean ## Follow-up - #932 — `conv_shared_memory_bytes` underestimates shared memory (missing `out_shared` buffer) - #933 — `_conv2d_1x1_wrapped_kernel` hardcodes `has_bias=True` - #934 — manifest `source.kernel` paths still reference removed `convolution/` directory (depends on this PR) Suggestions: bot reviewers suggested replacing manual output-copy loops with `T.copy` (lines 133, 513, 1019) — **rejected**: those loops contain bounds checks that `T.copy` would bypass, causing out-of-bounds writes when tile dimensions don't evenly divide the output. Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |