tilelang.language.proxy¶

Buffer/Tensor proxy in TileLang.

Attributes¶

Buffer

類別¶

`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.
`StridedTensorProxy`	Main tensor proxy class for global scope buffers, with strides supported.
`FragmentBufferProxy`	Proxy class for fragment memory buffers.
`SharedBufferProxy`	Proxy class for shared memory buffers.
`LocalBufferProxy`	Proxy class for local memory buffers.
`BaseTensor`

函式¶

`ptr`([dtype, storage_scope, is_size_var])	Create a TIR var that represents a pointer.
`make_tensor`(ptr, shape[, dtype, strides])

Module Contents¶

class tilelang.language.proxy.BufferProxy¶

Buffer proxy class for constructing tir buffer.

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

參數:

shape (tilelang._typing.ShapeType)
dtype (tilelang._typing.DType)

回傳型別:

tvm.tir.Buffer

__getitem__(keys)¶

回傳型別:: tvm.tir.Buffer

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

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

參數:

pointer_var (tvm.tir.Var) -- The pointer variable
shape (tilelang._typing.ShapeType) -- The shape of the buffer
dtype (tilelang._typing.DType) -- The data type of the buffer (default: float32)
strides (tuple[tvm.tir.PrimExpr, Ellipsis] | None)

回傳:

A buffer created from the given parameters

回傳型別:

tvm.tir.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)¶

參數:

shape (tilelang._typing.ShapeType)
dtype (tilelang._typing.DType)

回傳型別:

tvm.tir.Buffer

__getitem__(keys)¶

回傳型別:: tvm.tir.Buffer

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

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

參數:

pointer_var (tvm.tir.Var) -- The pointer variable
shape (tilelang._typing.ShapeType) -- The shape of the buffer
dtype (tilelang._typing.DType) -- The data type of the buffer (default: float32)
strides (tuple[tvm.tir.PrimExpr, Ellipsis] | None)

回傳:

A buffer created from the given parameters

回傳型別:

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, the tensor should be by default contiguous.

__call__(shape, dtype='float32', data=None, scope=None)¶

參數:

shape (tilelang._typing.ShapeType | tvm.tir.PrimExpr | int)
dtype (tilelang._typing.DType)

回傳型別:

tvm.tir.Buffer

class tilelang.language.proxy.StridedTensorProxy¶

Bases: BaseTensorProxy

Main tensor proxy class for global scope buffers, with strides supported.

This class implements the default tensor proxy with global memory scope, with the stride information required.

__call__(shape, strides, dtype='float32', scope=None)¶

參數:

shape (tilelang._typing.ShapeType)
strides (tuple[Any])
dtype (tilelang._typing.DType)

回傳型別:

tvm.tir.Buffer

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)¶

參數:

shape (tilelang._typing.ShapeType)
dtype (tilelang._typing.DType)

classmethod __class_getitem__(key)¶

__getitem__(key)¶

回傳型別:: Any

__setitem__(key, value)¶

回傳型別:: None

classmethod from_ptr(pointer_var, shape, dtype='float32', strides=None)¶

參數:

pointer_var (tvm.tir.Var)
shape (tilelang._typing.ShapeType)
dtype (tilelang._typing.DType)
strides (tuple[tvm.tir.PrimExpr, Ellipsis] | None)

回傳型別:

typing_extensions.Self

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

Create a TIR var that represents a pointer.

參數:

dtype (DType) -- 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.

回傳:

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

回傳型別:

PrimExpr

tilelang.language.proxy.make_tensor(ptr, shape, dtype='float32', strides=None)¶

參數:

ptr (tvm.tir.Var)
shape (tilelang._typing.ShapeType)
dtype (tilelang._typing.DType)
strides (tuple[tvm.tir.PrimExpr, Ellipsis] | None)

回傳型別:

tvm.tir.Buffer