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
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:
__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!