Mojo struct

TMALoadOp

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

Fields

• ​a_tma_ptr (UnsafePointer[TMATensorTile[a_type, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[a_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=a_swizzle]()]]):
• ​b_tma_ptr (UnsafePointer[TMATensorTile[b_type, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[b_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=b_swizzle]()]]):

Implemented traits

AnyType, LoadOp, UnknownDestructibility

Aliases

__del__is_trivial

alias __del__is_trivial = True

a_tma_desc_layout

alias a_tma_desc_layout = _tma_desc_tile_layout[a_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=a_swizzle]()

a_tma_layout

alias a_tma_layout = Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2))

a_tma_type

alias a_tma_type = TMATensorTile[a_type, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[a_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=a_swizzle]()]

args_type

alias args_type = TMALoadOpArgs[a_type, b_type, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[a_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=a_swizzle](), _tma_desc_tile_layout[b_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=b_swizzle]()]

b_tma_desc_layout

alias b_tma_desc_layout = _tma_desc_tile_layout[b_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=b_swizzle]()

b_tma_layout

alias b_tma_layout = Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2))

b_tma_type

alias b_tma_type = TMATensorTile[b_type, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[b_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=b_swizzle]()]

Methods

__init__

__init__(out self, args: TMALoadOpArgs[a_type, b_type, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[a_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=a_swizzle](), _tma_desc_tile_layout[b_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=b_swizzle]()])

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, Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), Layout.row_major(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), _tma_desc_tile_layout[a_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](0) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=a_swizzle](), _tma_desc_tile_layout[b_type, 2, Index(0 if (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">) == 0) ^ True) & ((block_tile_shape.__getitem__[3, DType.int64, Int](1) < 0) ^ (cluster_shape.__getitem__[3, DType.int64, Int](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})), "_mlir_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})), "_mlir_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}))), "_mlir_value">), block_tile_shape.__getitem__[3, DType.int64, Int](2)), swizzle_mode=b_swizzle]()]

Returns:

TMALoadOpArgs

__call__

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

Was this page helpful?

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

Thank you for helping us improve!

😔 What went wrong?

Some code doesn’t work

It includes inaccurate information

It's missing information I need

It was difficult to understand

Other

Submit