Enceladus v0.1.0 · alpha
Apple siliconMetalPython 3.11+

Triton-style GPU kernels, compiled for your Mac.

Enceladus compiles Python tile programs to Metal Shading Language while your program runs. You don't need Xcode or the Metal toolchain. Kernels accept NumPy arrays, PyTorch mps tensors, and MLX arrays.

uv add enceladus
0.37 msto compile a vector add from Python to MSL
3.4 µshost cost of an asynchronous launch
M1 and latermacOS 15 or later, any Apple GPU
import numpy as np

import enceladus
import enceladus.language as tl


@enceladus.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n
    x = tl.load(x_ptr + offs, mask=mask)
    y = tl.load(y_ptr + offs, mask=mask)
    tl.store(out_ptr + offs, x + y, mask=mask)


x = np.random.default_rng(0).standard_normal(100_000, dtype=np.float32)
y = np.ones_like(x)
out = np.empty_like(x)
add_kernel[(enceladus.cdiv(x.size, 1024),)](x, y, out, x.size, BLOCK=1024)
np.testing.assert_allclose(out, x + y)
The generated.metal tab is real compiler output for this kernel. Each of the 128 threads holds 8 of the 1,024 tile elements in registers.
OverviewAll pages
OverviewQuickstartProgramming modelMemory and synchronizationLanguage referenceDebuggingFramework interopPerformanceBenchmarksPorting from Triton

Write the tile, not the thread

You write an Enceladus kernel as a Python function that operates on tiles: blocks of values that one program computes on at once, like small NumPy arrays. The compiler decides which thread holds which element, lowers tl.dot to the GPU's SIMD-group matrix instructions, and hands the resulting MSL to Metal's newLibraryWithSource.

If you know Triton, you already know most of Enceladus. The decorator, the launch syntax, tl.constexpr, masks, and the autotuner work the same way.

no toolchain

Compiles at run time

Metal compiles the generated MSL in your process. Results are cached on disk, so later runs load in under 10 ms.

interop

Bring your arrays

Pass NumPy arrays, PyTorch mps tensors, MLX arrays, or enceladus.Tensor objects directly to a kernel.

ENCELADUS_INTERPRET=1

Debug on the CPU

The NumPy interpreter runs the same kernel with print(), pdb, and bounds-checked loads.

CompilationError

No silent miscompiles

Every unsupported construct is refused with the file, line, column, and a suggested fix.

tl.dot

Two matmul backends

Use simdgroup_matrix on every Apple GPU, or Metal 4 matmul2d for eligible loops.

@enceladus.autotune

Tuned per GPU

The autotuner compiles candidates in parallel, times them with GPU timestamps, and saves the winner.

Explore the docs

Get startedQuickstartInstall Enceladus and run a vector add, a softmax, a matmul, and flash attention. ConceptsProgramming modelPrograms, the grid, tiles and their layout across SIMD groups, and specialization. ConceptsMemory and synchronizationMasks, pointer tiles, tensor descriptors, streams, and when results are visible. ReferenceLanguage referenceEvery tl builtin and host function, the data types, and the environment variables. GuidesDebuggingCompilation errors, the CPU interpreter, MSL dumps, kernel.explain, and GPU capture. GuidesFramework interopPass PyTorch, MLX, and NumPy arrays to kernels, and share memory through DLPack. GuidesPerformanceBlock sizes, tl.dot backends, autotuning, and asynchronous launches. GuidesBenchmarksThroughput against MLX on an M4 Pro, and launch overhead. GuidesPorting from TritonThe steps to port a Triton kernel, what differs, and what isn't supported.

Project status

Enceladus is alpha software. Keep the following in mind:

  • All published performance numbers come from one M4 Pro. Wheels target macOS 15 but have been tested only on later releases.
  • No wheel is published for Python 3.14, so 3.14 builds from source.
  • Without ENCELADUS_VERIFY=1, the interpreter accepts some constructs that the compiler refuses, such as while and a tl.dot whose K isn't a multiple of 8.
  • MLX launches are always synchronous, at about 100 µs each.
  • atomic_add doesn't support float16 or bfloat16.

To report a problem, open an issue on GitHub.