tilelang.language.proxy¶

The language interface for tl programs.

Attributes¶

Buffer

Classes¶

BufferProxy

Buffer proxy class for constructing tir buffer.

BaseTensorProxy

Base proxy class for tensor types with configurable defaults.

TensorProxy

Main tensor proxy class for global scope buffers.

FragmentBufferProxy

Proxy class for fragment memory buffers.

SharedBufferProxy

Proxy class for shared memory buffers.

LocalBufferProxy

Proxy class for local memory buffers.

BaseTensor

Functions¶

ptr([dtype, storage_scope, is_size_var])

Create a TIR var that represents a pointer.

make_tensor(ptr, shape[, dtype])

Module Contents¶

class tilelang.language.proxy.BufferProxy¶

Buffer proxy class for constructing tir buffer.

__call__(shape, dtype='float32', data=None, strides=None, elem_offset=None, scope='global', align=0, offset_factor=0, buffer_type='', axis_separators=None)¶
Return type:

tvm.tir.Buffer

__getitem__(keys)¶
Return type:

tvm.tir.Buffer

from_ptr(pointer_var, shape, dtype='float32')¶

Create a buffer from a pointer, shape, and data type.

Parameters:

  • pointer_var (tvm.tir.Var) – The pointer variable

  • shape (tuple[tvm.tir.PrimExpr, Ellipsis]) – The shape of the buffer

  • dtype (str) – The data type of the buffer (default: float32)

Returns:

A buffer created from the given parameters

Return type:

Buffer

class tilelang.language.proxy.BaseTensorProxy¶

Base proxy class for tensor types with configurable defaults.

This class serves as a foundation for different tensor proxy types, providing customizable default values for scope, alignment, and offset factors. It implements the core functionality for creating TIR buffers with specific memory configurations.

default_scope = 'global'¶
default_align = 0¶
default_offset_factor = 0¶
__call__(shape, dtype='float32', data=None, strides=None, elem_offset=None, scope=None, align=None, offset_factor=None, buffer_type='', axis_separators=None)¶
Return type:

tvm.tir.Buffer

__getitem__(keys)¶
Return type:

tvm.tir.Buffer

from_ptr(pointer_var, shape, dtype='float32')¶

Create a buffer from a pointer, shape, and data type.

Parameters:

  • pointer_var (tvm.tir.Var) – The pointer variable

  • shape (tuple[tvm.tir.PrimExpr, Ellipsis]) – The shape of the buffer

  • dtype (str) – The data type of the buffer (default: float32)

Returns:

A buffer created from the given parameters

Return type:

tvm.tir.Buffer

class tilelang.language.proxy.TensorProxy¶

Bases: BaseTensorProxy

Main tensor proxy class for global scope buffers.

This class implements the default tensor proxy with global memory scope, inheriting all functionality from BaseTensorProxy without modifications.

class tilelang.language.proxy.FragmentBufferProxy¶

Bases: BaseTensorProxy

Proxy class for fragment memory buffers.

This class represents tensor proxies specifically for local fragment memory, typically used in GPU tensor core operations.

default_scope = 'local.fragment'¶
class tilelang.language.proxy.SharedBufferProxy¶

Bases: BaseTensorProxy

Proxy class for shared memory buffers.

This class represents tensor proxies for dynamic shared memory, commonly used in GPU shared memory operations.

default_scope = 'shared.dyn'¶
class tilelang.language.proxy.LocalBufferProxy¶

Bases: BaseTensorProxy

Proxy class for local memory buffers.

This class represents tensor proxies for local memory scope, typically used for temporary computations in GPU kernels.

default_scope = 'local'¶
tilelang.language.proxy.Buffer¶
class tilelang.language.proxy.BaseTensor(shape, dtype='float32', data=None, strides=None, elem_offset=None, scope=None, align=None, offset_factor=None, buffer_type='', axis_separators=None)¶
Parameters:

shape (Sequence[SupportsIndex])

classmethod __class_getitem__(key)¶
__getitem__(key)¶
Return type:

Any

__setitem__(key, value)¶
Return type:

None

classmethod from_ptr(pointer_var, shape, dtype='float32')¶
Parameters:

  • pointer_var (tvm.tir.Var)

  • shape (Sequence[tvm.tir.PrimExpr, Ellipsis])

  • dtype (str)

Return type:

typing_extensions.Self

tilelang.language.proxy.ptr(dtype=None, storage_scope='global', *, is_size_var=False)¶

Create a TIR var that represents a pointer.

Parameters:

  • dtype (str) – The data type of the pointer.

  • storage_scope (str) – The storage scope of the pointer.

  • is_size_var (bool) – Whether or not to return a SizeVar instead of Var.

Returns:

res – The new tir.Var with type handle or casted expression with type handle.

Return type:

PrimExpr

tilelang.language.proxy.make_tensor(ptr, shape, dtype='float32')¶
Parameters:

  • ptr (tvm.tir.Var)

  • shape (tuple[tvm.tir.PrimExpr, Ellipsis])

  • dtype (str)

Return type:

tvm.tir.Buffer