diff --git a/07-swizzling/swizzling.cu b/07-swizzling/swizzling.cu index 6374822..236beb5 100644 --- a/07-swizzling/swizzling.cu +++ b/07-swizzling/swizzling.cu @@ -29,9 +29,9 @@ __global__ void swizzling(void *__restrict__ Cptr, using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutO = typename Spec::SmemLayoutO; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kShmSizeA = Spec::kShmSizeA; constexpr int kShmSizeB = Spec::kShmSizeB; @@ -49,18 +49,18 @@ __global__ void swizzling(void *__restrict__ Cptr, Tensor mC = make_tensor(make_gmem_ptr((ComputeTypeC *)Cptr), make_shape(m, n), make_stride(n, Int<1>{})); // (M, N) Tensor mO = make_tensor(make_gmem_ptr((OutType *)Outptr), make_shape(m, n), make_stride(n, Int<1>{})); // (M, N) - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(0, 0, 0); - Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (kTileM, kTileK) - Tensor gB = local_tile(mB, tiler, coord, Step{}); // (kTileN, kTileK) - Tensor gC = local_tile(mC, tiler, coord, Step<_1, _1, X>{}); // (kTileM, kTileN) - Tensor gO = local_tile(mO, tiler, coord, Step<_1, _1, X>{}); // (kTileM, kTileN) + Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (kBlockM, kBlockK) + Tensor gB = local_tile(mB, tiler, coord, Step{}); // (kBlockN, kBlockK) + Tensor gC = local_tile(mC, tiler, coord, Step<_1, _1, X>{}); // (kBlockM, kBlockN) + Tensor gO = local_tile(mO, tiler, coord, Step<_1, _1, X>{}); // (kBlockM, kBlockN) - Tensor sA = make_tensor(make_smem_ptr((ComputeTypeA *)Aptr_smem), SmemLayoutA{}); // (kTileM, kTileK) - Tensor sB = make_tensor(make_smem_ptr((ComputeTypeB *)Bptr_smem), SmemLayoutB{}); // (kTileN, kTileK) - Tensor sC = make_tensor(make_smem_ptr((ComputeTypeC *)Cptr_smem), SmemLayoutC{}); // (kTileM, kTileN) - Tensor sO = make_tensor(make_smem_ptr((OutType *)Optr_smem), SmemLayoutO{}); // (kTileM, kTileN) + Tensor sA = make_tensor(make_smem_ptr((ComputeTypeA *)Aptr_smem), SmemLayoutA{}); // (kBlockM, kBlockK) + Tensor sB = make_tensor(make_smem_ptr((ComputeTypeB *)Bptr_smem), SmemLayoutB{}); // (kBlockN, kBlockK) + Tensor sC = make_tensor(make_smem_ptr((ComputeTypeC *)Cptr_smem), SmemLayoutC{}); // (kBlockM, kBlockN) + Tensor sO = make_tensor(make_smem_ptr((OutType *)Optr_smem), SmemLayoutO{}); // (kBlockM, kBlockN) typename Spec::TiledMMA tiled_mma; ThrMMA thr_mma = tiled_mma.get_slice(tid); @@ -140,9 +140,9 @@ __global__ void swizzling(void *__restrict__ Cptr, constexpr int kMmaTileN = Spec::kMmaTileN; constexpr int kMmaTileK = Spec::kMmaTileK; - constexpr int NTilesM = kTileM / kMmaTileM; - constexpr int NTilesN = kTileN / kMmaTileN; - constexpr int NTilesK = kTileK / kMmaTileK; + constexpr int NTilesM = kBlockM / kMmaTileM; + constexpr int NTilesN = kBlockN / kMmaTileN; + constexpr int NTilesK = kBlockK / kMmaTileK; #pragma unroll for (int m_tile = 0; m_tile < NTilesM; ++m_tile) { @@ -213,18 +213,18 @@ template + int kBlockM_ = 128, + int kBlockN_ = 128, + int kBlockK_ = 64> struct KernelSpec { using OutType = OutType_; using ComputeTypeA = ComputeTypeA_; using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; using MMA_op = std::conditional_t< std::is_same_v && std::is_same_v && @@ -246,7 +246,7 @@ struct KernelSpec { using MMA_shape = typename MMA_traits::Shape_MNK; static constexpr int kMmaThrExpandM = 2; - static constexpr int kMmaThrExpandN = 4; + static constexpr int kMmaThrExpandN = 2; static constexpr int kMmaThrExpandK = 1; static constexpr int kMmaValExpandM = 1; @@ -292,34 +292,24 @@ struct KernelSpec { using CopyO_S2G_atom = Copy_Atom; static constexpr int kThreadNum = size(TiledMMA{}); - static constexpr int kThreadsPerWarp = 32; - static constexpr int kTileM_Copy = cute::min(kThreadsPerWarp, kTileM); - static constexpr int kTileN_Copy = cute::min(kThreadsPerWarp, kTileN); - - // Here we omit cases that `kAlignedCopyItems < 1` - static constexpr int kAlignedCopyItemsA = - cute::min(128 / 8 / sizeof(ComputeTypeA), kTileK *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsB = - cute::min(128 / 8 / sizeof(ComputeTypeB), kTileK *kTileN_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsC = - cute::min(128 / 8 / sizeof(ComputeTypeC), kTileN *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsO = cute::min(128 / 8 / sizeof(OutType), kTileN *kTileM_Copy / kThreadNum); + static constexpr int kBlockK_Copy = cute::min(64, kBlockK) / 8; + static constexpr int kBlockN_Copy = cute::min(64, kBlockN) / 8; using TiledCopyA_G2S = decltype(make_tiled_copy(CopyA_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyB_G2S = decltype(make_tiled_copy(CopyB_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyC_G2S = decltype(make_tiled_copy(CopyC_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyA_S2R = decltype(make_tiled_copy_A(CopyA_S2R_atom{}, TiledMMA{})); using TiledCopyB_S2R = decltype(make_tiled_copy_B(CopyB_S2R_atom{}, TiledMMA{})); @@ -330,36 +320,36 @@ struct KernelSpec { using TiledCopyC_S2G = decltype(make_tiled_copy(CopyC_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyO_S2G = decltype(make_tiled_copy(CopyO_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using SmemLayoutAtomAB = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomC = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutA = - decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutB = - decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutC = - decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutO = - decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); static constexpr int kShmSizeA = cosize(SmemLayoutA{}) * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize(SmemLayoutB{}) * sizeof(ComputeTypeB); static constexpr int kShmSizeC = cosize(SmemLayoutC{}) * sizeof(ComputeTypeC); static constexpr int kShmSizeO = cosize(SmemLayoutO{}) * sizeof(OutType); - static constexpr int kShmSize = kShmSizeA + kShmSizeB + kShmSizeC; + static constexpr int kShmSize = cute::max(kShmSizeA + kShmSizeB + kShmSizeC, kShmSizeO); }; } // namespace spec @@ -462,14 +452,8 @@ torch::Tensor run_swizzling(const torch::Tensor a, const torch::Tensor b, std::o using Spec = spec::KernelSpec; - // cute::print_layout(typename decltype(kernel_spec)::SmemLayoutAtomAB{}); - // cute::print_layout(typename decltype(kernel_spec)::SmemLayoutA{}); - // if (sizeof(ComputeTypeC) == 2) { - // cute::print_latex(typename decltype(kernel_spec)::TiledCopyC_S2R{}); - // } - dim3 block = Spec::kThreadNum; - dim3 grid((N + Spec::kTileN - 1) / Spec::kTileN, (M + Spec::kTileM - 1) / Spec::kTileM); + dim3 grid((N + Spec::kBlockN - 1) / Spec::kBlockN, (M + Spec::kBlockM - 1) / Spec::kBlockM); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d Bytes\n", block.x, block.y, diff --git a/08-dynamic-mma/dynamic_mma.cu b/08-dynamic-mma/dynamic_mma.cu index 9f18bff..0663f9b 100644 --- a/08-dynamic-mma/dynamic_mma.cu +++ b/08-dynamic-mma/dynamic_mma.cu @@ -29,9 +29,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void dynamic_mma(void *__restrict using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutO = typename Spec::SmemLayoutO; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kShmSizeA = Spec::kShmSizeA; constexpr int kShmSizeB = Spec::kShmSizeB; @@ -51,7 +51,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void dynamic_mma(void *__restrict Tensor mC = make_tensor(make_gmem_ptr((ComputeTypeC *)Cptr), make_shape(M, N), make_stride(N, Int<1>{})); // (M, N) Tensor mO = make_tensor(make_gmem_ptr((OutType *)Outptr), make_shape(M, N), make_stride(N, Int<1>{})); // (M, N) - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -168,7 +168,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void dynamic_mma(void *__restrict copy_if(g2s_tiled_copy_c, tCpC_g2s, tCgC_g2s, tCsC_g2s); } - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); for (int ik = 0; ik < NTilesK; ++ik) { // Clear the smem tiles to account for predicated off loads @@ -302,18 +302,18 @@ template + int kBlockM_ = 128, + int kBlockN_ = 128, + int kBlockK_ = 64> struct KernelSpec { using OutType = OutType_; using ComputeTypeA = ComputeTypeA_; using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; using MMA_op = std::conditional_t< std::is_same_v && std::is_same_v && @@ -387,34 +387,24 @@ struct KernelSpec { using CopyO_S2G_atom = Copy_Atom; static constexpr int kThreadNum = size(TiledMMA{}); - static constexpr int kThreadsPerWarp = 32; - static constexpr int kTileM_Copy = cute::min(kThreadsPerWarp, kTileM); - static constexpr int kTileN_Copy = cute::min(kThreadsPerWarp, kTileN); - - // Here we omit cases that `kAlignedCopyItems < 1` - static constexpr int kAlignedCopyItemsA = - cute::min(128 / 8 / sizeof(ComputeTypeA), kTileK *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsB = - cute::min(128 / 8 / sizeof(ComputeTypeB), kTileK *kTileN_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsC = - cute::min(128 / 8 / sizeof(ComputeTypeC), kTileN *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsO = cute::min(128 / 8 / sizeof(OutType), kTileN *kTileM_Copy / kThreadNum); + static constexpr int kBlockK_Copy = cute::min(64, kBlockK) / 8; + static constexpr int kBlockN_Copy = cute::min(64, kBlockN) / 8; using TiledCopyA_G2S = decltype(make_tiled_copy(CopyA_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyB_G2S = decltype(make_tiled_copy(CopyB_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyC_G2S = decltype(make_tiled_copy(CopyC_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyA_S2R = decltype(make_tiled_copy_A(CopyA_S2R_atom{}, TiledMMA{})); using TiledCopyB_S2R = decltype(make_tiled_copy_B(CopyB_S2R_atom{}, TiledMMA{})); @@ -425,36 +415,36 @@ struct KernelSpec { using TiledCopyC_S2G = decltype(make_tiled_copy(CopyC_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyO_S2G = decltype(make_tiled_copy(CopyO_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using SmemLayoutAtomAB = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomC = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutA = - decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutB = - decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutC = - decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutO = - decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); static constexpr int kShmSizeC = cosize_v * sizeof(ComputeTypeC); static constexpr int kShmSizeO = cosize_v * sizeof(OutType); - static constexpr int kShmSize = kShmSizeA + kShmSizeB + kShmSizeC; + static constexpr int kShmSize = cute::max(kShmSizeA + kShmSizeB + kShmSizeC, kShmSizeO); }; } // namespace spec @@ -508,9 +498,9 @@ template constexpr bool needs_precisio return !std::is_same_v; } -template ; + using Spec = spec::KernelSpec; dim3 block = Spec::kThreadNum; - dim3 grid(cute::ceil_div(N, Spec::kTileN), cute::ceil_div(M, Spec::kTileM)); + dim3 grid(cute::ceil_div(N, Spec::kBlockN), cute::ceil_div(M, Spec::kBlockM)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d Bytes\n", block.x, block.y, diff --git a/09-pipelining/pipelining.cu b/09-pipelining/pipelining.cu index c509a7c..18ec335 100644 --- a/09-pipelining/pipelining.cu +++ b/09-pipelining/pipelining.cu @@ -25,9 +25,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void pipelining(void *__restrict_ using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutO = typename Spec::SmemLayoutO; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kShmSizeA = Spec::kShmSizeA; constexpr int kShmSizeB = Spec::kShmSizeB; constexpr int G2S_Stages = Spec::G2S_Stages; @@ -53,7 +53,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void pipelining(void *__restrict_ Tensor mC = make_tensor(make_gmem_ptr((ComputeTypeC *)Cptr), make_shape(M, N), make_stride(N, Int<1>{})); // (M, N) Tensor mO = make_tensor(make_gmem_ptr((OutType *)Outptr), make_shape(M, N), make_stride(N, Int<1>{})); // (M, N) - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -158,7 +158,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void pipelining(void *__restrict_ copy_if(g2s_tiled_copy_c, tCpC_g2s, tCgC_g2s, tCsC_g2s); } - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); clear(tAsA_g2s); clear(tBsB_g2s); @@ -349,9 +349,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -359,9 +359,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int G2S_Stages = G2S_Stages_; static_assert(G2S_Stages >= 2, "G2S_Stages should not be less than 2."); @@ -432,34 +432,24 @@ struct KernelSpec { using CopyO_S2G_atom = Copy_Atom; static constexpr int kThreadNum = size(TiledMMA{}); - static constexpr int kThreadsPerWarp = 32; - static constexpr int kTileM_Copy = cute::min(kThreadsPerWarp, kTileM); - static constexpr int kTileN_Copy = cute::min(kThreadsPerWarp, kTileN); - - // Here we omit cases that `kAlignedCopyItems < 1` - static constexpr int kAlignedCopyItemsA = - cute::min(128 / 8 / sizeof(ComputeTypeA), kTileK *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsB = - cute::min(128 / 8 / sizeof(ComputeTypeB), kTileK *kTileN_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsC = - cute::min(128 / 8 / sizeof(ComputeTypeC), kTileN *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsO = cute::min(128 / 8 / sizeof(OutType), kTileN *kTileM_Copy / kThreadNum); + static constexpr int kBlockK_Copy = cute::min(64, kBlockK) / 8; + static constexpr int kBlockN_Copy = cute::min(64, kBlockN) / 8; using TiledCopyA_G2S = decltype(make_tiled_copy(CopyA_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyB_G2S = decltype(make_tiled_copy(CopyB_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyC_G2S = decltype(make_tiled_copy(CopyC_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyA_S2R = decltype(make_tiled_copy_A(CopyA_S2R_atom{}, TiledMMA{})); using TiledCopyB_S2R = decltype(make_tiled_copy_B(CopyB_S2R_atom{}, TiledMMA{})); @@ -470,28 +460,28 @@ struct KernelSpec { using TiledCopyC_S2G = decltype(make_tiled_copy(CopyC_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyO_S2G = decltype(make_tiled_copy(CopyO_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using SmemLayoutAtomAB = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomC = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutA = - decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}, Int{}))); + decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}, Int{}))); using SmemLayoutB = - decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}, Int{}))); - using SmemLayoutC = decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}))); - using SmemLayoutO = decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(SmemLayoutAtomAB{}, make_shape(Int{}, Int{}, Int{}))); + using SmemLayoutC = decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}))); + using SmemLayoutO = decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}))); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -550,9 +540,9 @@ template constexpr bool needs_precisio return !std::is_same_v; } -template ; + using Spec = spec::KernelSpec; dim3 block = Spec::kThreadNum; - dim3 grid(cute::ceil_div(N, Spec::kTileN), cute::ceil_div(M, Spec::kTileM)); + dim3 grid(cute::ceil_div(N, Spec::kBlockN), cute::ceil_div(M, Spec::kBlockM)); constexpr int kShmSizeA = Spec::kShmSizeA; constexpr int kShmSizeB = Spec::kShmSizeB; diff --git a/09-pipelining/pipelining_no_reg_prefetch.cu b/09-pipelining/pipelining_no_reg_prefetch.cu index e9dd2fb..11872dd 100644 --- a/09-pipelining/pipelining_no_reg_prefetch.cu +++ b/09-pipelining/pipelining_no_reg_prefetch.cu @@ -25,9 +25,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void pipelining(void *__restrict_ using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutO = typename Spec::SmemLayoutO; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kShmSizeA = Spec::kShmSizeA; constexpr int kShmSizeB = Spec::kShmSizeB; constexpr int G2S_Stages = Spec::G2S_Stages; @@ -52,7 +52,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void pipelining(void *__restrict_ Tensor mC = make_tensor(make_gmem_ptr((ComputeTypeC *)Cptr), make_shape(M, N), make_stride(N, Int<1>{})); // (M, N) Tensor mO = make_tensor(make_gmem_ptr((OutType *)Outptr), make_shape(M, N), make_stride(N, Int<1>{})); // (M, N) - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -157,7 +157,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void pipelining(void *__restrict_ copy_if(g2s_tiled_copy_c, tCpC_g2s, tCgC_g2s, tCsC_g2s); } - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); clear(tAsA_g2s); clear(tBsB_g2s); @@ -330,9 +330,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -340,9 +340,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int G2S_Stages = G2S_Stages_; static_assert(G2S_Stages >= 2, "G2S_Stages should not be less than 2."); @@ -413,34 +413,24 @@ struct KernelSpec { using CopyO_S2G_atom = Copy_Atom; static constexpr int kThreadNum = size(TiledMMA{}); - static constexpr int kThreadsPerWarp = 32; - static constexpr int kTileM_Copy = cute::min(kThreadsPerWarp, kTileM); - static constexpr int kTileN_Copy = cute::min(kThreadsPerWarp, kTileN); - - // Here we omit cases that `kAlignedCopyItems < 1` - static constexpr int kAlignedCopyItemsA = - cute::min(128 / 8 / sizeof(ComputeTypeA), kTileK *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsB = - cute::min(128 / 8 / sizeof(ComputeTypeB), kTileK *kTileN_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsC = - cute::min(128 / 8 / sizeof(ComputeTypeC), kTileN *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsO = cute::min(128 / 8 / sizeof(OutType), kTileN *kTileM_Copy / kThreadNum); + static constexpr int kBlockK_Copy = cute::min(64, kBlockK) / 8; + static constexpr int kBlockN_Copy = cute::min(64, kBlockN) / 8; using TiledCopyA_G2S = decltype(make_tiled_copy(CopyA_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyB_G2S = decltype(make_tiled_copy(CopyB_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyC_G2S = decltype(make_tiled_copy(CopyC_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyA_S2R = decltype(make_tiled_copy_A(CopyA_S2R_atom{}, TiledMMA{})); using TiledCopyB_S2R = decltype(make_tiled_copy_B(CopyB_S2R_atom{}, TiledMMA{})); @@ -451,33 +441,33 @@ struct KernelSpec { using TiledCopyC_S2G = decltype(make_tiled_copy(CopyC_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyO_S2G = decltype(make_tiled_copy(CopyO_S2G_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using SmemLayoutAtomA = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomB = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomC = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomO = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutA = - decltype(tile_to_shape(SmemLayoutAtomA{}, make_shape(Int{}, Int{}, Int{}))); + decltype(tile_to_shape(SmemLayoutAtomA{}, make_shape(Int{}, Int{}, Int{}))); using SmemLayoutB = - decltype(tile_to_shape(SmemLayoutAtomB{}, make_shape(Int{}, Int{}, Int{}))); - using SmemLayoutC = decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}))); - using SmemLayoutO = decltype(tile_to_shape(SmemLayoutAtomO{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(SmemLayoutAtomB{}, make_shape(Int{}, Int{}, Int{}))); + using SmemLayoutC = decltype(tile_to_shape(SmemLayoutAtomC{}, make_shape(Int{}, Int{}))); + using SmemLayoutO = decltype(tile_to_shape(SmemLayoutAtomO{}, make_shape(Int{}, Int{}))); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -536,9 +526,9 @@ template constexpr bool needs_precisio return !std::is_same_v; } -template ; + using Spec = spec::KernelSpec; dim3 block = Spec::kThreadNum; - dim3 grid(cute::ceil_div(N, Spec::kTileN), cute::ceil_div(M, Spec::kTileM)); + dim3 grid(cute::ceil_div(N, Spec::kBlockN), cute::ceil_div(M, Spec::kBlockM)); constexpr int kShmSizeA = Spec::kShmSizeA; constexpr int kShmSizeB = Spec::kShmSizeB; diff --git a/10-gemm-api/gemm_api.cu b/10-gemm-api/gemm_api.cu index fa4f80a..bc875f9 100644 --- a/10-gemm-api/gemm_api.cu +++ b/10-gemm-api/gemm_api.cu @@ -32,9 +32,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -42,9 +42,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int G2S_Stages = G2S_Stages_; static_assert(G2S_Stages >= 2, "G2S_Stages should not be less than 2."); @@ -87,14 +87,7 @@ struct KernelSpec { using TiledMMA = decltype(make_tiled_mma(MMA_op{}, MMAThrLayout{}, MMATileLayout{})); static constexpr int kThreadNum = size(TiledMMA{}); - static constexpr int kThreadsPerWarp = 32; - static constexpr int kTileM_Copy = cute::min(kThreadsPerWarp, kTileM); - static constexpr int kTileN_Copy = cute::min(kThreadsPerWarp, kTileN); - - static constexpr int kAlignedCopyItemsA = - cute::min(128 / 8 / sizeof(ComputeTypeA), kTileK *kTileM_Copy / kThreadNum); - static constexpr int kAlignedCopyItemsB = - cute::min(128 / 8 / sizeof(ComputeTypeB), kTileK *kTileN_Copy / kThreadNum); + static constexpr int kBlockK_Copy = cute::min(64, kBlockK) / 8; using Copy_G2S_op = SM80_CP_ASYNC_CACHEGLOBAL; using CopyA_G2S_atom = Copy_Atom; @@ -102,14 +95,14 @@ struct KernelSpec { using TiledCopyA_G2S = decltype(make_tiled_copy(CopyA_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using TiledCopyB_G2S = decltype(make_tiled_copy(CopyB_G2S_atom{}, - make_layout(make_shape(Int{}, Int{}), - make_stride(Int{}, Int<1>{})), - make_layout(make_shape(Int<1>{}, Int{})))); + make_layout(make_shape(Int{}, Int{}), + make_stride(Int{}, Int<1>{})), + make_layout(make_shape(Int<1>{}, Int<8>{})))); using Copy_S2R_op_A = std::conditional_t; using Copy_S2R_op_B = std::conditional_t; @@ -118,8 +111,8 @@ struct KernelSpec { using CopyB_S2R_atom = Copy_Atom; using SmemLayoutAtomA = decltype(composition(Swizzle<3, 3, 3>{}, - make_layout(make_shape(Int<8>{}, Int{}), - make_stride(Int{}, Int<1>{})))); + make_layout(make_shape(Int<8>{}, Int{}), + make_stride(Int{}, Int<1>{})))); using SmemLayoutAtomB = SmemLayoutAtomA; ////////////////////////////////////////////////////////////////////////////////// @@ -150,7 +143,7 @@ struct KernelSpec { using ElementCompute = AccType_; using ArchTag = cutlass::arch::Sm80; using OperatorClass = cutlass::arch::OpClassTensorOp; - using TileShape = Shape, Int, Int>; + using TileShape = Shape, Int, Int>; using DispatchPolicy = cutlass::gemm::MainloopSm80CpAsync; @@ -289,9 +282,9 @@ template constexpr bool needs_precisio return !std::is_same_v; } -template ; + spec::KernelSpec; void *out_ptr = IsCvtPrecision ? out.data_ptr() : c.data_ptr(); diff --git a/10-gemm-api/gemm_api_sm90a.cu b/10-gemm-api/gemm_api_sm90a.cu index 9a17683..d387481 100644 --- a/10-gemm-api/gemm_api_sm90a.cu +++ b/10-gemm-api/gemm_api_sm90a.cu @@ -31,9 +31,9 @@ template + int kBlockM_, + int kBlockN_, + int kBlockK_> struct KernelSpec { // A matrix configuration using ElementA = ComputeTypeA_; @@ -60,7 +60,7 @@ struct KernelSpec { using ElementCompute = AccType_; using ArchTag = cutlass::arch::Sm90; using OperatorClass = cutlass::arch::OpClassTensorOp; - using TileShape = Shape, Int, Int>; + using TileShape = Shape, Int, Int>; using ClusterShape = Shape<_2, _1, _1>; using StageCount = cutlass::gemm::collective::StageCountAuto; using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto; @@ -219,9 +219,9 @@ template constexpr bool needs_precisio return !std::is_same_v; } -template ; + using Spec = spec::KernelSpec; void *out_ptr = IsCvtPrecision ? out.data_ptr() : c.data_ptr(); diff --git a/11-tma-load-store/tma_load_store.cu b/11-tma-load-store/tma_load_store.cu index 119a267..1dcdff8 100644 --- a/11-tma-load-store/tma_load_store.cu +++ b/11-tma-load-store/tma_load_store.cu @@ -61,9 +61,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void tma_load_store(__grid_consta using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutD = typename Spec::SmemLayoutD; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; extern __shared__ __align__(1024) uint8_t smem_raw[]; @@ -92,7 +92,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void tma_load_store(__grid_consta Tensor mC = tma_C.get_tma_tensor(make_shape(M, N)); Tensor mD = tma_D.get_tma_tensor(make_shape(M, N)); - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -115,7 +115,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void tma_load_store(__grid_consta group_modes<0, 2>(gD)); // (TMA,K_TILES) and (TMA) constexpr int tma_transaction_bytes = - kTileM * kTileK * sizeof(ComputeTypeA) + kTileN * kTileK * sizeof(ComputeTypeB); + kBlockM * kBlockK * sizeof(ComputeTypeA) + kBlockN * kBlockK * sizeof(ComputeTypeB); typename Spec::TiledMMA tiled_mma; ThrMMA thr_mma = tiled_mma.get_slice(tid); @@ -140,7 +140,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void tma_load_store(__grid_consta // MAINLOOP // - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); int warp_idx = cutlass::canonical_warp_idx_sync(); int lane_predicate = cute::elect_one_sync(); @@ -186,7 +186,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void tma_load_store(__grid_consta if constexpr (!IsGemm) { // Load C matrix with TMA - constexpr int tma_transaction_load_c_bytes = kTileM * kTileN * sizeof(ComputeTypeC); + constexpr int tma_transaction_load_c_bytes = kBlockM * kBlockN * sizeof(ComputeTypeC); using SharedStorage_C = SharedStorage_C; SharedStorage_C &smem_c = *reinterpret_cast(smem_raw); uint64_t &tma_load_c_mbarrier = smem_c.tma_barrier[0]; @@ -251,18 +251,18 @@ template + int kBlockM_ = 128, + int kBlockN_ = 128, + int kBlockK_ = 64> struct KernelSpec { using OutType = OutType_; using ComputeTypeA = ComputeTypeA_; using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; using MMA_op = std::conditional_t< std::is_same_v && std::is_same_v && @@ -326,22 +326,22 @@ struct KernelSpec { // Before Layout A: Sw<3,4,3> o _0 o ((_8,_16),(_64,_1)):((_64,_512),(_1,_0)) // After Layout A: Sw<3,4,3> o smem_ptr[16b](unset) o ((_8,_16),(_64,_1)):((_64,_512),(_1,_0)) using SmemLayoutA = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); // Also we can use: // using SmemLayoutAtomA = decltype(composition( // Swizzle<3, 4, 3>{}, // smem_ptr_flag_bits::value>{}, - // make_layout(make_shape(Int<8>{}, Int{}), - // make_stride(Int{}, Int<1>{})))); + // make_layout(make_shape(Int<8>{}, Int{}), + // make_stride(Int{}, Int<1>{})))); // using SmemLayoutA = decltype(tile_to_shape(SmemLayoutAtomA{}, - // make_shape(Int{}, Int{}), + // make_shape(Int{}, Int{}), // Step<_1,_2>{})); using SmemLayoutB = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutC = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutD = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -400,9 +400,9 @@ template constexpr torch::ScalarType to_torch_scalar_type() { throw std::runtime_error("Unsupported type!"); } -template ; + using Spec = spec::KernelSpec; dim3 block = Spec::kThreadNum; - dim3 grid(cute::ceil_div(N, Spec::kTileN), cute::ceil_div(M, Spec::kTileM)); + dim3 grid(cute::ceil_div(N, Spec::kBlockN), cute::ceil_div(M, Spec::kBlockM)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d Bytes\n", block.x, block.y, @@ -472,7 +472,7 @@ torch::Tensor run_tma_load_store(const torch::Tensor a, const torch::Tensor b, s auto smem_layout_D = typename Spec::SmemLayoutD{}; auto tma_A = make_tma_copy(SM90_TMA_LOAD{}, mA, smem_layout_A); - // auto tma_A = make_tma_atom(SM90_TMA_LOAD{}, mA, smem_layout_A, make_shape(kTileM,kTileK)); + // auto tma_A = make_tma_atom(SM90_TMA_LOAD{}, mA, smem_layout_A, make_shape(kBlockM,kBlockK)); auto tma_B = make_tma_copy(SM90_TMA_LOAD{}, mB, smem_layout_B); auto tma_C = make_tma_copy(SM90_TMA_LOAD{}, mC, smem_layout_C); auto tma_D = make_tma_copy(SM90_TMA_STORE{}, mD, smem_layout_D); diff --git a/12-tma-multicast-reduce/tma_multicast_reduce.cu b/12-tma-multicast-reduce/tma_multicast_reduce.cu index 23ba542..8a4c58d 100644 --- a/12-tma-multicast-reduce/tma_multicast_reduce.cu +++ b/12-tma-multicast-reduce/tma_multicast_reduce.cu @@ -74,9 +74,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutD = typename Spec::SmemLayoutD; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; extern __shared__ __align__(1024) uint8_t smem_raw[]; @@ -105,7 +105,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void Tensor mC = tma_C.get_tma_tensor(make_shape(M, N)); Tensor mD = tma_D.get_tma_tensor(make_shape(M, N)); - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -141,7 +141,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void group_modes<0, 2>(gD)); // (TMA,K_TILES) and (TMA) constexpr int tma_transaction_bytes = - kTileM * kTileK * sizeof(ComputeTypeA) + kTileN * kTileK * sizeof(ComputeTypeB); + kBlockM * kBlockK * sizeof(ComputeTypeA) + kBlockN * kBlockK * sizeof(ComputeTypeB); typename Spec::TiledMMA tiled_mma; ThrMMA thr_mma = tiled_mma.get_slice(tid); @@ -166,7 +166,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void // MAINLOOP // - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); int warp_idx = cutlass::canonical_warp_idx_sync(); int lane_predicate = cute::elect_one_sync(); @@ -242,7 +242,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void if constexpr (!IsGemm) { // Load C matrix with TMA - constexpr int tma_transaction_load_c_bytes = kTileM * kTileN * sizeof(ComputeTypeC); + constexpr int tma_transaction_load_c_bytes = kBlockM * kBlockN * sizeof(ComputeTypeC); using SharedStorage_C = SharedStorage_C; SharedStorage_C &smem_c = *reinterpret_cast(smem_raw); uint64_t &tma_load_c_mbarrier = smem_c.tma_barrier[0]; @@ -308,18 +308,18 @@ template + int kBlockM_ = 128, + int kBlockN_ = 128, + int kBlockK_ = 64> struct KernelSpec { using OutType = OutType_; using ComputeTypeA = ComputeTypeA_; using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; using MMA_op = std::conditional_t< std::is_same_v && std::is_same_v && @@ -382,13 +382,13 @@ struct KernelSpec { using TiledCopyD_R2S = decltype(make_tiled_copy_C(CopyD_R2S_atom{}, TiledMMA{})); using SmemLayoutA = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutB = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutC = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); using SmemLayoutD = decltype(tile_to_shape( - GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); + GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}), Step<_1, _2>{})); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -454,9 +454,9 @@ template constexpr bool if_use_tma_sto return false; } -template ; + using Spec = spec::KernelSpec; auto cluster_shape = typename Spec::ClusterShape{}; auto [cls_x, cls_y, cls_z] = cluster_shape; dim3 block = Spec::kThreadNum; dim3 cluster(cls_x, cls_y, cls_z); - dim3 grid(cute::ceil_div(N, Spec::kTileN), cute::ceil_div(M, Spec::kTileM)); + dim3 grid(cute::ceil_div(N, Spec::kBlockN), cute::ceil_div(M, Spec::kBlockM)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Cluster Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d " @@ -537,7 +537,7 @@ torch::Tensor run_tma_multicast_reduce(const torch::Tensor a, const torch::Tenso using TMA_Load_CopyB = std::conditional_t(cluster_shape) == 1, SM90_TMA_LOAD, SM90_TMA_LOAD_MULTICAST>; auto tma_A = make_tma_copy(TMA_Load_CopyA{}, mA, smem_layout_A, Int(cluster_shape)>{}); - // auto tma_A = make_tma_atom(TMA_Load_CopyA{}, mA, smem_layout_A, make_shape(kTileM, kTileK), + // auto tma_A = make_tma_atom(TMA_Load_CopyA{}, mA, smem_layout_A, make_shape(kBlockM, kBlockK), // Int(cluster_shape)>{}); auto tma_B = make_tma_copy(TMA_Load_CopyB{}, mB, smem_layout_B, Int(cluster_shape)>{}); diff --git a/13-warpgroup-mma/warpgroup_mma.cu b/13-warpgroup-mma/warpgroup_mma.cu index 35c3ae4..87e0ec2 100644 --- a/13-warpgroup-mma/warpgroup_mma.cu +++ b/13-warpgroup-mma/warpgroup_mma.cu @@ -75,9 +75,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutD = typename Spec::SmemLayoutD; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kStages = Spec::kStages; constexpr int kThreadNum = Spec::kThreadNum; constexpr int kWarpgroupNum = Spec::kWarpgroupNum; @@ -102,7 +102,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void Tensor mC = tma_C.get_tma_tensor(make_shape(M, N)); Tensor mD = tma_D.get_tma_tensor(make_shape(M, N)); - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -138,7 +138,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void group_modes<0, 2>(gD)); // (TMA,K_TILES) and (TMA) constexpr int tma_transaction_bytes = - kTileM * kTileK * sizeof(ComputeTypeA) + kTileN * kTileK * sizeof(ComputeTypeB); + kBlockM * kBlockK * sizeof(ComputeTypeA) + kBlockN * kBlockK * sizeof(ComputeTypeB); typename Spec::TiledMMA tiled_mma; ThrMMA thr_mma = tiled_mma.get_slice(tid); @@ -153,7 +153,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void // PREFETCH // - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); int warpgroup_idx = cutlass::canonical_warp_group_idx(); int warpgroup_tid = tid % cutlass::NumThreadsPerWarpGroup; @@ -279,7 +279,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void if constexpr (!IsGemm) { // Load C matrix with TMA - constexpr int tma_transaction_load_c_bytes = kTileM * kTileN * sizeof(ComputeTypeC); + constexpr int tma_transaction_load_c_bytes = kBlockM * kBlockN * sizeof(ComputeTypeC); using SharedStorage_C = SharedStorage_C; SharedStorage_C &smem_c = *reinterpret_cast(smem_raw); uint64_t &tma_load_c_mbarrier = smem_c.tma_barrier[0]; @@ -345,9 +345,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -355,9 +355,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int kStages = kStages_; // TN => K Major for both A & B @@ -391,7 +391,7 @@ struct KernelSpec { using MMA_op = decltype(cute::GMMA::ss_op_selector, Int, Int>, + Shape, Int, Int>, GmmaMajorA, GmmaMajorB>()); @@ -435,13 +435,13 @@ struct KernelSpec { using TiledCopyD_R2S = decltype(make_tiled_copy_C(CopyD_R2S_atom{}, TiledMMA{})); using SmemLayoutA = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutB = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutC = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); using SmemLayoutD = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -507,9 +507,9 @@ template constexpr bool if_use_tma_sto return false; } -template ; + using Spec = spec::KernelSpec; auto cluster_shape = typename Spec::ClusterShape{}; @@ -561,7 +561,7 @@ torch::Tensor run_warpgroup_mma(const torch::Tensor a, const torch::Tensor b, st dim3 block = Spec::kThreadNum; dim3 cluster(cls_x, cls_y, cls_z); - dim3 grid(cute::ceil_div(N, Spec::kTileN), cute::ceil_div(M, Spec::kTileM)); + dim3 grid(cute::ceil_div(N, Spec::kBlockN), cute::ceil_div(M, Spec::kBlockM)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Cluster Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d " diff --git a/13-warpgroup-mma/warpgroup_mma_multicast.cu b/13-warpgroup-mma/warpgroup_mma_multicast.cu index ecc2068..ae94ea7 100644 --- a/13-warpgroup-mma/warpgroup_mma_multicast.cu +++ b/13-warpgroup-mma/warpgroup_mma_multicast.cu @@ -90,9 +90,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutD = typename Spec::SmemLayoutD; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kStages = Spec::kStages; constexpr int kThreadNum = Spec::kThreadNum; constexpr int kWarpgroupNum = Spec::kWarpgroupNum; @@ -117,7 +117,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void Tensor mC = tma_C.get_tma_tensor(make_shape(M, N)); Tensor mD = tma_D.get_tma_tensor(make_shape(M, N)); - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -153,7 +153,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void group_modes<0, 2>(gD)); // (TMA,K_TILES) and (TMA) constexpr int tma_transaction_bytes = - kTileM * kTileK * sizeof(ComputeTypeA) + kTileN * kTileK * sizeof(ComputeTypeB); + kBlockM * kBlockK * sizeof(ComputeTypeA) + kBlockN * kBlockK * sizeof(ComputeTypeB); typename Spec::TiledMMA tiled_mma; ThrMMA thr_mma = tiled_mma.get_slice(tid); @@ -166,7 +166,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void // PREFETCH // - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); int warpgroup_idx = cutlass::canonical_warp_group_idx(); int warpgroup_tid = tid % cutlass::NumThreadsPerWarpGroup; @@ -307,7 +307,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void if constexpr (!IsGemm) { // Load C matrix with TMA - constexpr int tma_transaction_load_c_bytes = kTileM * kTileN * sizeof(ComputeTypeC); + constexpr int tma_transaction_load_c_bytes = kBlockM * kBlockN * sizeof(ComputeTypeC); using SharedStorage_C = SharedStorage_C; SharedStorage_C &smem_c = *reinterpret_cast(smem_raw); uint64_t &tma_load_c_mbarrier = smem_c.tma_barrier[0]; @@ -378,9 +378,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -388,9 +388,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int kStages = kStages_; // TN => K Major for both A & B @@ -424,7 +424,7 @@ struct KernelSpec { using MMA_op = decltype(cute::GMMA::ss_op_selector, Int, Int>, + Shape, Int, Int>, GmmaMajorA, GmmaMajorB>()); @@ -468,13 +468,13 @@ struct KernelSpec { using TiledCopyD_R2S = decltype(make_tiled_copy_C(CopyD_R2S_atom{}, TiledMMA{})); using SmemLayoutA = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutB = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutC = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); using SmemLayoutD = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -540,9 +540,9 @@ template constexpr bool if_use_tma_sto return false; } -template ; + using Spec = spec::KernelSpec; auto cluster_shape = typename Spec::ClusterShape{}; @@ -594,8 +594,8 @@ torch::Tensor run_warpgroup_mma(const torch::Tensor a, const torch::Tensor b, st dim3 block = Spec::kThreadNum; dim3 cluster(cls_x, cls_y, cls_z); - dim3 grid(cute::max(cute::ceil_div(N, Spec::kTileN), Spec::kClusterDimX), - cute::max(cute::ceil_div(M, Spec::kTileM), Spec::kClusterDimY)); + dim3 grid(cute::max(cute::ceil_div(N, Spec::kBlockN), Spec::kClusterDimX), + cute::max(cute::ceil_div(M, Spec::kBlockM), Spec::kClusterDimY)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Cluster Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d " diff --git a/14-warp-specialization/warp_specialization.cu b/14-warp-specialization/warp_specialization.cu index a067e98..e485aa7 100644 --- a/14-warp-specialization/warp_specialization.cu +++ b/14-warp-specialization/warp_specialization.cu @@ -98,9 +98,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutD = typename Spec::SmemLayoutD; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kStages = Spec::kStages; constexpr int kClusterSize = Spec::kClusterSize; constexpr int kProducerRegisters = Spec::kProducerRegisters; @@ -110,7 +110,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c int tid = threadIdx.x; int bidx = blockIdx.x; int bidy = blockIdx.y; - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); int warpgroup_idx = cutlass::canonical_warp_group_idx(); int thread_idx_in_warpgroup = tid % cutlass::NumThreadsPerWarpGroup; @@ -154,7 +154,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c Tensor mC = tma_C.get_tma_tensor(make_shape(M, N)); Tensor mD = tma_D.get_tma_tensor(make_shape(M, N)); - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -222,7 +222,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c group_modes<0, 2>(sB), group_modes<0, 2>(gB)); // (TMA,K_TILES) and (TMA,kStages) constexpr int tma_transaction_bytes = - kTileM * kTileK * sizeof(ComputeTypeA) + kTileN * kTileK * sizeof(ComputeTypeB); + kBlockM * kBlockK * sizeof(ComputeTypeA) + kBlockN * kBlockK * sizeof(ComputeTypeB); int smem_write_pipe = 0; int mma_phase_bit = 1; @@ -340,7 +340,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c if constexpr (!IsGemm) { // Load C matrix with TMA - constexpr int tma_transaction_load_c_bytes = kTileM * kTileN * sizeof(ComputeTypeC); + constexpr int tma_transaction_load_c_bytes = kBlockM * kBlockN * sizeof(ComputeTypeC); using SharedStorage_C = SharedStorage_C; SharedStorage_C &smem_c = *reinterpret_cast(smem_raw); uint64_t &tma_load_c_mbarrier = smem_c.tma_barrier[0]; @@ -412,9 +412,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -422,9 +422,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int kStages = kStages_; static constexpr cute::GMMA::Major GmmaMajorA = cute::GMMA::Major::K; @@ -433,7 +433,7 @@ struct KernelSpec { using MMA_op = decltype(cute::GMMA::ss_op_selector, Int, Int>, + Shape, Int, Int>, GmmaMajorA, GmmaMajorB>()); @@ -484,13 +484,13 @@ struct KernelSpec { using TiledCopyD_R2S = decltype(make_tiled_copy_C(CopyD_R2S_atom{}, TiledMMA{})); using SmemLayoutA = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutB = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutC = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); using SmemLayoutD = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -556,9 +556,9 @@ template constexpr bool if_use_tma_sto return false; } -template ; + using Spec = spec::KernelSpec; auto cluster_shape = typename Spec::ClusterShape{}; @@ -610,8 +610,8 @@ torch::Tensor run_warp_specialization(const torch::Tensor a, const torch::Tensor dim3 block = Spec::kThreadNum; dim3 cluster(cls_x, cls_y, cls_z); - dim3 grid(cute::max(cute::ceil_div(N, Spec::kTileN), Spec::kClusterDimX), - cute::max(cute::ceil_div(M, Spec::kTileM), Spec::kClusterDimY)); + dim3 grid(cute::max(cute::ceil_div(N, Spec::kBlockN), Spec::kClusterDimX), + cute::max(cute::ceil_div(M, Spec::kBlockM), Spec::kClusterDimY)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Cluster Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d " diff --git a/14-warp-specialization/warp_specialization_pipeline_api.cu b/14-warp-specialization/warp_specialization_pipeline_api.cu index 3c13cdb..c85aff2 100644 --- a/14-warp-specialization/warp_specialization_pipeline_api.cu +++ b/14-warp-specialization/warp_specialization_pipeline_api.cu @@ -84,9 +84,9 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c using SmemLayoutC = typename Spec::SmemLayoutC; using SmemLayoutD = typename Spec::SmemLayoutD; - constexpr int kTileM = Spec::kTileM; - constexpr int kTileN = Spec::kTileN; - constexpr int kTileK = Spec::kTileK; + constexpr int kBlockM = Spec::kBlockM; + constexpr int kBlockN = Spec::kBlockN; + constexpr int kBlockK = Spec::kBlockK; constexpr int kStages = Spec::kStages; constexpr int kClusterSize = Spec::kClusterSize; constexpr int kProducerRegisters = Spec::kProducerRegisters; @@ -96,7 +96,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c int tid = threadIdx.x; int bidx = blockIdx.x; int bidy = blockIdx.y; - int NTilesK = ceil_div(K, kTileK); + int NTilesK = ceil_div(K, kBlockK); int warpgroup_idx = cutlass::canonical_warp_group_idx(); int thread_idx_in_warpgroup = tid % cutlass::NumThreadsPerWarpGroup; @@ -137,7 +137,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c Tensor mC = tma_C.get_tma_tensor(make_shape(M, N)); Tensor mD = tma_D.get_tma_tensor(make_shape(M, N)); - auto tiler = make_tile(Int{}, Int{}, Int{}); + auto tiler = make_tile(Int{}, Int{}, Int{}); auto coord = make_coord(bidy, bidx, _); Tensor gA = local_tile(mA, tiler, coord, Step<_1, X, _1>{}); // (BLK_M, BLK_K, K_TILES) @@ -151,7 +151,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c Tensor sD = make_tensor(make_smem_ptr((OutType *)Dptr_smem), SmemLayoutD{}); // (BLK_M, BLK_N) constexpr int tma_transaction_bytes = - kTileM * kTileK * sizeof(ComputeTypeA) + kTileN * kTileK * sizeof(ComputeTypeB); + kBlockM * kBlockK * sizeof(ComputeTypeA) + kBlockN * kBlockK * sizeof(ComputeTypeB); enum class WarpGroupRole { Producer = 0, Consumer0 = 1, Consumer1 = 2 }; @@ -294,7 +294,7 @@ __global__ __launch_bounds__(Spec::kThreadNum) void warp_specialization(__grid_c if constexpr (!IsGemm) { // Load C matrix with TMA - constexpr int tma_transaction_load_c_bytes = kTileM * kTileN * sizeof(ComputeTypeC); + constexpr int tma_transaction_load_c_bytes = kBlockM * kBlockN * sizeof(ComputeTypeC); using SharedStorage_C = SharedStorage_C; SharedStorage_C &smem_c = *reinterpret_cast(smem_raw); uint64_t &tma_load_c_mbarrier = smem_c.tma_barrier[0]; @@ -366,9 +366,9 @@ template struct KernelSpec { using OutType = OutType_; @@ -376,9 +376,9 @@ struct KernelSpec { using ComputeTypeB = ComputeTypeB_; using ComputeTypeC = ComputeTypeC_; - static constexpr int kTileM = kTileM_; - static constexpr int kTileN = kTileN_; - static constexpr int kTileK = kTileK_; + static constexpr int kBlockM = kBlockM_; + static constexpr int kBlockN = kBlockN_; + static constexpr int kBlockK = kBlockK_; static constexpr int kStages = kStages_; static constexpr cute::GMMA::Major GmmaMajorA = cute::GMMA::Major::K; @@ -387,7 +387,7 @@ struct KernelSpec { using MMA_op = decltype(cute::GMMA::ss_op_selector, Int, Int>, + Shape, Int, Int>, GmmaMajorA, GmmaMajorB>()); @@ -438,13 +438,13 @@ struct KernelSpec { using TiledCopyD_R2S = decltype(make_tiled_copy_C(CopyD_R2S_atom{}, TiledMMA{})); using SmemLayoutA = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutB = decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, - make_shape(Int{}, Int{}, Int{}))); + make_shape(Int{}, Int{}, Int{}))); using SmemLayoutC = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); using SmemLayoutD = - decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); + decltype(tile_to_shape(GMMA::Layout_K_SW128_Atom{}, make_shape(Int{}, Int{}))); static constexpr int kShmSizeA = cosize_v * sizeof(ComputeTypeA); static constexpr int kShmSizeB = cosize_v * sizeof(ComputeTypeB); @@ -510,9 +510,9 @@ template constexpr bool if_use_tma_sto return false; } -template ; + using Spec = spec::KernelSpec; auto cluster_shape = typename Spec::ClusterShape{}; @@ -564,8 +564,8 @@ torch::Tensor run_warp_specialization(const torch::Tensor a, const torch::Tensor dim3 block = Spec::kThreadNum; dim3 cluster(cls_x, cls_y, cls_z); - dim3 grid(cute::max(cute::ceil_div(N, Spec::kTileN), Spec::kClusterDimX), - cute::max(cute::ceil_div(M, Spec::kTileM), Spec::kClusterDimY)); + dim3 grid(cute::max(cute::ceil_div(N, Spec::kBlockN), Spec::kClusterDimX), + cute::max(cute::ceil_div(M, Spec::kBlockM), Spec::kClusterDimY)); int shm_size = Spec::kShmSize; printf("Block Size: (%d, %d, %d) | Cluster Size: (%d, %d, %d) | Grid Size: (%d, %d, %d) | Shared Memory Size: %d "