[QST]How to make_tiled_copy for C layout after permutation?

[zhang110112110]

How to make_tiled_copy from register to global memory for this layout? m:32 n:64, threads: 128
(row0)t0v0...v7, t1v0-v7, t2v0-v7, t3v0-v7, t64v0...v7, t65v0-v7, t66v0-v7, t67v0-v7 (fp16)
t4v0...v7, t5v0-v7, t6v0-v7, t7v0-v7, t68v0...v7, t69v0-v7, t70v0-v7, t71v0-v7
...
(row8)t0v8...v15, t1v8...v15, t2vv8...v15 t3v8...v15 t64v8...v15 t65vv8...v15, t66v8...v15 t67v8...v15
t4v8...v15 t5vv8...v15 t6v8...v15,t7vv8...v15, t68vv8...v15, t69v8...v15, t70v8...v15, t71v8...v15
...
(row16)t32v0...v7, t33v0-v7, t34v0-v7, t35v0-v7, t96v0...v7, t97v0-v7, t98v0-v7, t99v0-v7 (fp16)
..
(row31)t60v8...v15, t61v8...v15, t62v8...v15, t63v8...v15 t124v8...v15 t125vv8...v15, t126v8...v15 t127v8...v15
permutation layout:

i have already transpose v2v3 and get 128B contiguous register for each threads. how to make_tiled_copy write 128B contiguous to global memory for this layout.

[zhang110112110]

using R2GCopyAtomC = Copy_Atom<UniversalCopy<uint128_t>, __half>;
using thr_layout = decltype(make_layout(make_shape(Int<32>{}, Int<4>{}),
make_stride(Int<4>{}, Int<1>{})));
using val_layout = decltype(make_layout(make_shape(Int<1>{},Int<8>{}),
make_stride(Int<8>{},Int<1>{}))); // layout is wrong
using R2GCopyC = decltype(make_tiled_copy(R2GCopyAtomC{},
thr_layout{},
val_layout{}));
auto tile_g = make_shape(Int<32>{}, Int<64>{});
auto gdC = flat_divide(gD, tile_g);
R2GCopyC r2g_tiled_copy_c;
auto r2g_thr_copy_c = r2g_tiled_copy_c.get_slice(threadIdx.x);
auto tCr_flat = cute::flatten(tCrD); // 1D连续
auto tCr_reshaped = cute::make_tensor<__half>(cute::make_layout(
cute::make_shape(Int<8>{}, Int<1>{}, Int<2>{}), // shape (8,2)
cute::make_stride(Int<1>{}, Int<8>{}, Int<8>{}) // stride
)
);
auto* dst_vec = reinterpret_cast<__half2*>(tCr_reshaped.data());
auto tCr2g = r2g_thr_copy_c.retile_S(tCr_reshaped);
auto tCg2r2 = r2g_thr_copy_c.partition_D(gdC);