tilelang.contrib.cutedsl.mbar

Simple wrappers that delegate to cutlass.cute.arch implementations. We use the existing implementations from cutlass rather than reinventing the wheel.

Functions

mbarrier_wait(mbar_ptr, phase[, timeout_ns, loc, ip])	Waits on a mbarrier with a specified phase.
mbarrier_cp_async_arrive(mbar_ptr, *[, loc, ip])
fence_proxy_async()
fence_barrier_init()

Module Contents

tilelang.contrib.cutedsl.mbar.mbarrier_wait(mbar_ptr, phase, timeout_ns=10000000, *, loc=None, ip=None)

Waits on a mbarrier with a specified phase.

Parameters:

• mbar_ptr (cutlass.cute.typing.Pointer)

• phase (cutlass.cute.typing.Int)

• timeout_ns (cutlass.cute.typing.Int)

Return type:

None

tilelang.contrib.cutedsl.mbar.mbarrier_cp_async_arrive(mbar_ptr, *, loc=None, ip=None)

Parameters:

mbar_ptr (cutlass.cute.typing.Pointer)

Return type:

None

tilelang.contrib.cutedsl.mbar.fence_proxy_async()

tilelang.contrib.cutedsl.mbar.fence_barrier_init()