Mojo struct

TMALoadOp

struct TMALoadOp[a_type: DType, b_type: DType, block_tile_shape: IndexList[3], cluster_shape: IndexList[3], a_swizzle: TensorMapSwizzle = TensorMapSwizzle(3), b_swizzle: TensorMapSwizzle = TensorMapSwizzle(3)]

Fields

• ​a_tma_ptr (UnsafePointer[TMATensorTile[a_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](0) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](0) < 0) ^ (cluster_shape.__getitem__[::Indexer](0) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()]]):
• ​b_tma_ptr (UnsafePointer[TMATensorTile[b_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](1) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](1) < 0) ^ (cluster_shape.__getitem__[::Indexer](1) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()]]):

Implemented traits

AnyType, LoadOp, UnknownDestructibility

Aliases

__del__is_trivial

alias __del__is_trivial = True

a_tma_desc_layout

alias a_tma_desc_layout = _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()

a_tma_layout

alias a_tma_layout = row_major(0 if (cluster_shape.__getitem__[::Indexer](0) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](0) < 0) ^ (cluster_shape.__getitem__[::Indexer](0) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">), block_tile_shape.__getitem__[::Indexer](2))

a_tma_type

alias a_tma_type = TMATensorTile[a_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](0) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](0) < 0) ^ (cluster_shape.__getitem__[::Indexer](0) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()]

args_type

alias args_type = TMALoadOpArgs[a_type, b_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](0) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](0) < 0) ^ (cluster_shape.__getitem__[::Indexer](0) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({0}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), row_major(0 if (cluster_shape.__getitem__[::Indexer](1) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](1) < 0) ^ (cluster_shape.__getitem__[::Indexer](1) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType(), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()]

b_tma_desc_layout

alias b_tma_desc_layout = _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()

b_tma_layout

alias b_tma_layout = row_major(0 if (cluster_shape.__getitem__[::Indexer](1) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](1) < 0) ^ (cluster_shape.__getitem__[::Indexer](1) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">), block_tile_shape.__getitem__[::Indexer](2))

b_tma_type

alias b_tma_type = TMATensorTile[b_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](1) == 0) else (div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](1) < 0) ^ (cluster_shape.__getitem__[::Indexer](1) < 0))) else div_s(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:@stdlib::@builtin::@int::@Int cond(eq(#lit.struct.extract<:@stdlib::@builtin::@int::@Int apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> rebind(:!lit.generator<[1]("self": !lit.struct<@stdlib::@utils::@index::@IndexList<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}>>, "idx": !lit.ref<@stdlib::@builtin::@int::@Int, imm *[0,0]> read_mem) -> !lit.struct<@stdlib::@builtin::@int::@Int>> @stdlib::@utils::@index::@IndexList::@"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:@stdlib::@builtin::@int::@Int {3}, :@stdlib::@builtin::@dtype::@DType {:dtype si64}, :trait<@stdlib::@builtin::@int::@Indexer> @stdlib::@builtin::@int::@Int>), cluster_shape, store_to_mem({1}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()]

Methods

__init__

__init__(out self, args: TMALoadOpArgs[a_type, b_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](0) == 0) else (div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](0) < 0) ^ (cluster_shape.__getitem__[::Indexer](0) < 0))) else div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), row_major(0 if (cluster_shape.__getitem__[::Indexer](1) == 0) else (div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](1) < 0) ^ (cluster_shape.__getitem__[::Indexer](1) < 0))) else div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType(), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()])

to_kernel_args

static to_kernel_args(a: LayoutTensor[a_type, layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], b: LayoutTensor[b_type, layout, origin, address_space=address_space, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], ctx: DeviceContext) -> TMALoadOpArgs[a_type, b_type, row_major(0 if (cluster_shape.__getitem__[::Indexer](0) == 0) else (div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](0) < 0) ^ (cluster_shape.__getitem__[::Indexer](0) < 0))) else div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({0})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({0}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), row_major(0 if (cluster_shape.__getitem__[::Indexer](1) == 0) else (div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1}))), "value">) + -1) if (((rem_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1}))), "value">) == 0) ^ True) & ((block_tile_shape.__getitem__[::Indexer](1) < 0) ^ (cluster_shape.__getitem__[::Indexer](1) < 0))) else div_s(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), block_tile_shape, store_to_mem({1})), "value">, #lit.struct.extract<:_stdlib::_builtin::_int::_Int cond(eq(#lit.struct.extract<:_stdlib::_builtin::_int::_Int apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1})), "value">, 0), {1}, apply(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm #lit.comptime.origin> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> rebind(:!lit.generator<[1]("self": !lit.struct<_stdlib::_utils::_index::_IndexList<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}>>, "idx": !lit.ref<_stdlib::_builtin::_int::_Int, imm *[0,0]> read_mem) -> !lit.struct<_stdlib::_builtin::_int::_Int>> _stdlib::_utils::_index::_IndexList::_"__getitem__[::Indexer](::IndexList[$0, $1],$2)"<:_stdlib::_builtin::_int::_Int {3}, :_stdlib::_builtin::_dtype::_DType {:dtype si64}, :trait<_stdlib::_builtin::_int::_Indexer> _stdlib::_builtin::_int::_Int>), cluster_shape, store_to_mem({1}))), "value">), block_tile_shape.__getitem__[::Indexer](2)), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType(), _tma_desc_tile_layout[::DType,::Int,::IndexList[$1, ::DType()]

Returns:

TMALoadOpArgs

__call__

__call__(self, a_smem_tile: LayoutTensor[dtype, layout, origin, address_space=AddressSpace(3), element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], b_smem_tile: LayoutTensor[dtype, layout, origin, address_space=AddressSpace(3), element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], m: SIMD[uint32, 1], n: SIMD[uint32, 1], k: SIMD[uint32, 1], ref [3] mbar: SharedMemBarrier)

Was this page helpful?

Thank you! We'll create more content like this.

Thank you for helping us improve!

😔 What went wrong?

Some code doesn’t work

It includes inaccurate information

It's missing information I need

It was difficult to understand

Other

Submit