tilelang.cuda.language.kernel¶

CUDA dialect of T.Kernel: the common launch plus CUDA launch annotations.

Functions¶

Kernel(*blocks[, threads, prelude, cluster_dims])

Construct a kernel launch frame for CUDA: a grid of thread blocks.

Module Contents¶

tilelang.cuda.language.kernel.Kernel(*blocks, threads=None, prelude=None, cluster_dims=None)¶

Construct a kernel launch frame for CUDA: a grid of thread blocks.

Code inside the launch operates at the block level: T.Parallel, T.copy and friends are mapped onto threads by the compiler. T.get_thread_binding() exposes threadIdx for thread-level code. The keyword arguments are recorded at trace time and materialized by the CUDA pipeline once the target is known.

Parameters:

  • *blocks (int | PrimExpr) – Grid extent along each axis (1-3 dimensions, gridDim.(x|y|z)). The launch yields one block index per axis (blockIdx.(x|y|z)).

  • threads (int | list[int] | tuple[int, ...], optional) – Threads per block: a count for blockDim.x or up to three per-dimension extents for blockDim.(x|y|z). Defaults to 128 when omitted.

  • prelude (str, optional) – CUDA source injected before the generated kernel, e.g. #include lines or helper functions.

  • cluster_dims (int | tuple[int, int, int] | list[int], optional) – Thread block cluster shape (SM90+). 2 or (2, 1, 1) launches 2-CTA clusters via cudaLaunchKernelEx. T.ClusterKernel is the same launch with a required cluster_dims.

Return type:

tilelang.language.kernel.KernelLaunchFrame

Examples

with T.Kernel(T.ceildiv(N, 128), threads=128) as bx:
    ...

with T.Kernel(grid_x, grid_y, threads=(64, 2)) as (bx, by):
    tx, ty = T.get_thread_bindings()
    ...