Skip to content

Releases: pytorch/helion

v1.2.0

Choose a tag to compare

@oulgen oulgen released this 22 Jun 16:27
d55389b

What's Changed

  • [Pallas] Add Pallas interpret CI job (revival of #1938) by @norx1991 in #2491
  • Disable rope-bwd from nightly perf CI for now by @choijon5 in #2494
  • torch_tpu: update pin to f863c2bacd52aaef02599a0e80f53a24fe9431c0 by @cota in #2528
  • Improve CuTe RMS norm reduction path by @oulgen in #2540
  • [pallas] support dynamic shapes in indirect gather lowering by @cota in #2493
  • Enable >32-thread persistent reduction on CuTe for RMS norm perf by @oulgen in #2541
  • Cache CuTe launcher imports and dtype mapping in hot path by @oulgen in #2542
  • Add CuTe reduction-kernel autotuner seed by @oulgen in #2543
  • Fuse two-pass loads of x in CuTe reduction kernels by @oulgen in #2544
  • Add wide-chunk reduction seed and tidy fuse-two-pass-loads helper by @oulgen in #2545
  • [Pallas] correct RNG seed buffer argument ordering in launcher by @AmesingFlank in #2539
  • [cutedsl] move CuTe state out of DeviceFunction by @jansel in #2497
  • [cutedsl] split CuTe tcgen05 config logic by @jansel in #2498
  • [cutedsl] replace CuTe golden codegen tests by @jansel in #2499
  • [cutedsl] handle CuTe edge epilogue autotune by @jansel in #2500
  • [cutedsl] support flat edge TMA stores with scalar fallback by @jansel in #2501
  • [cutedsl] enable clustered edge TMA and seed configs by @jansel in #2502
  • [cutedsl] add hybrid TMA store pipeline for edge tiles by @jansel in #2503
  • [cutedsl] tune edge K-tail epilogue placements by @jansel in #2504
  • [cutedsl] split scheduler edge epilogues by @jansel in #2505
  • [cutedsl] add clc scheduler diagnostics and lifecycle fixes by @jansel in #2506
  • [cutedsl] add aux TMA load mode by @jansel in #2507
  • [cutedsl] seed aux TMA and clc persistence candidates by @jansel in #2508
  • [cutedsl] retile and pack aux epilogue by @jansel in #2509
  • [cutedsl] vectorize and stage aux epilogue stores by @jansel in #2510
  • [cutedsl] add staged scheduler mailbox diagnostic by @jansel in #2511
  • [cutedsl] skip static full tma fallback by @jansel in #2512
  • [cutedsl] add tcgen05 lifecycle context and kloop ownership by @jansel in #2513
  • [cutedsl] introduce pure matmul object model for tcgen05 by @jansel in #2514
  • [cutedsl] add explicit tcgen05 epilogue tile and two-cta kloop by @jansel in #2515
  • [cutedsl] cache cute launch arguments by @jansel in #2516
  • [cutedsl] add guarded flat tcgen05 role coordinates by @jansel in #2517
  • [cutedsl] add tcgen05 tvm ffi launch with seed configs by @jansel in #2518
  • [cutedsl] introduce target1 clc scheduler object by @jansel in #2519
  • [cutedsl] add target1 direct entry path by @jansel in #2520
  • [cutedsl] Update benchmarking code by @jansel in #2521
  • [Pallas] int32 gather lowering by @thcmbs in #2550
  • [Autotuner] Cap inner-tile block size by outer tile by @thcmbs in #2478
  • Use fixed seed for nightly benchmarks by @choijon5 in #2561
  • torch_tpu: update pin to 5281ef191211f27bde9144cb17755c4f6a89f8fa by @cota in #2562
  • Emit vectorized 128-bit loads for CuTe inner reductions by @oulgen in #2546
  • [Pallas] Fix mask broadcast on size-1 tensor dimensions by @AmesingFlank in #2556
  • Enable vec loads for bf16/fp16 CuTe reductions via Uint16 bitcast by @oulgen in #2547
  • [cutedsl] extend tcgen05 direct-entry path to additional targets by @jansel in #2551
  • [cutedsl] add tcgen05_aux_stages and tcgen05_consumer_regs autotune knobs by @jansel in #2552
  • [cutedsl] stage dynamic_persistent scheduler object infrastructure by @jansel in #2553
  • [cutedsl] Expand hillclimber benchmarking to 32 targets by @jansel in #2554
  • Fix XPU CI SYCL runtime setup by @choijon5 in #2566
  • (Onboarding task) examples: add XSA (exclusive self-attention) kernel by @calebmkim in #2564
  • Vec-aware two-pass load fusion for CuTe reductions by @oulgen in #2548
  • [autoreview] Add --codex-no-sandbox flag for hosts without bubblewrap by @choijon5 in #2568
  • Support tensor_descriptor + epilogue subtile by @ethche in #2557
  • [CI] Bump Pyrefly, ruff, cute by @oulgen in #2569
  • [cute] skip test_xsa on cute by @oulgen in #2581
  • [cute] Softmax perf: vec hoist, fuser alias, hoist warp reduce, heuristics by @oulgen in #2570
  • [cute] V=8 fp16/bf16 via 4+4 split; multi-row investigation tests by @oulgen in #2571
  • [cute] Merge sibling constexpr V-loops + elide warp-reduce dtype wraps by @oulgen in #2572
  • [cute] Warp-per-row layout for softmax-shape reductions by @oulgen in #2573
  • Reject tensor_descriptor indexing when block size exceeds tensor dim (#2555) by @mengluy0125 in #2555
  • [cute] LICM for reciprocals: hoist 1/divisor out of inner loops by @oulgen in #2574
  • [Pallas] Fix synchronize_device skipping sync for tuple returns on TPU by @thcmbs in #2580
  • [Pallas] remove jax_export_ignore_forward_compatibility setting by @cota in #2452
  • [cute] Cleaner LICM: alias DCE + FMA-friendly scale hoist by @oulgen in #2575
  • [Pallas] Add support for non zero dim in gather by @thcmbs in #2558
  • [cute] Software-pipeline inner vec loads to hide HBM latency by @oulgen in #2576
  • [cute] Optional carried-only gate for the load pipeline pass by @oulgen in #2577
  • [cute] Rewrite online softmax_two_pass → equivalent 3-pass form by @oulgen in #2578
  • [test] Split cute compiler-pass tests by pass, drop kernel name from filenames by @oulgen in #2579
  • [CI] bump cutlass by @oulgen in #2610
  • A jagged_hstu_attention example that works on Pallas TPU by @AmesingFlank in #2218
  • torch_tpu: update pin to adbe52d4a6082d23bc40f630b335d37379b604d2 by @cota in #2602
  • test_examples: skip test_jagged_dense_add on Pallas by @cota in #2607
  • [metal] Auto-cap threadgroup size via CuteND/CuteFlattenedTileStrategy by @aditvenk in #1855
  • [cosmetic/UI] Quieter console output and nicer progress formatting by @choijon5 in #2598
  • [Pallas] Propagate inner tile alignment min_size to bounding outer tiles by @thcmbs in #2559
  • [Pallas] Suppress Pallas RNG x64 warning spam in random tests by @thcmbs in #2617
  • [metal] Add DeviceIR MPPGraph lowering for aten matmul by @aditvenk in #2472
  • [Pallas] Avoid bool reshape lowering by @thcmbs in #2618
  • ci: bump TPU jax/jaxlib pin to 0.10.1 by @thcmbs in #2621
  • [Pallas] Fix indirect gather by @thcmbs in #2619
  • [Pallas] Add generic fallback scatter by @thcmbs in #2620
  • Migrate H100 CI to OSDC runners (pytorch/pytorch cu128 base image) by @huydhn in #2625
  • torch_tpu: update pin to dd5b8dd76b1eb0175db812c35b0d2f52ca1f4e3b by @cota in #2633
  • [cute] Fix benchmark errors on CI for cute by @oulgen in #2622
  • [cute] Require cuda13+ by @oulgen in #2623
  • [compile] add ast backend customization hook by @hinriksnaer in #2624
  • torch_tpu: update pin to e2250c4179c5022f7d0a4007dcc746f60fde64bb by @cota in https://github.com/pytorch...
Read more

v1.1.0

Choose a tag to compare

@oulgen oulgen released this 20 May 23:32
d807139

What's Changed

  • [Pallas] Use TPU-aware synchronize_device across autotuner and testing by @norx1991 in #1944
  • Refactor if-else branches, unblock test_if_new_variable_in_static_range for Pallas TPU by @AmesingFlank in #1935
  • [Pallas] Register pallas_loop_type only when inner loops exist by @norx1991 in #1915
  • [Pallas] Fix emit_pipeline program_id mapping with loop_order reordering by @norx1991 in #1916
  • [Pallas] Add expected-failure test for non-power-of-2 RDIM size by @norx1991 in #1945
  • Update docs: add missing API entries by @choijon5 in #1948
  • [Benchmark] Add compile time measurement to CI benchmarks by @choijon5 in #1952
  • use is_symm_mem_tensor if the API is available by @shunting314 in #1933
  • [Pallas] When not doing tiling for entire kernel, use explicit BlockSpecs instead of None by @AmesingFlank in #1960
  • Fix FROM_BEST_AVAILABLE matching with hl.specialize() after #1883 by @fulvius31 in #1940
  • [Pallas] Fix indexing scalars using SMEM memory space by @AmesingFlank in #1955
  • [Pallas] Add xfail tests for scalar .begin index not collapsing dims by @norx1991 in #1971
  • Skip test_hl_rand_mixed_argument_order on MTIA due to unaligned address crash by @karthickai in #1977
  • Fix _supports_maxnreg() to guard against non-CUDA backends by @karthickai in #1981
  • Skip TestRandomPhiloxParity class on MTIA (#1979) by @karthickai in #1982
  • [Pallas] Fix acccesing tensors using index from hl.grid() by @AmesingFlank in #1956
  • [Pallas] Support using traced size-1 tensor as condition predicate, unblocking test_if_arg_indexed_scalar by @AmesingFlank in #1957
  • [Pallas] Fix atomic_add dtype cast and VMEM preload for fori/pipeline launchers by @thcmbs in #1966
  • [Pallas] Use exact RDIM size instead of next-power-of-2 by @norx1991 in #1954
  • [Pallas] Add support for accessing tensors with the pattern of tile.index + offset /.id/.begin/.end by @AmesingFlank in #1968
  • [metal] Reuse Inductor's MetalOverrides for MSL expression emission by @aditvenk in #1853
  • [metal] Add Metal codegen handlers for load, store, and mask_to by @aditvenk in #1854
  • [Pallas] Enable a subset of test_grid tests for Pallas by @AmesingFlank in #1985
  • [Pallas] Add a test for accessing tensors with hl.grid() index + offset by @AmesingFlank in #1988
  • Relax rms_norm example tolerance for Pallas bf16 by @thcmbs in #1983
  • [autotuner] Introduce BenchmarkProvider abstraction for kernel benchmarking by @hinriksnaer in #1928
  • [Pallas] Use HBM BlockSpecs for output-only tensors to save VMEM by @norx1991 in #1984
  • Emit autotune failure summary warnings by @allgather in #1994
  • [Pallas] Remove unused is_device_loop variable in _pallas_index_str by @AmesingFlank in #1990
  • [Pallas] Fix BlockSpecs for 2D tl.grid([m, n]), unblocking test_scalar_access_hl_grid_2d by @AmesingFlank in #1986
  • Reduce some autotuner overhead without changing kernel behavior by @svdrecbd in #1885
  • [compile] Add pre_codegen hook to Backend ABC by @hinriksnaer in #1976
  • Adding AMD Mi350x machines to CI with new labels. by @umechand-amd in #1835
  • [Pallas] Fix accessing tensor via hl.grid() index within a device loop, unblocking test_scalar_access_hl_grid_2d_nested by @AmesingFlank in #1989
  • Support mtia in LocalAutotuneCache by @Hamlin-Li in #1996
  • removed stale comments by @hinriksnaer in #1997
  • Fix invalid default config for kernels with large tensor numel by @fulvius31 in #1839
  • chore: Bump actions/github-script from 8 to 9 by @dependabot[bot] in #2000
  • Fix _benchmark dropping configs that fail compilation by @fulvius31 in #1942
  • [metal] Add MSL AST walker for Python-to-C++ translation by @aditvenk in #1794
  • Enable tensor_descriptor based atomic ops by @ethche in #1953
  • [metal] Add @metal_jit decorator for AST-to-MSL compilation by @aditvenk in #1991
  • Epilogue subtiling: store indexing fix, example, and tuple output support in run_example by @choijon5 in #1907
  • [Compiler] Added backend registry by @hinriksnaer in #1967
  • [metal] Wire @metal_jit into MetalBackend and simplify launcher by @aditvenk in #1992
  • [Pallas] Skip trivial reduction mask when RDIM size equals actual dim by @norx1991 in #1993
  • [Pallas] Fix TPU min_dot_size for matmul autotuning by @norx1991 in #1999
  • chore: Bump actions/upload-pages-artifact from 4 to 5 by @dependabot[bot] in #2015
  • Make _get_tile_with_offset_info accept a torch.fx.Node as arg, instead of the entire CodegenState by @AmesingFlank in #2005
  • [metal] Skip tests better to pass in internal CI by @aditvenk in #2016
  • Add to_code() as backend-agnostic alias for to_triton_code() by @norx1991 in #2012
  • [Pallas] Fix fori_loop multi-dim index decomposition with nested loops by @norx1991 in #1917
  • custom config filter by @shunting314 in #1847
  • [Pallas] Reject 64-bit input tensors and fix tiling ZeroDivisionError by @norx1991 in #1950
  • Bump CI libtpu to 0.0.40 by @AmesingFlank in #2017
  • [Pallas] Make test_tensor_access_tile_index_offset more meaningful and reveal codegen issue in Pallas backend by @AmesingFlank in #2006
  • [Pallas] [NFC] Dedup _jnp_dtype_map helper by @thcmbs in #2021
  • [Pallas] Add a plan_tiling pre_codegen pass to make more consistent tiling and indexing decisions by @AmesingFlank in #2007
  • [Compiler] Add reserved_launch_param_names to Backend ABC by @hinriksnaer in #1970
  • [Pallas] Remove non-tiled fallback paths which are no longer used after #2007 by @AmesingFlank in #2026
  • [TPU][Pallas] Add lower bound analytical VMEM estimation and OOM guard for Pallas launchers by @yarongmu-google in #2024
  • [Pallas] Fix codegen for slice indexing when there are squeezed dimensions by @AmesingFlank in #2027
  • [CI] Pin Triton-to-tile-IR to last known-good commit by @norx1991 in #2028
  • [Pallas] Add non-DMA fori_loop fallback for DMA-unaligned inner blocks by @thcmbs in #1969
  • Skip test_matmul_smaller_than_min_dot_size on MTIA (#2030) by @will-cromar in #2030
  • docs: add missing autofunction directives for 9 API functions by @Bodlux in #2036
  • [AMD ROCM] matrix_instr_nonkdim restricted to 16 by @umechand-amd in #2032
  • [Autotuner] Move benchmarking implementation into BenchmarkProvider by @hinriksnaer in #2029
  • [Bug][Autotuner] Autotuner failed to clone mutated input arg by @xiaohongchen1991 in #2042
  • [TPU][Pallas] Fix OOM on large reductions by delegating chunking to Mosaic by @yarongmu-google in #2033
  • Remove ref baseline kernel count for test_clone_with_multiple_views_one_mutated as it depends on PyTorch version by @choijon5 in #2037
  • docs: remove remaining mentions of HELION_USE_AUTOTUNE by @cota in #2038
  • Fix test broken by #2033: use env.backend_name by @norx1991 in #2046
  • Simplify CI: update PyTorch releases to 2.11, use bundled Triton by @choijon5 in #2047
  • [Pallas] Under interpret mode, Use float16 as HALF_DTYPE because bfloat16 is not supported on CPU by @AmesingFlank in #2050
  • [AMD ROCm] Use AMD Triton backend for min_dot_size instead of NVIDIA by @choijon5 in #2048
  • [Pallas] Adjust block size constraints by analyzing subscript exprs on tensors, unblock rms_nor...
Read more

v1.0.0

Choose a tag to compare

@oulgen oulgen released this 03 Apr 16:57
a6ba80b

What's Changed

  • Add PyTorch Conference Europe 2026 events by @choijon5 in #1812
  • Fix LFBO surrogate training data duplication by @fulvius31 in #1806
  • Move imports to top of file by @jansel in #1804
  • [cutedsl] Generalize shape-chain lowering and fix nested matmul argreduce by @jansel in #1787
  • Fix docs build failure by @jansel in #1814
  • Capping the number of warps to avoid non-performing configs and reducing the autotuning time. by @umechand-amd in #1674
  • Add pytorch-probot config to enable retryBot for Helion by @huydhn in #1808
  • Update torch_tpu pin by @oulgen in #1816
  • avoid nvshmem symm-mem backend by @shunting314 in #1750
  • increase signal pad size for dist matmul kernels by @shunting314 in #1753
  • move distributed runtime utils out of examples/ folder by @shunting314 in #1771
  • Increasing block size dimensions to avoid configs which are slow and poor candidates. by @umechand-amd in #1677
  • Cleanup the tests skips on ROCM. by @umechand-amd in #1820
  • Fix Pallas buffer donation: only donate in-place mutated tensors by @thcmbs in #1802
  • Gracefully handle the error, "failed to translate module to LLVM IR" by @umechand-amd in #1793
  • fix to prevent incorrect output in test_print by @umechand-amd in #1821
  • HELION_FORCE_AUTOTUNE: skip cache read but write back result by @fulvius31 in #1815
  • Support block_ptr/TensorDescriptor with extra_mask for loads by @hinriksnaer in #1768
  • [autotuner] Add benchmark_batch as unified entry point for autotuner benchmarking by @hinriksnaer in #1810
  • restrict to persistent pid for dist kernels by @shunting314 in #1772
  • remove hl.signal/wait by @shunting314 in #1791
  • cleanup the autotuner factory function by @shunting314 in #1770
  • properly codegen hl.triton_kernel by @shunting314 in #1797
  • move compile time dist APIs to _dist_utils.py by @shunting314 in #1799
  • Use a larger timeout setting to autotune distributed kernels by @shunting314 in #1800
  • [autotuner] removed unused duplicate function check_config_consistancy by @hinriksnaer in #1829
  • chore: Bump actions/deploy-pages from 4 to 5 by @dependabot[bot] in #1826
  • [pallas-tpu] Add tests for previously untested examples by @v0i0 in #1665
  • [cutedsl] Enable more tests by @jansel in #1822
  • Unskipping tests on ROCM by @umechand-amd in #1831
  • [Helion + torch.compile] Fix dynamic shapes in HOP path and enable test_dynamic_shapes_basic by @yf225 in #1836
  • fix regress caused by large block size in dist autotuner by @shunting314 in #1834
  • Revert "Add scheduled workflow to rerun GPU health check failures (#1683)" by @huydhn in #1833
  • rendezvous only once for matmul reduce scatter by @shunting314 in #1824
  • Refactor: use backend *_expr methods in PointerIndexingStrategy by @aditvenk in #1850
  • Fix #1081: Skip failed configs in parallel benchmark by @tianrengao in #1673
  • Fix pre-commit pyrefly on macOS by sharing ignore-missing-imports logic by @aditvenk in #1858
  • Revert "Capping the number of warps to avoid non-performing configs and reducing the autotuning time. (#1674)" by @jansel in #1860
  • [autotuning] Refactor PrecompileFuture construction out of BaseSearch by @hinriksnaer in #1817
  • [cutedsl] Fix launch axis mapping and serialized reductions by @jansel in #1840
  • [Helion + torch.compile] Always use HOP path, gate fusion with torch_compile_fusion kernel setting by @yf225 in #1837
  • [cutedsl] Fix CuTe shape/view lowering and loop-carried matmul accumulation by @jansel in #1856
  • [Pallas] Fix block_spec_info grid/loop_order mismatch by @norx1991 in #1788
  • add script to make tpu wheels by @v0i0 in #1818
  • chore: Bump actions/configure-pages from 5 to 6 by @dependabot[bot] in #1882
  • Fix Pallas BMM block size alignment via new dim_matches method by @thcmbs in #1828
  • Remove FROM_DEFAULT, unify initial population on FROM_BEST_AVAILABLE with pad control by @fulvius31 in #1809
  • [Helion + torch.compile] Reject static_shapes=True with dynamic=True to prevent wrong results by @yf225 in #1864
  • [Helion + torch.compile] Add bench_compile_config protocol for autotuning by @yf225 in #1865
  • Add pallas_attention kernel and default-loop-type test by @v0i0 in #1866
  • Skip Triton-specific batch dim constraint for Pallas matmul by @v0i0 in #1867
  • [autotune] Migrate precompile future to a dedicated file location by @hinriksnaer in #1872
  • [Pallas] Enable test_use_block_size_var_without_hl_tile Pallas test which already passes, removing triton-specific code checks by @AmesingFlank in #1880
  • torch_tpu: point to new repo location by @cota in #1886
  • [pallas-tpu] treat sublane constraint like lane constraint by @v0i0 in #1887
  • Add epilogue subtiling pass by @choijon5 in #1838
  • [Pallas] Enable test_batch_softmax for Pallas backend by @norx1991 in #1889
  • fix ref eager subtile tests by @v0i0 in #1899
  • use post unfiication symbols during indexing by @v0i0 in #1896
  • Fix _output_indices to recognize views of inputs by @v0i0 in #1868
  • Add emit_pipeline loop-carried state for Pallas attention by @v0i0 in #1869
  • Add fori_loop loop-carried state and refactor shared helpers by @v0i0 in #1870
  • [metal] Add MetalBackend method overrides for codegen by @aditvenk in #1852
  • add CodegenDict to better handle common codegen function by @shunting314 in #1902
  • [Helion + torch.compile] Add extra_cache_key() and is_cacheable() to _AutotunableKernel protocol by @yf225 in #1903
  • Skip overridden config indices during autotuner neighbor generation by @fulvius31 in #1884
  • [Pallas] Fix if branches which modify nonlocal variables, unblocking test_grpo_loss_fwd by @AmesingFlank in #1892
  • [Pallas] Add HelionPallasPrinter to avoid triton_helpers in codegen by @norx1991 in #1900
  • [pallas] Add interpret=True support for CPU reference mode by @v0i0 in #1841
  • [Helion + torch.compile] Move extra_cache_key to its own field in LooseAutotuneCacheKey by @yf225 in #1906
  • [Helion + torch.compile] Add fusion-aware autotuning tests (temporarily skipped) by @yf225 in #1904
  • [Autotuner] Handle CUDA OOM errors gracefully during autotuning by @yf225 in #1912
  • [Helion + torch.compile] Switch to fusion-aware autotuning by @yf225 in #1801
  • Cap maxnreg to fit register file budget for num_warps by @v0i0 in #1918
  • use the user provide process group for autotuning by @shunting314 in #1823
  • disable the flex attention reference in the attention example on tpus by @v0i0 in #1929
  • Refactor RNG op handling to be backend-independent by @jansel in #1871
  • [cutedsl] Fix tile-origin indexing and atomic/matmul correctness by @jansel in #1873
  • Speed up tests by @jansel in #1897
  • add max_mismatch_pct argument to run_example by @shunting314 in #1909
  • scaled matmul reduce scatter by @shunting314 in #1842
  • Use warps_to_threads() instead of hardcoded warp size in maxnreg cap by @v0i0 in #1930
  • [Helion + torch.co...
Read more

v0.3.3

Choose a tag to compare

@oulgen oulgen released this 24 Mar 20:39
d4003c1

What's Changed

  • bump cublas for b200 benchmarks by @v0i0 in #1666
  • Remove need to pass device=... into torch functions by @hinriksnaer in #1657
  • [docs] Updated docs to communicate fast math support by @hinriksnaer in #1669
  • Fix #762: Replace internal assertion with user-facing error in set_pid by @tianrengao in #1670
  • [Helion + torch.compile] Add expected_num_kernels validation to torch.compile tests by @yf225 in #1655
  • [pallas-tpu] Fix segfault xfails: softmax tests now pass, reclassify others by @v0i0 in #1721
  • Reject data-dependent output shapes in infer_output_spec by @gmagogsfm in #1722
  • [pallas-tpu] Add test_long_sum_manual to verify range bound fix by @v0i0 in #1732
  • [pallas-tpu] fix default configs for TPU examples by @v0i0 in #1731
  • Move Triton-specific implementations from Backend base class to TritonBackend. by @norx1991 in #1728
  • Fix lint in ref_mode.py by @jansel in #1681
  • Add missing rebenchmark, finishing phase, and effort profile wiring to DESurrogateHybrid by @fulvius31 in #1680
  • Faster expected generation in test_indexing.py by @jansel in #1682
  • [cutedsl] Implement layout planning phase by @jansel in #1664
  • Revert "[pallas-tpu] fix default configs for TPU examples (#1731)" by @norx1991 in #1740
  • [pallas-tpu] Fix Pallas test_add by making non-contiguous inputs contiguous in pallas launcher. by @norx1991 in #1737
  • [Helion + torch.compile] Refactor HelionTemplateBuffer to use TemplateBuffer base class by @yf225 in #1723
  • precompile in the current process by @shunting314 in #1730
  • [CI] Fix Pyrefly lint error in template_buffer.py by @yf225 in #1746
  • [Autotuner] Add autotune_baseline_accuracy_check_fn for custom accuracy checks by @yf225 in #1733
  • Add Dockerfile by @jansel in #1748
  • Add scripts/runpod.py by @jansel in #1749
  • [cutedls] Initial mma support by @jansel in #1742
  • Fix doubled test output in non-distributed CI jobs by @norx1991 in #1741
  • Add hl.jagged_tile by @nullplay in #1651
  • Fix logging to be compatible with pytest by @bringlein in #1734
  • Add autotune_initial_population_strategy kernel setting by @bringlein in #1735
  • [CI] Increase atol for test_squeeze_and_excitation_net_fwd on B200 by @yf225 in #1752
  • Unpin H100 nightly torch and Triton versions by @v0i0 in #1654
  • Enable pyrefly on macOS with ignore-missing-imports by @aditvenk in #1760
  • [metal] Register "metal" backend with minimal MetalBackend and launcher by @aditvenk in #1761
  • [metal] Respect force_tile_mask() in NDTileStrategy mask generation by @aditvenk in #1762
  • Remove past hackathon event from README by @choijon5 in #1780
  • Remove deadcode _clone_tree and _assert_args_close by @choijon5 in #1781
  • Fix PermutationFragment.encode() returning wrong value by @choijon5 in #1779
  • Fix atomic_max ref to return previous value by @choijon5 in #1782
  • Fix lints by @jansel in #1775
  • Add runpod SKILL.md by @jansel in #1776
  • [Helion + torch.compile] Add store/load transform hooks and prologue/epilogue fusion codegen by @yf225 in #1724
  • [Helion + torch.compile] Enable torch.compile fusion tests by @yf225 in #1727
  • [Helion + torch.compile] Simplify _remap_or_resolve for compound sympy expressions by @yf225 in #1785
  • Add scheduled workflow to rerun GPU health check failures by @v0i0 in #1683
  • Fix device sync in generic benchmarking functions for TPU/Pallas by @norx1991 in #1773
  • Removing skips and in some cases adding skipIfNotCUDA for cuda only features. by @umechand-amd in #1790
  • APIs to debug distributed kernel by @shunting314 in #1743
  • [cutedsl] tcgen05 MMA support by @jansel in #1777
  • [cutedsl] Fix broadcast and reshape-backed matmul lowering by @jansel in #1783
  • [cutedsl] Fix packed-RHS lowering and add general stack/reshape views by @jansel in #1784
  • Fix ROCm failures on main by @jansel in #1805
  • add kernel-filter to select kernel for allreduce-rmsnorm by @shunting314 in #1744
  • helion distributed kernel autotuning by @shunting314 in #1532
  • [Helion + torch.compile] Add ref baseline kernel count checks by @yf225 in #1786

New Contributors

Full Changelog: v0.3.2...v0.3.3

v0.3.2

Choose a tag to compare

@oulgen oulgen released this 13 Mar 21:33
59927cc

What's Changed

  • [pallas-tpu] Enable test_geglu, test_swiglu on Pallas; fix softmax/kl_div/gdn xfails by @v0i0 in #1637
  • [pallas-tpu] make dot work: use jnp.matmul and f32 accumulator by @v0i0 in #1647
  • [pallas-tpu] Enable broadcasting and control flow tests for Pallas backend by @v0i0 in #1643
  • Use ephemeral Triton cache during autotuning to eliminate loser artifacts by @fulvius31 in #1649
  • Try public bazel cache for torch_tpu build by @v0i0 in #1658
  • fix cudagraph grouped gemm fail in benchmarks by @v0i0 in #1662
  • Use matmul fwd direclty in autograd for performance by @tianrengao in #1045
  • Recreate foreign SymInts in helion's own ShapeEnv by @gmagogsfm in #1671

Full Changelog: v0.3.1...v0.3.2

v0.3.1

Choose a tag to compare

@oulgen oulgen released this 12 Mar 21:23
4c4b33b

What's Changed

  • [Autotuner] Use chunked comparison in autotuner accuracy checks to reduce peak memory by @yf225 in #1538
  • one_shot_allreduce_bias_rmsnorm: alloc&rendezvous symm memory only once by @shunting314 in #1525
  • Add Horace as Core Maintainer by @oulgen in #1552
  • add two-shot all redcue rms norm kernel by @shunting314 in #1526
  • Update bazel build command to use remote caching by @malfet in #1550
  • Add BlockSpec support and proper indexing to Pallas launcher by @oulgen in #1548
  • Add Pallas codegen for prims.iota by @oulgen in #1556
  • [CI] Fix AMD error on test_default_block_sizes_high_dim_with_reduction by @yf225 in #1546
  • [Autotuner] Eliminate self._original_args clone to reduce peak memory by @yf225 in #1547
  • sync bench/rebench result across ranks by @shunting314 in #1542
  • Skip failing test on MTIA by @Myrthan in #1558
  • [Cache] Defer some initialization in autotuner to skip unnecessary work on cache hits by @fulvius31 in #1557
  • Enable Pallas backend for test_control_flow with lax.cond by @oulgen in #1554
  • Add Pallas atomic ops -- everything is atomic by default by @oulgen in #1560
  • Add Pallas emit_pipeline by @oulgen in #1561
  • Move do_bench/sync_object from _testing to autotuner/benchmarking to fix docs build by @oulgen in #1562
  • Improve launcher overhead by @oulgen in #1563
  • fix mock in debug utils - hopefully make ci less flaky by @v0i0 in #1568
  • fix fp8 autotuner fail for ci by @v0i0 in #1569
  • fix some indexing by @v0i0 in #1453
  • chore: Bump actions/upload-artifact from 6 to 7 by @dependabot[bot] in #1579
  • chore: Bump actions/download-artifact from 7 to 8 by @dependabot[bot] in #1580
  • Build TorchTPU in opt mode by @oulgen in #1573
  • Add Pallas fori_loop with async DMA and pallas_loop_type config by @oulgen in #1574
  • Enable waves_per_eu tunable on RDNA GPUs by @fulvius31 in #1571
  • [cutedsl] Cross-warp reductions by @jansel in #1543
  • sync seeds across ranks by @shunting314 in #1555
  • [cutedsl] Add more reduction types by @jansel in #1565
  • [cutedsl] Tuple reductions by @jansel in #1566
  • Fix CPU tests by @jansel in #1582
  • Add Pallas reduction support and enable test_reductions for TPU by @oulgen in #1549
  • Tell AGENTS to not run git push by @oulgen in #1592
  • partition aot tuning by config validity by @v0i0 in #1581
  • bump pytorch to 2.10 in benchmark nightly to unblock grouped_gemm tritonbench by @v0i0 in #1597
  • fix benchmark parsing by @v0i0 in #1596
  • make sure persistent setup code uses index type for total pids by @v0i0 in #1595
  • Fix issue with branching and static_range by @jansel in #1586
  • Update AGENTS.md assertExpectedJournal reference by @jansel in #1587
  • [cutedsl] Enable more test files by @jansel in #1583
  • [cutedsl] Fix issue with sympy printer by @jansel in #1585
  • [Cache] Fix cross-backend cache poisoning by adding backend to cache key by @fulvius31 in #1593
  • [Helion + torch.compile] Add additional inductor fusion tests by @yf225 in #1594
  • [Autotuner] Add FROM_BEST_AVAILABLE initial population strategy by @fulvius31 in #1365
  • Unwrap single-element list in hl.tile() to match scalar behavior by @blake-snc in #1559
  • Enable test_examples by @oulgen in #1598
  • Generalize autotuning infrastructure to support Pallas/TPU backend by @oulgen in #1591
  • Fix symbolic variable specialization when indexing tensor in host block by @yf225 in #1575
  • Add CLAUDE.md as symlink to AGENTS.md by @oulgen in #1610
  • update grouped gemm tb signature by @v0i0 in #1608
  • Update logo and add events for hackathon and PLDI tutorial by @choijon5 in #1611
  • Add claude skill for TPU development by @oulgen in #1612
  • Enable test random and rnd by @oulgen in #1607
  • [Helion + torch.compile] Temporarily skip torch.compile fusion test cases to wait for cross-repo changes by @yf225 in #1613
  • add ci health check (hopefully catches cudaDeviceUnavailable) by @v0i0 in #1614
  • [Helion + torch.compile] Skip test_symint_return_from_tensor_shape temporarily by @yf225 in #1615
  • update rocm in benchmark now that we bumped pytorch by @v0i0 in #1609
  • Remove Triton CPU backend and all CPU references by @oulgen in #1616
  • [Internal CI] Use MTIA-aligned tensor shape in test_tile_single_element_list by @yf225 in #1618
  • Refactor device_ir passes by @jansel in #1606
  • Add @skipIfCudaSharedMemoryLessThan for failing RTX5090 tests by @jansel in #1617
  • another benchmark parsing fix by @v0i0 in #1621
  • Disable forking to fix failing test by @jansel in #1619
  • Support for Advanced Control Files in Autotuner and Configs by @ptorru in #1576
  • Fix scalar tensor indexing crash by @hinriksnaer in #1620
  • [Docs] Update autotuner and TileIR backend documentation by @fulvius31 in #1600
  • benchmarking all_gather_matmul by @shunting314 in #1605
  • Pin H100 nightly and skip flakey pallas test by @jansel in #1624
  • [CI] Relax tolerances in test_hl_arange_non_power_of_2 by @yf225 in #1627
  • Decrease sizes for test_batch_softmax_block_ptr by @jansel in #1628
  • fix dynamic shapes handling tensor descriptors conservatively by @v0i0 in #1604
  • [Docs] Fix Helion repo link by @choijon5 in #1629
  • Rename Helion Puzzles to Tutorials; fix broken tutorial examples; move to Markdown by @yf225 in #1631
  • fix size hint lint from nightly by @v0i0 in #1630
  • Introduce MTIA autotuning knobs by @kile01 in #1572
  • change to gfx942.1 runner like pytorch, hopefully shorter queue by @v0i0 in #1632
  • Set TRITON_STORE_BINARY_ONLY=1 during autotuning to reduce cache size by @fulvius31 in #1590
  • Enable fast triton sigmoid by @hinriksnaer in #1564
  • Fix fake impl inference with unbacked SymInts by @gmagogsfm in #1626
  • [pallas-tpu] Add HALF_DTYPE constant for backend-portable half-precision tests by @v0i0 in #1636
  • add gpu health check to tests as well by @v0i0 in #1634
  • Properly generate floor division using // for constexpr by @PaulZhang12 in #1625
  • Add --file-header option for custom headers in AOT generated files by @v0i0 in #1645
  • [pallas-tpu] Bump libtpu to 0.0.37, jax/jaxlib to 0.9.1 by @v0i0 in #1641
  • [pallas-tpu] Enable pallas backend for view/reshape tests in test_views.py by @v0i0 in #1638
  • fix a10g lerp nightly error by @v0i0 in #1652
  • [pallas-tpu] remove shape based block specs by @v0i0 in #1633
  • Fix lerp decomp in _make_fx for PyTorch nightly by @v0i0 in #1653
  • [pallas-tpu] Fix BlockSpec regressions from codegen BlockSpecs PR by @v0i0 in #1656
  • aot tuning standalone output by @v0i0 in #1650
  • [pallas-tpu] ...
Read more

v0.3.0

Choose a tag to compare

@oulgen oulgen released this 24 Feb 18:49
30ebc91

What's Changed

  • [xpu][test]Skip Failed xpu tests by @Stonepia in #1337
  • [Autotuner] Introduce Triton-TileIR Backend Support by @qelk123 in #1250
  • test: removing skip condition for ROCm since tests are passing locally. by @umechand-amd in #1340
  • Removing skipIfRocm property from the tests which are passing. by @umechand-amd in #1341
  • Fix softmax puzzle by @nicolo-domyn in #1344
  • Diversity-aware batch selection for LFBO Pattern Search by @ethche in #1350
  • [xpu][fix] fix errors related to maxnreg on xpu by @jianyizh in #1347
  • AOT autotuning runner and cache and heuristics gener… by @v0i0 in #1278
  • handle ToFloat for symints / from non-static shapes by @v0i0 in #1357
  • [Benchmark CI] Set log level to INFO to reduce log volume by @yf225 in #1360
  • Support List Type kernel cache by @mengluy0125 in #1359
  • Improve autotuning log readability by @hinriksnaer in #1362
  • Skip tensor_descriptor and eviction policy for AMD by @mengluy0125 in #1349
  • Fix pyrefly type error in surrogate pattern search by @oulgen in #1364
  • Add test for torch.nn.Parameter by @jansel in #1371
  • Skip tests on lower memory cards by @jansel in #1370
  • [TileIR][Doc] add doc for CUDA Tile IR backend by @qelk123 in #1352
  • [Docs] Put API Reference on the docs earlier by @oulgen in #1374
  • Increase Top Nav options to 7 by @sekyondaMeta in #1375
  • v0i0/fix is none by @v0i0 in #1369
  • disable tma with int64 by @v0i0 in #1372
  • fix indexer expansion by @v0i0 in #1367
  • [Helion-Inductor Fusion] Add extensive unit tests for mutation/aliasing edge cases by @yf225 in #1358
  • [Helion + torch.compile] Fix Dynamo error when same Helion kernel is called twice in same compile region by @yf225 in #1378
  • improve signal handling for the aot runner by @v0i0 in #1366
  • Use unittest.skip instead of pytest.mark.xfail by @yf225 in #1380
  • [CI] Speedup tests via parallelism by @oulgen in #1338
  • Support List indexing with non-literal indices by @mengluy0125 in #1377
  • Throw error if TRITON_INTERPRET=1 and HELION_INTERPRET=1 are used together by @parsshar-RH in #1379
  • [autotuner] filter and clone only mutated tensors by @bringlein in #1252
  • Update pyrefly to 0.51.1 by @oulgen in #1386
  • Update ruff to 0.15.0 by @oulgen in #1388
  • [CI] Increase atol for test_hl_arange_non_power_of_2 by @yf225 in #1390
  • Set an adaptive compile timeout for autotuning by @ethche in #1384
  • [Helion + torch.compile] Add allow_torch_compile_fusion flag to unit tests (no-op for now) by @yf225 in #1382
  • [Helion + torch.compile] Add Helion kernel Dynamo variable and HOP by @yf225 in #1383
  • AOT oneshot workflow fix by @tyler-romero in #1381
  • [CI] Fix NCCL error in distributed CI by @yf225 in #1396
  • Fix zero-size dim handling due to upstream PyTorch changes by @yf225 in #1398
  • Fix hl.rand/hl.randint crash when dim is specialized to size 1 by @karthickai in #1401
  • Update num_warps and num_stages range for AMD by @yvonne-lab in #1368
  • [CI] Fix B200 flaky test errors by deferring CUDA initialization to test execution time for pytest-xdist compatibility by @yf225 in #1399
  • Add generic kernel config adapter by @drisspg in #1385
  • Disable autocast during kernel compilation to fix mixed-precision training by @yf225 in #1402
  • Add nvfp4 gemm by @oulgen in #1403
  • Fix unpacking of torch.sort/topk return types in Helion kernel by @yf225 in #1406
  • Skip O(N^2) argsort in torch.sort lowering when indices is unused by @yf225 in #1408
  • [CI] Adjust atol for squeeze_and_excitation_net_bwd unit tests by @yf225 in #1409
  • [CI] Adjust atol for test_dtype_cast_preserved_before_second_dot test by @yf225 in #1411
  • Fix torch.cumsum overwriting its input when used in Helion kernel by @yf225 in #1410
  • Skip hl.rand/hl.randint specialize test on MTIA by @karthickai in #1414
  • Skip "too many blocks in cooperative launch" error during autotuning by @yf225 in #1417
  • [CI] Fix flaky OOM error for test_int32_offset_out_of_range_error on H100 by @yf225 in #1419
  • [CI] Adjust atol for test_hl_arange_non_power_of_2 test by @yf225 in #1421
  • [Helion + torch.compile] Add multi-output and scalar return support by @yf225 in #1395
  • Improve DeviceIR.get_root() to handle multiple reduction dimensions by @yf225 in #1434
  • Introduce Backend abstraction for multi-backend code generation by @oulgen in #1426
  • Route codegen decorator, constexpr, and imports through Backend by @oulgen in #1427
  • Route dtype string conversions through Backend abstraction by @oulgen in #1428
  • Rename triton_index_type to backend-agnostic index_type by @oulgen in #1429
  • dump triton failures by @v0i0 in #1407
  • Fix PointwiseLowering rank mismatch for size-1 reductions by @yf225 in #1424
  • [Helion + torch.compile] Fix aliasing, mutation detection, and nested return support by @yf225 in #1432
  • for now rewrite gdn fwd h without hl.grid for fusion by @v0i0 in #1391
  • Add PallasBackend class and make backend configurable via Settings by @oulgen in #1447
  • sync compute_repeat output for distributed kernels by @shunting314 in #1435
  • raise tolerance for test_hl_dot_out_dtype_argument by @shunting314 in #1457
  • Skip too many resource error in autotuner by @oulgen in #1458
  • Preserve symbolic shapes in HOP output_spec during tracing by @gmagogsfm in #1454
  • Refactor "start_precompile_and_check_for_hangs" by @hinriksnaer in #1422
  • Add finishing phase to autotuner for config minimization by @v0i0 in #1319
  • [Helion + torch.compile] Add unit test for returning SymInt from Helion kernel by @yf225 in #1459
  • Workaround var_to_val deprecated warning by @jansel in #1461
  • Make triton an optional dependency so Pallas backend works without triton installed by @oulgen in #1452
  • Add CuteBackend boilerplate by @jansel in #1462
  • Allow 'common' as fallback backend implementation by @jansel in #1463
  • Refactor backend hooks for launch, program IDs, and autotune by @jansel in #1464
  • Add Pallas codegen for tile, load, store and runtime launcher by @oulgen in #1448
  • Add Pallas ATen lowerings, backend-aware tile strategy, and creation ops by @oulgen in #1449
  • Add Pallas backend tests for add kernel by @oulgen in #1450
  • Add Pallas backend to CI test matrix by @oulgen in #1451
  • Add Cute pointwise add codegen, launcher, and tests by @jansel in #1465
  • Wire up inductor op_overrides for the cute backend by @jansel in #1467
  • Update pyproject with cute-cu12/cute-cu13 by @jansel in #1468
  • Add tests for cute scalar and kwargs by @jansel in #1469
  • Ignore some linter warnings by @jansel in #1470
  • Skip most tests when HELION_BACKEND!=triton and add .expected_{cute,pallas} files by @jansel in #1471
  • [Cache] Add bac...
Read more

v0.2.10

Choose a tag to compare

@oulgen oulgen released this 13 Jan 19:45
bd64602

What's Changed

Full Changelog: v0.2.9...v0.2.10

v0.2.9

Choose a tag to compare

@oulgen oulgen released this 31 Dec 00:57
2a2d2b2

What's Changed

Full Changelog: v0.2.8...v0.2.9

v0.2.8

Choose a tag to compare

@oulgen oulgen released this 23 Dec 20:10
768326f

What's Changed

  • Support passing Triton function object to hl.triton_kernel() by @yf225 in #1263
  • chore: Bump actions/download-artifact from 6 to 7 by @dependabot[bot] in #1270
  • chore: Bump actions/upload-artifact from 5 to 6 by @dependabot[bot] in #1271
  • [Distributed] one_shot_allreduce_bias_rmsnorm example by @yf225 in #1266
  • [Distributed] matmul_reduce_scatter example by @yf225 in #1269
  • feat(benchmarks): add shapes to json output by @fulvius31 in #1273
  • [Autotuner] Log the 'started' state to CSV, for easier user monitoring of kernel hanging at runtime by @yf225 in #1279
  • default pattern search by @v0i0 in #1259
  • Set LFBOPatternSearch as default by @ethche in #1280
  • fix surrogate search for singleton population by @v0i0 in #1281
  • Ignore bzl files in git. by @Myrthan in #1282
  • chunk fused_linear_jsd by @v0i0 in #1277
  • Fix buggy interation between XYZProgramIDs and L2GroupingProgramIDs by @jansel in #1288
  • Fix bug with torch.rand_like compile error by @jansel in #1289
  • [autotuner] print path to generated Triton code after selection of kernel by @bringlein in #1285
  • Add support for torch.gather by @jansel in #1290
  • [docs] Add more example to docs by @oulgen in #1301
  • [lint] remove dead ignores by @oulgen in #1302
  • Add proper error handling for torch.split and torch.tensor_split in device loops by @oulgen in #1297
  • Fix BlockReductionStrategy to use existing index variables for argmax/argmin operations by @oulgen in #1298
  • Skip more failing tests on cpu backend by @oulgen in #1304
  • [CI] Fix broken notebook by @oulgen in #1305
  • Fix shape inference for tile indexing on size-1 dimensions and use broadcast_to for block_ptr by @oulgen in #1299
  • Fix codegen broadcasting for tile indexing on size-1 tensor dimensions by @oulgen in #1300
  • Enable tests on py314 by @oulgen in #1306

New Contributors

Full Changelog: v0.2.7...v0.2.8