// Host-only checks for make_tiled_copy. No GPU instruction is executed.
/* Compile at the CUTLASS root:
g++ -x c++ -std=c++17 -Iinclude -I/usr/local/cuda/include \
-I/usr/local/cuda/include/cccl cute_make_tiled_copy_reference_trace.cu \
-o /tmp/cute_make_tiled_copy_reference_trace
*/
#include <cute/tensor.hpp>
#include <cute/atom/copy_traits_sm75.hpp>
#include <cute/atom/copy_traits_sm80.hpp>
#include <cute/atom/mma_traits_sm80.hpp>
#include <array>
#include <cstdio>
#include <cstdlib>
#include <type_traits>
using namespace cute;
void require(bool condition, char const* message) {
if (!condition) {
std::fprintf(stderr, "FAILED: %s\n", message);
std::exit(1);
}
template<class T>
void show(char const* name, T const& value) {
std::printf("%s = ", name);
print(value);
print("\n");
template<class TC, class Thr, class Val>
void trace(TC const& tc, Thr thr, Val val) {
auto mn = raked_product(thr, val);
show("thr_layout", thr);
show("val_layout", val);
show("logical_product", logical_product(thr, val));
show("layout_mn", mn);
show("coalesce(layout_mn)", coalesce(mn));
show("right_inverse(layout_mn)", right_inverse(mn));
show("Tiler_MN", typename TC::Tiler_MN{});
show("TiledLayout_TV", typename TC::TiledLayout_TV{});
show("AtomLayoutSrc", typename TC::AtomLayoutSrc{});
show("AtomLayoutDst", typename TC::AtomLayoutDst{});
show("AtomLayoutRef", typename TC::AtomLayoutRef{});
show("source_to_reference", right_inverse(typename TC::AtomLayoutRef{}).compose(typename TC::AtomLayoutSrc{}));
show("get_layoutS_TV", tc.get_layoutS_TV());
show("get_layoutD_TV", tc.get_layoutD_TV());
template<class TC>
void trace_partition_source() {
auto atom_tv = zipped_divide(typename TC::TiledLayout_TV{},
make_shape(typename TC::AtomNumThr{}, typename TC::AtomNumVal{}));
auto endpoint_to_ref = right_inverse(typename TC::AtomLayoutRef{}).compose(typename TC::AtomLayoutSrc{});
auto endpoint_tv = atom_tv.compose(endpoint_to_ref, _);
auto zipped = zip(endpoint_tv);
show("partition_S/atom_layout_TV", atom_tv);
show("partition_S/trg_layout_TV", endpoint_tv);
show("partition_S/zip(trg_layout_TV)", zipped);
show("partition_S/coalesce with profile", coalesce(zipped, Shape<_1,Shape<_1,_1>>{}));
show("partition_S/coalesce without profile", coalesce(zipped));
void verify_small_copy() {
std::puts("=== reading example: 32 threads, 4 half per thread, Ref=Src ===");
using SmallAtom = Copy_Atom<SM80_CP_ASYNC_CACHEALWAYS<uint64_t>, cutlass::half_t>;
auto thr_layout = Layout<Shape<_4,_8>,Stride<_8,_1>>{};
auto val_layout = Layout<Shape<_1,_4>>{};
auto tc = make_tiled_copy(SmallAtom{}, thr_layout, val_layout);
trace(tc, thr_layout, val_layout);
trace_partition_source<decltype(tc)>();
auto identity = make_identity_tensor(Shape<_4,_32>{});
auto mn = raked_product(thr_layout, val_layout);
auto tv = typename decltype(tc)::TiledLayout_TV{};
auto inv = right_inverse(mn);
alignas(16) std::array<cutlass::half_t, 4*40> storage{};
auto padded = make_tensor(make_smem_ptr(storage.data()),
Layout<Shape<_4,_32>,Stride<Int<40>,_1>>{});
std::array<int,128> visits{};
for (int t=0; t<32; ++t) {
auto ps = tc.get_slice(t).partition_S(identity);
auto pd = tc.get_slice(t).partition_D(identity);
auto padded_d = tc.get_slice(t).partition_D(padded);
for (int v=0; v<4; ++v) {
int m = t/8, n = 4*(t%8)+v;
int q = m+4*n;
require(ps(v) == make_coord(m,n), "reading example source coordinates");
require(pd(v) == make_coord(m,n), "reading example destination coordinates");
require(int(tv(make_coord(t,v))) == q, "reading example reference TV");
require(int(mn(q)) == t+32*v, "reading example TV encoding");
require(int(inv(t+32*v)) == q, "reading example inverse mapping");
require(&padded_d(v)-storage.data() == 40*m+n, "reading example padded destination");
++visits[32*m+n];
for (int count : visits) require(count == 1, "reading example exact coverage");
std::puts("Verified all 128 reading-example coordinates, inverse mappings, and padded destination offsets.\n");
template<class MMA>
void trace_mma_a_reference(MMA const& mma) {
auto thr_layout = mma.get_thr_layout_vmnk();
auto ref_A = make_layout(make_shape(tile_size<0>(mma), tile_size<2>(mma)));
auto t_tile = make_tile(mma.template permutation_mnk<0>(),
mma.template permutation_mnk<2>());
auto t_tensor = logical_divide(ref_A, t_tile);
auto zipped_t_tensor = zipped_divide(ref_A, t_tile);
auto a_tile = make_tile(make_layout(size<0>(typename MMA::AtomShape_MNK{})),
make_layout(size<2>(typename MMA::AtomShape_MNK{})));
auto logical_a_tensor = logical_divide(t_tensor, a_tile);
auto a_tensor = zipped_divide(t_tensor, a_tile);
auto tv_tensor = a_tensor.compose(typename MMA::AtomLayoutA_TV{}, _);
auto thr_tile = make_tile(_, make_tile(make_layout(size<1>(thr_layout)),
make_layout(size<3>(thr_layout))));
auto before_atom_tv = zipped_divide(a_tensor, thr_tile);
auto logical_thr_tensor = logical_divide(tv_tensor, thr_tile);
auto thr_tensor = zipped_divide(tv_tensor, thr_tile);
// The tutorial uses a dimension-only permutation. In this configuration,
// omitting the first divide preserves the element mapping.
auto direct_a_tensor = zipped_divide(ref_A, a_tile);
auto direct_tv_tensor = direct_a_tensor.compose(typename MMA::AtomLayoutA_TV{}, _);
auto without_first_divide = zipped_divide(direct_tv_tensor, thr_tile);
auto atile = make_tile(_, make_tile(
make_layout(make_shape(size<1>(thr_layout), size<2>(thr_layout)),
make_stride(Int<1>{}, Int<0>{})), _));
auto extended_threads = make_layout(thr_layout, complement(thr_layout));
auto inverse_threads = right_inverse(extended_threads);
auto thridx_2_thrid = composition(
make_layout(make_shape(size(thr_layout), Int<1>{}),
make_stride(Int<1>{}, Int<0>{})), inverse_threads);
auto with_n_warps = thr_tensor.compose(atile, _);
auto result = with_n_warps.compose(thridx_2_thrid, _);
show("get_layoutA/ref_A", ref_A);
show("get_layoutA/t_tile", t_tile);
show("get_layoutA/t_tensor", t_tensor);
show("get_layoutA/zipped_divide(ref_A,t_tile)", zipped_t_tensor);
show("get_layoutA/logical_divide(t_tensor,a_tile)", logical_a_tensor);
show("get_layoutA/a_tensor", a_tensor);
show("get_layoutA/tv_tensor", tv_tensor);
show("get_layoutA/thr_tile", thr_tile);
show("get_layoutA/thread_divide_before_atom_TV", before_atom_tv);
show("get_layoutA/logical_divide(tv_tensor,thr_tile)", logical_thr_tensor);
show("get_layoutA/thrfrg_A(ref_A)", thr_tensor);
show("get_layoutA/without_first_divide_dimension_only", without_first_divide);
show("get_layoutA/atile", atile);
show("get_layoutA/complement(ThrLayoutVMNK)", complement(thr_layout));
show("get_layoutA/inverse_extended_threads", inverse_threads);
show("get_layoutA/thridx_2_thrid", thridx_2_thrid);
show("get_layoutA/with_N_warps", with_n_warps);
show("get_layoutA/final_TV", result);
require(int(size<0>(thr_tensor)) == 64, "A thread domain before N-warp expansion");
require(int(size<0>(with_n_warps)) == 128, "A thread domain after N-warp expansion");
auto official = mma.get_layoutA_TV();
auto direct = mma.thrfrg_A(ref_A);
for (int t=0; t<128; ++t) {
int lane = t%32, wm = (t/32)%2;
require(int(thridx_2_thrid(t)) == t, "thread-ID adapter is identity in this configuration");
for (int v=0; v<8; ++v) {
int m = lane/4 + 8*((v/2)%2) + 16*wm;
int k = 2*(lane%4) + v%2 + 8*(v/4);
auto local_tv = make_coord(lane+32*wm, v);
auto full_tv = make_coord(t, v);
require(int(thr_tensor(local_tv)) == m+32*k, "thrfrg_A before N-warp expansion");
require(int(without_first_divide(local_tv)) == int(thr_tensor(local_tv)),
"first divide is elementwise redundant for this dimension-only permutation");
require(int(direct(local_tv)) == m+32*k, "manual thrfrg_A matches the source method");
require(int(with_n_warps(full_tv)) == m+32*k, "atile broadcasts A across N warps");
require(int(result(full_tv)) == int(official(full_tv)), "stepwise get_layoutA_TV composition");
require(int(result(make_coord(109,6))) == 347, "worked example t=109,v=6");
require(int(result(make_coord(45,6))) == 347, "worked example shares A across N warps");
std::puts("Verified all 1024 get_layoutA_TV composition paths and the t=109,v=6 example.\n");
void verify_thread_inverse_completion() {
auto l = Layout<Shape<_2,_2>, Stride<_1,_4>>{};
auto c = complement(l);
auto completed = make_layout(l,c);
auto inverse = right_inverse(completed);
auto project = make_layout(make_shape(size(l),_1{}), make_stride(_1{},_0{}));
auto adapter = composition(project,inverse);
require(int(size(right_inverse(l))) == 2, "direct inverse stops at the thread-ID gap");
require(int(size(inverse)) == 8, "completed inverse covers all eight IDs");
for (int t=0; t<8; ++t) {
require(int(completed(inverse(t))) == t, "completed thread-ID inverse round trip");
require(int(project(inverse(t))) == int(inverse(t))%4,
"projection discards the auxiliary coordinate, including nonzero ones");
for (int i=0; i<4; ++i) {
require(int(l(i)) == i%2+4*(i/2), "four-thread gapped layout formula");
require(int(adapter(l(i))) == i, "all four original coordinates are recovered");
show("thread inverse/gapped L", l);
show("thread inverse/complement(L)", c);
show("thread inverse/completed layout", completed);
show("thread inverse/right_inverse(L)", right_inverse(l));
show("thread inverse/right_inverse(completed)", inverse);
show("thread inverse/project", project);
// Exercise the same construction through the actual TiledMMA interface.
auto sparse = make_tiled_mma(SM80_16x8x16_F16F16F16F16_TN{},
Layout<Shape<_2,_2>,Stride<_1,_4>>{}, Tile<_32,_32,_16>{});
auto dense = make_tiled_mma(SM80_16x8x16_F16F16F16F16_TN{},
Layout<Shape<_2,_2>>{}, Tile<_32,_32,_16>{});
auto sparse_threads = sparse.get_thr_layout_vmnk();
auto sparse_completed = make_layout(sparse_threads,complement(sparse_threads));
auto sparse_project = make_layout(make_shape(size(sparse_threads),_1{}),
make_stride(_1{},_0{}));
auto sparse_adapter = composition(sparse_project,right_inverse(sparse_completed));
auto sparse_tv = sparse.get_layoutA_TV();
auto dense_tv = dense.get_layoutA_TV();
require(int(size(sparse_threads)) == 128, "sparse layout has 128 original thread coordinates");
require(int(size(right_inverse(sparse_threads))) == 64,
"direct sparse MMA inverse stops after the first 64 IDs");
require(int(size(right_inverse(sparse_completed))) == 256,
"complement completes the 256-ID ambient range");
require(int(size(right_inverse(dense.get_thr_layout_vmnk()))) == 128,
"official dense thread layout needs no nontrivial completion");
for (int i=0; i<128; ++i) {
int t = int(sparse_threads(i));
require(t == i%64+128*(i/64), "sparse MMA valid thread-ID ranges");
require(int(sparse_adapter(t)) == i, "sparse physical ID recovers the compact VMNK index");
require(int(sparse_tv(make_coord(t,v))) == int(dense_tv(make_coord(i,v))),
"sparse and dense MMA A coordinates agree at every valid thread value");
show("thread inverse/sparse MMA L", sparse_threads);
show("thread inverse/sparse MMA complement", complement(sparse_threads));
std::puts("Verified completed inverses and projection for all 8 ambient IDs; all 128 valid sparse MMA threads preserve 1024 A values.\n");
void verify_mma_a_permutation() {
using PermM = Layout<Shape<_2,_16>, Stride<_16,_1>>;
auto mma = make_tiled_mma(SM80_16x8x16_F16F16F16F16_TN{},
Layout<Shape<_2,_2>>{}, Tile<PermM,_32,_16>{});
using MMA = decltype(mma);
auto t_tensor = logical_divide(ref_A,
make_tile(mma.template permutation_mnk<0>(), mma.template permutation_mnk<2>()));
auto direct_tv = zipped_divide(ref_A, a_tile).compose(typename MMA::AtomLayoutA_TV{}, _);
auto without_first_divide = zipped_divide(direct_tv, thr_tile);
auto with_first_divide = mma.thrfrg_A(ref_A);
require(int(tile_size<0>(mma)) == 32 && int(tile_size<2>(mma)) == 16,
"permutation keeps the reference shape 32 by 16");
for (int u=0; u<32; ++u) {
int permuted_m = 16*(u%2)+u/2;
require(int(PermM{}(u)) == permuted_m, "flattened permutation formula");
for (int k=0; k<16; ++k) {
require(int(ref_A(make_coord(u,k))) == u+32*k,
"size alone does not apply the permutation to ref_A");
require(int(t_tensor(make_coord(make_coord(u,0),make_coord(k,0)))) == permuted_m+32*k,
"first divide applies the full permutation");
for (int t=0; t<64; ++t) {
int lane = t%32, wm = t/32;
int u = lane/4 + 8*((v/2)%2) + 16*wm;
auto tv = make_coord(t,v);
require(int(with_first_divide(tv)) == 16*(u%2)+u/2+32*k,
"permuted MMA-A grouping");
require(int(without_first_divide(tv)) == u+32*k,
"omitting first divide loses the permutation");
require(int(with_first_divide(make_coord(0,2))) == 4,
"permutation example thread 0 value 2 maps to A(4,0)");
require(int(without_first_divide(make_coord(0,2))) == 8,
"omitting first divide maps thread 0 value 2 to A(8,0)");
show("permuted get_layoutA/ref_A", ref_A);
show("permuted get_layoutA/t_tensor", t_tensor);
std::puts("Verified all 512 reordered A coordinates and 512 local thread values; omitting the first divide loses PermM.\n");
void verify_mma_b_n_repetition(MMA const& mma) {
auto ref_B = make_layout(make_shape(tile_size<1>(mma), tile_size<2>(mma)));
auto t_tile = make_tile(mma.template permutation_mnk<1>(),
auto t_tensor = logical_divide(ref_B, t_tile);
auto b_tile = make_tile(make_layout(size<1>(typename MMA::AtomShape_MNK{})),
auto b_tensor = zipped_divide(t_tensor, b_tile);
auto tv_tensor = b_tensor.compose(typename MMA::AtomLayoutB_TV{}, _);
auto thr_tile = make_tile(_, make_tile(make_layout(size<2>(thr_layout)),
auto b_tv = mma.get_layoutB_TV();
auto c_tv = mma.get_layoutC_TV();
auto identity_b = make_identity_tensor(Shape<_32,_16>{});
auto identity_c = make_identity_tensor(Shape<_32,_32>{});
using LD = Copy_Atom<SM75_U32x4_LDSM_N, cutlass::half_t>;
auto copy_b = make_tiled_copy_impl(LD{}, b_tv, Shape<_32,_16>{});
auto full_identity_b = make_identity_tensor(Shape<_128,_64>{});
alignas(16) std::array<cutlass::half_t,128*64> b_storage{};
auto b_smem = make_tensor(make_smem_ptr(b_storage.data()),
Layout<Shape<_128,_64>,Stride<_64,_1>>{});
show("MMA-B/AtomLayoutB_TV", typename MMA::AtomLayoutB_TV{});
show("MMA-B/ref_B", ref_B);
show("MMA-B/b_tensor", b_tensor);
show("MMA-B/tv_tensor", tv_tensor);
show("MMA-B/logical_divide(tv_tensor,thr_tile)", logical_thr_tensor);
show("MMA-B/thrfrg_B(ref_B)", thr_tensor);
show("MMA-B/get_layoutB_TV", b_tv);
show("MMA-B/partition_B shape", shape(mma.get_slice(0).partition_B(identity_b)));
show("MMA-C/partition_C shape", shape(mma.get_slice(0).partition_C(identity_c)));
require(int(size<1,0>(b_tensor)) == 4, "B has four N-direction Atom blocks");
require(int(size<1,1,0>(thr_tensor)) == 2, "two B N repeats remain after assigning two N warps");
std::array<int,512> b_visits{};
std::array<int,1024> c_visits{};
int lane=t%32, wm=(t/32)%2, wn=t/64;
auto pb=mma.get_slice(t).partition_B(identity_b);
auto pc=mma.get_slice(t).partition_C(identity_c);
auto copy_d=copy_b.get_slice(t).partition_D(identity_b);
require(int(size(pb)) == 8, "four B values per MMA times two N repeats");
require(int(size(pc)) == 8, "four accumulators per MMA times two N repeats");
for (int repeat_n=0; repeat_n<2; ++repeat_n) {
for (int atom_v=0; atom_v<4; ++atom_v) {
int n=lane/4+8*wn+16*repeat_n;
int k=2*(lane%4)+atom_v%2+8*(atom_v/2);
int v=atom_v+4*repeat_n;
require(pb(atom_v,repeat_n,0) == make_coord(n,k), "B N-warp and N-repeat coordinate formula");
require(copy_d(v) == pb(atom_v,repeat_n,0),
"one eight-value Copy destination contains two four-value MMA B fragments");
require(int(b_tv(make_coord(t,v))) == n+32*k, "B reference TV includes both N expansions");
require(int(thr_tensor(make_coord(lane+32*wn,v))) == n+32*k,
"B local thread layout ignores the M warp coordinate");
++b_visits[n+32*k];
int cm=lane/4+8*(atom_v/2)+16*wm;
int cn=2*(lane%4)+atom_v%2+8*wn+16*repeat_n;
require(pc(atom_v,0,repeat_n) == make_coord(cm,cn), "C N-warp and N-repeat coordinate formula");
require(int(c_tv(make_coord(t,v))) == cm+32*cn, "C reference TV covers the full 32x32 tile");
++c_visits[cm+32*cn];
auto thr_mma = mma.get_slice(t);
auto thr_copy = copy_b.get_slice(t);
auto regs = thr_mma.partition_fragment_B(b_smem);
auto copy_regs = thr_copy.retile_D(regs);
auto mma_coords = thr_mma.partition_B(full_identity_b);
auto copy_coords = thr_copy.partition_D(full_identity_b);
auto source_coords = thr_copy.partition_S(full_identity_b);
auto source = thr_copy.partition_S(b_smem);
if (t == 0) {
show("MMA-B/register layout", regs.layout());
show("MMA-B/retile_D register layout", copy_regs.layout());
show("MMA-B/Copy partition_D shape", shape(copy_coords));
show("MMA-B/Copy partition_S shape", shape(source));
std::puts("MMA-B/lane 0 Copy values -> (Atom value, N repeat) -> B coordinate:");
std::printf(" %d -> (%d,%d) -> ",v,v%4,v/4);
print(copy_d(v));
require(int(size<0>(regs)) == 4 && int(size<1>(regs)) == 8,
"full B fragment has four Atom values and eight N repeats");
require(int(size<0>(copy_regs)) == 8 && int(size<1>(copy_regs)) == 4,
"Copy B groups eight values and four N repeats");
for (int repeat_n=0; repeat_n<4; ++repeat_n) {
for (int repeat_k=0; repeat_k<4; ++repeat_k) {
int atom_v = v%4, mma_n = 2*repeat_n+v/4;
require(©_regs(v,repeat_n,repeat_k) == ®s(atom_v,mma_n,repeat_k),
"B retile_D pairs MMA fragments by aliasing their original registers");
require(copy_coords(v,repeat_n,repeat_k) == mma_coords(atom_v,mma_n,repeat_k),
"Copy and MMA B coordinates agree after regrouping");
int sn = lane%8+8*wn+16*(lane/16)+32*repeat_n;
int sk = v+8*((lane/8)%2)+16*repeat_k;
require(source_coords(v,repeat_n,repeat_k) == make_coord(sn,sk),
"B source rows match the grouped Copy destination");
require(&source(v,repeat_n,repeat_k)-b_storage.data() == 64*sn+sk,
"B row-major shared-memory source offset");
require(&source(v,repeat_n,repeat_k) == &source(0,repeat_n,repeat_k)+v,
"B ldmatrix source lane supplies eight contiguous half values");
for (int visits:b_visits) require(visits == 2, "B is shared across two M warps");
for (int visits:c_visits) require(visits == 1, "four warps and two N repeats cover C exactly once");
std::puts("Verified all 1024 B values with two N repeats and twofold M replication, and exact coverage of all 1024 C elements.\n");
std::puts("Verified all 16384 B register aliases, Copy/MMA coordinate correspondences, and row-major source offsets for two-fragment Copy grouping.\n");
void verify_mma_a_copy() {
std::puts("=== ldmatrix: explicit reference TV from the actual sgemm_sm80 MMA ===");
trace_mma_a_reference(mma);
verify_mma_b_n_repetition(mma);
auto reference_tv = mma.get_layoutA_TV();
auto tc = make_tiled_copy_impl(LD{}, reference_tv, Shape<_32,_16>{});
show("MMA/AtomLayoutA_TV", typename decltype(mma)::AtomLayoutA_TV{});
show("MMA/ThrLayoutVMNK", mma.get_thr_layout_vmnk());
show("MMA/tile_shape", tile_shape(mma));
show("MMA/get_layoutA_TV", mma.get_layoutA_TV());
show("MMA-A/Tiler_MN", typename decltype(tc)::Tiler_MN{});
show("MMA-A/TiledLayout_TV", typename decltype(tc)::TiledLayout_TV{});
show("MMA-A/get_layoutS_TV", tc.get_layoutS_TV());
show("MMA-A/get_layoutD_TV", tc.get_layoutD_TV());
show("MMA-A/AtomLayoutSrc", typename decltype(tc)::AtomLayoutSrc{});
show("MMA-A/AtomLayoutDst", typename decltype(tc)::AtomLayoutDst{});
show("MMA-A/AtomLayoutRef", typename decltype(tc)::AtomLayoutRef{});
auto identity = make_identity_tensor(Shape<_32,_16>{});
auto dst_tv = tc.get_layoutD_TV();
std::array<int,512> s_visits{}, d_visits{};
int lane = t%32, warp_m = (t/32)%2;
auto pa = mma.get_slice(t).partition_A(identity);
int dm = lane/4 + 8*((v/2)%2) + 16*warp_m;
int dk = 2*(lane%4) + v%2 + 8*(v/4);
int sm = lane%16 + 16*warp_m;
int sk = v + 8*(lane/16);
require(pa(v) == make_coord(dm,dk), "MMA A coordinate formula");
require(pd(v) == pa(v), "Copy destination matches MMA A value order");
require(ps(v) == make_coord(sm,sk), "MMA A source row formula");
require(int(tv(make_coord(t,v))) == dm+32*dk, "MMA A reference TV formula");
require(int(dst_tv(make_coord(t,v))) == int(tv(make_coord(t,v))),
"ldmatrix Ref=Dst preserves MMA A reference mapping");
++s_visits[sm+32*sk];
++d_visits[dm+32*dk];
for (int i=0; i<512; ++i) {
require(s_visits[i] == 2, "A is loaded twice across the two N warps");
require(d_visits[i] == 2, "A is consumed twice across the two N warps");
// Keep the tutorial's MMA, Copy Atom, and CTA dimensions. Simplify storage
// to a two-dimensional row-major view for this layout-only explanation.
auto smem_layout = Layout<Shape<_128,_64>,Stride<_64,_1>>{};
alignas(16) std::array<cutlass::half_t,128*64> storage{};
auto smem = make_tensor(make_smem_ptr(storage.data()), smem_layout);
show("MMA-A/simplified row-major shared layout", smem_layout);
auto full_identity = make_identity_tensor(Shape<_128,_64>{});
auto thr_copy = tc.get_slice(t);
auto regs = thr_mma.partition_fragment_A(smem);
auto mma_coords = thr_mma.partition_A(full_identity);
auto copy_coords = thr_copy.partition_D(full_identity);
auto source_coords = thr_copy.partition_S(full_identity);
auto source = thr_copy.partition_S(smem);
show("MMA-A/register layout", regs.layout());
show("MMA-A/retile_D register layout", copy_regs.layout());
show("MMA-A/partition_S shape", shape(thr_copy.partition_S(smem)));
for (int m=0; m<4; ++m) {
for (int k=0; k<4; ++k) {
require(©_regs(v,m,k) == ®s(v,m,k), "A retile_D aliases the same register");
require(copy_coords(v,m,k) == mma_coords(v,m,k), "A copy and MMA agree across CTA repeats");
int sm = t%32%16 + 16*((t/32)%2) + 32*m;
int sk = v + 8*((t%32)/16) + 16*k;
require(source_coords(v,m,k) == make_coord(sm,sk), "A source coordinates across CTA repeats");
require(&source(v,m,k)-storage.data() == 64*sm+sk,
"A source address in simplified row-major storage");
require(&source(v,m,k) == &source(0,m,k)+v,
"Each ldmatrix source lane supplies eight contiguous half values");
std::puts("Verified all 1024 MMA-A source/destination values, twofold N replication, 16384 CTA register aliases, and 16384 row-major source addresses.");
// Deliberately misuse the source TV as a reference TV, although Ref=Dst.
auto wrong = make_tiled_copy_impl(LD{}, tc.get_layoutS_TV(),
typename decltype(tc)::Tiler_MN{});
auto wrong_source = wrong.get_slice(0).partition_S(identity);
std::printf("wrong source lane 0 coordinates:");
for (int u=0; u<8; ++u) {
require(wrong_source(u) == make_coord(u/2,u%2), "wrong source TV must change the requested A row");
print(" ");
print(wrong_source(u));
print("\n\n");
// Introduce the convenience constructor after the explicit layout checks.
auto via_helper = make_tiled_copy_A(LD{}, mma);
static_assert(std::is_same<decltype(tc),decltype(via_helper)>::value,
"make_tiled_copy_A must preserve the explicit reference-TV construction");
std::puts("make_tiled_copy_A produces the same TiledCopy type as the explicit reference-TV construction.\n");
int main() {
verify_small_copy();
verify_mma_a_copy();
verify_thread_inverse_completion();
verify_mma_a_permutation();
std::puts("=== cp.async: actual sgemm_sm80 tutorial layouts, Ref=Src ===");
using CP = Copy_Atom<SM80_CP_ASYNC_CACHEALWAYS<uint128_t>, cutlass::half_t>;
auto cp_thr = Layout<Shape<_16,_8>,Stride<_8,_1>>{};
auto cp_val = Layout<Shape<_1,_8>>{};
auto cp = make_tiled_copy(CP{}, cp_thr, cp_val);
trace(cp, cp_thr, cp_val);
trace_partition_source<decltype(cp)>();
auto cp_shape = Shape<_16,_64>{};
auto cp_identity = make_identity_tensor(cp_shape);
auto cp_s = cp.get_layoutS_TV();
auto cp_d = cp.get_layoutD_TV();
auto cp_mn = raked_product(cp_thr, cp_val);
std::array<int, 1024> cp_visits{};
alignas(16) std::array<cutlass::half_t, 16*72> storage{};
Layout<Shape<_16,_64>,Stride<Int<72>,_1>>{});
for (int t = 0; t < 128; ++t) {
auto ps = cp.get_slice(t).partition_S(cp_identity);
auto pd = cp.get_slice(t).partition_D(cp_identity);
auto padded_d = cp.get_slice(t).partition_D(padded);
for (int v = 0; v < 8; ++v) {
int m = t/8, n = 8*(t%8)+v;
int q = m+16*n;
require(int(cp_s(make_coord(t,v))) == q, "cp source TV");
require(int(cp_d(make_coord(t,v))) == q, "cp destination TV");
require(int(cp_mn(q)) == t+128*v, "cp coordinate to TV encoding");
require(ps(v) == make_coord(m,n), "cp source identity partition");
require(pd(v) == make_coord(m,n), "cp destination identity partition");
require(&padded_d(v)-storage.data() == 72*m+n, "cp padded destination physical address");
++cp_visits[m*64+n];
for (int count : cp_visits) require(count == 1, "cp exact coverage");
std::puts("Verified all 1024 cp.async source/destination coordinates and padded destination offsets.");
std::puts("\n=== right_inverse (8,16):(16,1) subexample ===");
auto small_thr = Layout<Shape<_4,_8>,Stride<_8,_1>>{};
auto small_val = Layout<Shape<_1,_4>>{};
auto small_mn = raked_product(small_thr, small_val);
auto small_inv = right_inverse(small_mn);
show("layout_mn", small_mn);
show("coalesce(layout_mn)", coalesce(small_mn));
show("right_inverse(layout_mn)", small_inv);
for(int i=0; i<128; ++i) {
require(int(small_inv(i)) == 16*(i%8)+i/8, "small inverse formula");
require(int(small_mn(small_inv(i))) == i, "small right inverse round trip");
std::puts("PASS: make_tiled_copy layout checks only; no GPU instruction executed.");