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> |
|
|
ChongLi
|
5837ea0ee1
|
[Feat][CONV_POOL] Infer conv pool shape metadata from inputs (#1677)
Closes #1673 ## Summary - Refactor Conv1d/2d/3d APIs to infer tensor-derived metadata from `forward(input, weight)`: input shape, output channels, weight-derived kernel size, and dtype. - Refactor AvgPool1d/2d/3d APIs to infer input shape and dtype from `forward(input)` while keeping pooling semantics such as `kernel_size`, `stride`, `padding`, `ceil_mode`, `count_include_pad`, and `divisor_override` in the constructor. - Remove the old explicit constructor metadata path for tensor-derived values; removed arguments now fail naturally as unexpected keyword arguments. - Add lazy kernel caching and forward-bound roofline metadata for Conv and AvgPool, and update tests and benchmarks to use the inferred API. - Preserve Conv padding as a constructor semantic parameter and keep resolved runtime padding separate from the user-provided padding mode/value. ## Test plan - [x] `python -m py_compile tileops/ops/convolution.py tileops/ops/pool.py tests/ops/test_convolution.py tests/ops/test_pool.py benchmarks/ops/bench_convolution.py benchmarks/ops/bench_pool.py` - [x] `python -m pytest tests/ops/test_convolution.py tests/ops/test_pool.py -q --tb=short` - [x] `python scripts/validate_manifest.py --family convolution` - [x] `python scripts/validate_manifest.py --family pool` ## Benchmark - Not run. Benchmark constructors were updated to use the input-inferred API, but this PR does not include benchmark result changes. ## Regression - Verified Conv pre-forward `_infer_output_shapes` works for Conv1d/2d/3d. - Verified AvgPool constructors require explicit `kernel_size`. |
|
|
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 | |
|
|
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
|
a15dd9c880
|
[Feat][convolution] Align conv fwd op entrypoints (#1518)
Closes #1507 ## Summary - Adds manifest-aligned Conv2d/Conv3d forward op classes and bias variants. - Removes legacy Conv2dOp/Conv3dOp public exports and updates tests/benchmarks to use the new classes. - Keeps dilation, groups, and layout follow-ups out of this pass; tracked separately in #1519, #1520, and #1521. - Tracks current benchmark follow-ups separately in #1522 and #1523. ## Test plan - [x] pre-commit passed - [x] python -c "from tileops.ops import Conv2dFwdOp, Conv2dBiasFwdOp, Conv3dFwdOp, Conv3dBiasFwdOp; print('imports ok')" - [x] python scripts/validate_manifest.py - [x] python -m pytest tests/ops/test_convolution.py --collect-only -q - [x] python -m pytest tests/ops/test_convolution.py -vvs (41 passed) ## Additional context - scripts/validate.sh is not present in this checkout, so the skill-specific pre/post validation command could not be run. - python -m pytest benchmarks/ops/bench_convolution.py -vvs does not pass yet: Conv1d seanet-k3-s1-d2-fp16 fails autotuning, and Conv2d/Conv3d benchmark cases are skipped. These are tracked in #1522 and #1523. - Current string padding support was checked for Conv2d/Conv3d with padding="valid" and padding="same"; both construct successfully for representative supported shapes. |
|
|
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. |
|
|
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. |
|
|
Ang Gao
|
d5b3279a16
|
[CI] Add temporary TileLang 0.1.9 skips (#1043)
## Summary - bump TileLang references from `0.1.8` to `0.1.9` - add a centralized temporary TileLang 0.1.9 skip set in `tests/conftest.py` - update existing temporary skip messages from the old `5f70374c` tracker to #1039 - keep the suite green while the skipped buckets are fixed and restored under #1041 The skip set currently covers 10 file-level skips, 2 nodeid-prefix skips, and 71 exact nodeid skips. In response to review, `test_mha.py` and `test_mha_decode.py` are no longer skipped at file level, so MHA forward and the short-context MHA decode case continue to run. `test_gqa.py` remains file-level skipped because narrowing it exposed a smoke-only hang in `test_gqa_fwd[1-1024-8-4-64-False-dtype0-False]` under TileLang 0.1.9. Closes #1042 Refs #1039 Refs #1041 ## Validation - `git diff --check` - `python3 -m ruff check tests/conftest.py tests/test_ci_venv_hash.py pyproject.toml` inside `tileops-runner:nightly-tl019-fullstack-no-tileops-ldfix` - TileLang 0.1.9 full validation with this skip set: `2160 passed, 232 skipped, 196 warnings in 40.21s` |
|
|
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. |
|
|
Ang Gao
|
45cb188970
|
[Test][CI] Temporarily skip known failures under TileLang 5f70374c (#1001)
## Summary - temporarily skip the currently known failing buckets under the validated TileLang `5f70374c48bb1d52e9260e995bf7adbcd64341e5` environment - keep this PR narrowly focused on test scope reduction only - leave published dependency / README updates to follow-up issue #1000 ## Why After validating the TileLang `5f70374c48bb1d52e9260e995bf7adbcd64341e5` environment tracked in #999, the remaining full-suite failures are no longer generic runtime instability. They have been isolated into a bounded set of known failing buckets. This PR keeps CI / validation usable in the short term by temporarily marking those buckets, while the underlying kernels/operators are fixed separately. Detailed error-cascade isolation remains documented in #980. ## Included temporary skips - gated deltanet low-precision decode failures - GLA decode and multi-step correctness coverage, where fp32 currently returns NaNs and fp16/bf16 remain tracked as low-precision launch failures - attention decode / backward failure buckets still failing under TileLang `5f70374c` - selected convolution / fp8 / rope / mhc_pre failures ## Validation Command: ```bash CUDA_VISIBLE_DEVICES=0 \ TILELANG_CLEANUP_TEMP_FILES=1 \ TMPDIR=/home/ga/tltmp \ /home/ga/TileOPs-upstream-cu128-t210/.venv-cu128-t210-ded6/bin/python -m pytest -q tests ``` Result: ```text 2180 passed, 89 skipped, 16 warnings in 162.99s ``` ## Follow-ups - #980 error cascade / CUDA poisoning investigation record - #999 validated TileLang `5f70374c` env bring-up tracking - #1000 published TileLang dependency and README follow-up |
|
|
Cao Ying
|
84fa4b5a70
|
[Refactor][Benchmark] Narrow gen_inputs return types across workload protocols (#970)
## Summary - Convert `Tuple[...]` → `tuple[...]` and `List[...]` → `list[...]` (PEP 585) repo-wide, remove unused `typing` imports - Widen protocol/base `gen_inputs()` from `-> Any` to `-> tuple[Any, ...]` (`WorkloadBase`, `InputGeneratingWorkload`); narrow `BinaryWorkload` and `_UnaryWorkload` to `-> tuple[torch.Tensor, ...]` - Widen `bench_kernel` args and `profile` *inputs to `Any` to match existing non-tensor usage - Add runtime guard in `bench_kernel` against non-tuple args (bare tensor silently iterates over rows) - Fix `FusedTopKTest.gen_inputs()` returning bare tensor instead of tuple - Fix `Optional` annotations for nullable returns in `mean_pooling.py` and `mamba.py` - Add missing return type annotations in `batch_norm.py`, `moe.py` Closes #956 ## Test plan - [x] `pre-commit run --all-files` passes - [x] No `Tuple` or `List` imports from `typing` remaining repo-wide - [x] All `workloads/**/*.py` `gen_inputs()` methods have explicit return type annotations - [x] `bench_kernel` raises `TypeError` on non-tuple args ## Follow-up No follow-up issues or suggestions. --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
ChongLi
|
93340add26
|
[Test][TESTING] Add missing per-dtype smoke coverage (#952)
## Summary - add representative `smoke` cases for each supported dtype in the issue-listed test fixtures - keep the smoke cases at the front of each parametrized list and preserve `tune=False` - update the current-file equivalents for the issue paths, including `tests/ops/test_engram.py` and `tests/ops/test_pool.py` Closes #945 ## Test plan - `python -m pytest tests --collect-only -q` - `python -m pytest tests/ops/test_activation.py tests/ops/test_fused_gated.py tests/ops/test_binary_arith.py tests/ops/test_engram.py tests/ops/test_pool.py --collect-only -q` - `python scripts/test_node_delta.py tests/ops/test_activation.py tests/ops/test_fused_gated.py tests/ops/test_binary_arith.py tests/ops/test_engram.py tests/ops/test_pool.py` ## Regression - smoke-only relabeling stays scoped to representative dtype-dispatch coverage - all added smoke cases keep `tune=False` - collected `smoke` cases grow from `1072` on `main` to `1597` on this branch (`+525`) - total collected test nodes grow from `2383` on `main` to `2410` on this branch (`+27`, `+1.1%`) ## Additional context - the issue references `tests/ops/test_engram_fwd.py` and `tests/ops/test_avg_pool1d.py`; in the current workspace those correspond to `tests/ops/test_engram.py` and `tests/ops/test_pool.py` ## Follow-up No follow-up issues. Suggestion: `WelfordNonAlignedMultiDimFixture` `flat63_fp16` was promoted to smoke but adds shape coverage not dtype coverage — move back to `full`, keep `flat63_bf16` as the bf16 smoke case. |
|
|
Cao Ying
|
6a61ba14d9
|
[Refactor] Align kernels/tests/benchmarks structure with ops layout (#928)
## Summary Align `kernels/`, `tests/`, and `benchmarks/` file structure 1:1 with the post-#890 `ops/` layout. Eliminates unnecessary directory nesting and consolidates trivially small files. **Changes (166 files, 11 commits):** 1. Consolidate attention kernel directories into `kernels/attention/` 2. Rename `kernels/conv/` to `kernels/convolution/` 3. Drop `_fwd` suffix from mamba kernel filenames 4. Flatten `kernels/norm/` subdirectories 5. Flatten `kernels/reduction/` subdirectories 6. Flatten single-file kernel directories 7. Consolidate tests and benchmarks to mirror kernel layout 8. Delete stale `CHANGELOG.md` Closes #924 ## Test plan - [x] **AC-1**: All `__init__.py`, `ops_manifest.yaml`, and internal imports updated for moved files -- Manifest path audit: checked 208 paths, 0 missing. Runtime validation passed. - [x] **AC-2**: `from tileops.kernels import ...` resolves for all moved kernels -- 73 symbols exported and resolved with 0 failures. - [x] **AC-3**: Full test suite passes -- `pytest tests` -> 2361 passed, 22 skipped, 0 failed (236.78s). - [x] **AC-4**: No orphaned directories remain after migration -- `find` produced no empty directories. ## Follow-up - #929 — Update stale kernel path reference in docs/manifest.md (doc drift from this refactor) --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |