Merge pull request #600 from fsx950223/grouped_api

add grouped flash attention api

example/32_batched_gemm_scale_softmax_gemm/CMakeLists.txt

@@ -9,6 +9,7 @@ add_example_executable(example_grouped_multihead_attention_forward grouped_multi
 add_example_executable(example_batched_multihead_attention_forward batched_multihead_attention_forward.cpp)
 add_example_executable(example_batched_multihead_attention_backward batched_multihead_attention_backward.cpp)
 add_example_executable(example_batched_multihead_attention_train batched_multihead_attention_train.cpp)
+add_example_executable(example_grouped_multihead_attention_backward_fp16 grouped_multihead_attention_backward_fp16.cpp)
 
 add_custom_target(example_gemm_scale_softmax_gemm)
 add_dependencies(example_gemm_scale_softmax_gemm example_batched_gemm_scale_softmax_gemm_xdl_fp16)

example/32_batched_gemm_scale_softmax_gemm/grouped_multihead_attention_backward_fp16.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+/*
+Backprop for Gemm + Softmax + Gemm fused operation, where forward prop is defined as:
+
+  Y_g_m_o = Softmax(alpha * Q_g_m_k * K_g_k_n) * V_g_n_o
+
+Computation graph:
+
+          K^T                   V
+          |                     |
+          |                     |
+    Q --- * ----- Softmax ----- * --> Y
+              S             P
+
+Kernel inputs:
+
+    Q, K, V, Y, dY, per-row softmax stats (LSE)
+
+Kernel outputs:
+
+    dQ, dK, dV
+
+*/
+
+#define PRINT_HOST 0
+#define USING_MASK 0
+#define DIM 64
+
+#include <iostream>
+#include <numeric>
+#include <initializer_list>
+#include <cstdlib>
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_specialization.hpp"
+#include "ck/tensor_operation/gpu/device/impl/device_grouped_multihead_attention_backward_xdl_cshuffle_v2.hpp"
+#include "ck/tensor_operation/gpu/device/impl/device_grouped_multihead_attention_backward_xdl_cshuffle_pt1.hpp"
+#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
+
+#include "ck/library/utility/check_err.hpp"
+#include "ck/library/utility/device_memory.hpp"
+#include "ck/library/utility/host_tensor.hpp"
+#include "ck/library/utility/host_tensor_generator.hpp"
+#include "ck/library/reference_tensor_operation/cpu/reference_batched_gemm.hpp"
+#include "ck/library/reference_tensor_operation/cpu/reference_softmax.hpp"
+#include "ck/library/reference_tensor_operation/cpu/reference_dropout.hpp"
+
+
+template <ck::index_t... Is>
+using S = ck::Sequence<Is...>;
+
+using F16 = ck::half_t;
+using F32 = float;
+using U16 = unsigned short;
+
+using PassThrough = ck::tensor_operation::element_wise::PassThrough;
+using Scale       = ck::tensor_operation::element_wise::Scale;
+
+using QKVElementOp = PassThrough;
+using YElementOp   = PassThrough;
+
+using DataType         = F16;
+using GemmDataType     = F16;
+using AccDataType      = F32;
+using ShuffleDataType  = F32;
+using LSEDataType      = F32;
+using ZDataType        = U16;
+using Acc0BiasDataType = ck::Tuple<>;
+using Acc1BiasDataType = ck::Tuple<>;
+
+static constexpr ck::index_t NumDimG = 2;
+static constexpr ck::index_t NumDimM = 1;
+static constexpr ck::index_t NumDimN = 1;
+static constexpr ck::index_t NumDimK = 1;
+static constexpr ck::index_t NumDimO = 1;
+
+static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::MNKOPadding;
+#if USING_MASK
+static constexpr auto MaskingSpec =
+    ck::tensor_operation::device::MaskingSpecialization::MaskOutUpperTriangle;
+#else
+static constexpr auto MaskingSpec =
+    ck::tensor_operation::device::MaskingSpecialization::MaskDisabled;
+#endif
+
+static constexpr auto TensorSpecQ = ck::tensor_operation::device::TensorSpecialization::Default;
+static constexpr auto TensorSpecK = ck::tensor_operation::device::TensorSpecialization::Default;
+static constexpr auto TensorSpecV = ck::tensor_operation::device::TensorSpecialization::Default;
+static constexpr auto TensorSpecY = ck::tensor_operation::device::TensorSpecialization::Default;
+
+#if DIM >=128
+using DeviceGemmInstance = ck::tensor_operation::device::DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_V2<
+        NumDimG,
+        NumDimM,
+        NumDimN,
+        NumDimK,
+        NumDimO,
+        DataType,
+        GemmDataType,
+        ZDataType,
+        LSEDataType,
+        Acc0BiasDataType,
+        Acc1BiasDataType,
+        AccDataType,
+        ShuffleDataType,
+        QKVElementOp,
+        QKVElementOp,
+        Scale,
+        QKVElementOp,
+        YElementOp,
+        GemmSpec,
+        TensorSpecQ,
+        TensorSpecK,
+        TensorSpecV,
+        TensorSpecY,
+        1,
+        256,
+        128,         // MPerBlock
+        128,         // NPerBlock
+        64,          // KPerBlock
+        128,         // Gemm1NPerBlock
+        32,          // Gemm1KPerBlock
+        8,           // AK1
+        8,           // BK1
+        2,           // B1K1
+        32,          // MPerXDL
+        32,          // NPerXDL
+        1,           // MXdlPerWave
+        4,           // NXdlPerWave
+        4,           // Gemm1NXdlPerWave
+        2,           // Gemm2NXdlPerWave
+        S<4, 64, 1>, // ABlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<4, 64, 1>, // BBlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<8, 32, 1>, // B1BlockTransfer
+        S<0, 2, 1>,
+        S<0, 2, 1>,
+        1,
+        4,
+        2,
+        false,
+        1,              // CShuffleMXdlPerWavePerShuffle
+        4,              // CShuffleNXdlPerWavePerShuffle
+        S<1, 32, 1, 8>, // CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
+        8,              // CShuffleBlockTransferScalarPerVector_NPerBlock
+        MaskingSpec>;   // MaskingSpecialization
+#elif DIM >64
+using DeviceGemmInstance = ck::tensor_operation::device::DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_V2<
+        NumDimG,
+        NumDimM,
+        NumDimN,
+        NumDimK,
+        NumDimO,
+        DataType,
+        GemmDataType,
+        ZDataType,
+        LSEDataType,
+        Acc0BiasDataType,
+        Acc1BiasDataType,
+        AccDataType,
+        ShuffleDataType,
+        QKVElementOp,
+        QKVElementOp,
+        Scale,
+        QKVElementOp,
+        YElementOp,
+        GemmSpec,
+        TensorSpecQ,
+        TensorSpecK,
+        TensorSpecV,
+        TensorSpecY,
+        1,
+        256,
+        128,         // MPerBlock
+        128,         // NPerBlock
+        64,          // KPerBlock
+        64,          // Gemm1NPerBlock
+        64,          // Gemm1KPerBlock
+        8,           // AK1
+        8,           // BK1
+        2,           // B1K1
+        32,          // MPerXDL
+        32,          // NPerXDL
+        1,           // MXdlPerWave
+        4,           // NXdlPerWave
+        2,           // Gemm1NXdlPerWave
+        2,           // Gemm2NXdlPerWave
+        S<4, 64, 1>, // ABlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<4, 64, 1>, // BBlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<8, 32, 1>, // B1BlockTransfer
+        S<0, 2, 1>,
+        S<0, 2, 1>,
+        1,
+        2,
+        2,
+        false,
+        1,              // CShuffleMXdlPerWavePerShuffle
+        2,              // CShuffleNXdlPerWavePerShuffle
+        S<1, 32, 1, 8>, // CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
+        8,              // CShuffleBlockTransferScalarPerVector_NPerBlock
+        MaskingSpec>;   // MaskingSpecialization
+#elif DIM >32
+using DeviceGemmInstance = ck::tensor_operation::device::DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_PT1<
+        NumDimG,
+        NumDimM,
+        NumDimN,
+        NumDimK,
+        NumDimO,
+        DataType,
+        GemmDataType,
+        ZDataType,
+        LSEDataType,
+        Acc0BiasDataType,
+        Acc1BiasDataType,
+        AccDataType,
+        ShuffleDataType,
+        QKVElementOp,
+        QKVElementOp,
+        Scale,
+        QKVElementOp,
+        YElementOp,
+        GemmSpec,
+        TensorSpecQ,
+        TensorSpecK,
+        TensorSpecV,
+        TensorSpecY,
+        1,
+        256,
+        128,         // MPerBlock
+        128,         // NPerBlock
+        64,          // KPerBlock
+        64,          // Gemm1NPerBlock
+        32,          // Gemm1KPerBlock
+        8,           // AK1
+        8,           // BK1
+        2,           // B1K1
+        32,          // MPerXDL
+        32,          // NPerXDL
+        1,           // MXdlPerWave
+        4,           // NXdlPerWave
+        2,           // Gemm1NXdlPerWave
+        2,           // Gemm2NXdlPerWave
+        S<4, 64, 1>, // ABlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<4, 64, 1>, // BBlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<8, 32, 1>, // B1BlockTransfer
+        S<0, 2, 1>,
+        S<0, 2, 1>,
+        1,
+        4,
+        2,
+        false,
+        1,              // CShuffleMXdlPerWavePerShuffle
+        2,              // CShuffleNXdlPerWavePerShuffle
+        S<1, 32, 1, 8>, // CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
+        8,              // CShuffleBlockTransferScalarPerVector_NPerBlock
+        MaskingSpec>;   // MaskingSpecialization
+#else
+using DeviceGemmInstance = ck::tensor_operation::device::DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_PT1<
+        NumDimG,
+        NumDimM,
+        NumDimN,
+        NumDimK,
+        NumDimO,
+        DataType,
+        GemmDataType,
+        ZDataType,
+        LSEDataType,
+        Acc0BiasDataType,
+        Acc1BiasDataType,
+        AccDataType,
+        ShuffleDataType,
+        QKVElementOp,
+        QKVElementOp,
+        Scale,
+        QKVElementOp,
+        YElementOp,
+        GemmSpec,
+        TensorSpecQ,
+        TensorSpecK,
+        TensorSpecV,
+        TensorSpecY,
+        1,
+        256,
+        128,         // MPerBlock
+        128,         // NPerBlock
+        32,          // KPerBlock
+        32,          // Gemm1NPerBlock
+        32,          // Gemm1KPerBlock
+        8,           // AK1
+        8,           // BK1
+        2,           // B1K1
+        32,          // MPerXDL
+        32,          // NPerXDL
+        1,           // MXdlPerWave
+        4,           // NXdlPerWave
+        1,           // Gemm1NXdlPerWave
+        1,           // Gemm2NXdlPerWave
+        S<4, 64, 1>, // ABlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<4, 64, 1>, // BBlockTransfer
+        S<1, 0, 2>,
+        S<1, 0, 2>,
+        2,
+        8,
+        8,
+        true,
+        S<8, 32, 1>, // B1BlockTransfer
+        S<0, 2, 1>,
+        S<0, 2, 1>,
+        1,
+        4,
+        2,
+        false,
+        1,              // CShuffleMXdlPerWavePerShuffle
+        1,              // CShuffleNXdlPerWavePerShuffle
+        S<1, 64, 1, 4>, // CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
+        8,              // CShuffleBlockTransferScalarPerVector_NPerBlock
+        MaskingSpec>;   // MaskingSpecialization
+#endif
+// Ref Gemm0: S = alpha * Q * K^T
+// fp16 in, fp32 out
+using ReferenceGemm0Instance = ck::tensor_operation::host::ReferenceBatchedGemm<DataType,
+                                                                                DataType,
+                                                                                AccDataType,
+                                                                                AccDataType,
+                                                                                PassThrough,
+                                                                                PassThrough,
+                                                                                Scale>;
+
+// Ref Softmax: P = Softmax(S)
+// fp32 in, fp16 out
+using ReferenceSoftmaxInstance =
+    ck::tensor_operation::host::ReferenceSoftmax<AccDataType, DataType, AccDataType>;
+
+// Ref Gemm1: Y = P * V
+// fp16 in, fp16 out
+using ReferenceGemm1Instance = ck::tensor_operation::host::ReferenceBatchedGemm<DataType,
+                                                                                DataType,
+                                                                                DataType,
+                                                                                AccDataType,
+                                                                                PassThrough,
+                                                                                PassThrough,
+                                                                                PassThrough>;
+
+// Ref Gemm for backward pass
+// fp16 in, fp16 out
+using ReferenceGemmGradInstance = ck::tensor_operation::host::ReferenceBatchedGemm<DataType,
+                                                                                   DataType,
+                                                                                   DataType,
+                                                                                   AccDataType,
+                                                                                   PassThrough,
+                                                                                   PassThrough,
+                                                                                   Scale>;
+
+using ReferenceDropoutInstance =
+    ck::tensor_operation::host::ReferenceDropout<ushort, DataType, DataType>;
+
+template <typename TensorQ,
+          typename TensorK,
+          typename TensorV,
+          typename TensorS,
+          typename TensorP,
+          typename TensorZ,
+          typename TensorY,
+          typename TensorLSE = TensorP>
+void run_attention_fwd_host(const TensorQ& q_g_m_k,
+                            const TensorK& k_g_n_k,
+                            const TensorV& v_g_n_o,
+                            const float alpha,
+                            TensorS& s_g_m_n,
+                            TensorP& p_g_m_n,
+                            TensorY& y_g_m_o,
+                            TensorLSE& lse_g_m,
+                            TensorP& p_drop_g_m_n,
+                            TensorZ& z_g_m_n,
+                            ushort p_dropout_in_16bits,
+                            float rp_dropout)
+{
+    // S = alpha * Q * K^T
+    auto k_g_k_n            = k_g_n_k.Transpose({0, 2, 1});
+    auto ref_gemm0          = ReferenceGemm0Instance{};
+    auto ref_gemm0_invoker  = ref_gemm0.MakeInvoker();
+    auto ref_gemm0_argument = ref_gemm0.MakeArgument(
+        q_g_m_k, k_g_k_n, s_g_m_n, PassThrough{}, PassThrough{}, Scale{alpha});
+
+    ref_gemm0_invoker.Run(ref_gemm0_argument);
+
+    // masking
+#if USING_MASK
+    auto N          = s_g_m_n.GetLengths()[2];
+    const auto mask = DeviceGemmInstance::C0MatrixMask(N);
+    s_g_m_n.ForEach([&](auto& self, auto idx) {
+        if(mask.IsMaskedElement(idx[1], idx[2]))
+            self(idx) = -ck::NumericLimits<float>::Infinity();
+    });
+#endif
+
+    // P = Softmax(S)
+    auto ref_softmax          = ReferenceSoftmaxInstance{};
+    auto ref_softmax_invoker  = ref_softmax.MakeInvoker();
+    auto ref_softmax_argument = ref_softmax.MakeArgument(s_g_m_n, p_g_m_n, 1, 0, {2}, &lse_g_m);
+
+    ref_softmax_invoker.Run(ref_softmax_argument);
+
+    // P_dropped
+    auto ref_dropout         = ReferenceDropoutInstance{};
+    auto ref_dropout_invoker = ref_dropout.MakeInvoker();
+    auto ref_dropout_argment =
+        ref_dropout.MakeArgument(z_g_m_n, p_g_m_n, p_drop_g_m_n, p_dropout_in_16bits, rp_dropout);
+    ref_dropout_invoker.Run(ref_dropout_argment);
+
+    // Y = P_dropout * V
+    auto ref_gemm1          = ReferenceGemm1Instance{};
+    auto ref_gemm1_invoker  = ref_gemm1.MakeInvoker();
+    auto ref_gemm1_argument = ref_gemm1.MakeArgument(
+        p_drop_g_m_n, v_g_n_o, y_g_m_o, PassThrough{}, PassThrough{}, PassThrough{});
+
+    ref_gemm1_invoker.Run(ref_gemm1_argument);
+}
+
+int run(int argc, char* argv[])
+{
+    bool do_verification = true;
+    int init_method      = 2; // method 1 will have slightly higher error; TODO: to investigate
+    bool time_kernel     = true;
+
+    // Overall QKV matrices shape
+    // y_g_m_o = Softmax(alpha * Q_g_m_k * K_g_k_n) * V_g_n_o
+    // y_g0_g1_m_o = reshape(y_g_m_o, [G0, G1, M, O])
+    // y_g0_m_g1_o = permute(y_g0_g1_m_o, [0, 2, 1, 3])
+    float alpha                     = 1.f / std::sqrt(DIM);
+    float p_drop                    = 0.2;
+    float p_dropout                 = 1 - p_drop;
+    uint16_t p_dropout_in_16bits    = uint16_t(std::floor(p_dropout * 65535.0));
+    float rp_dropout                = 1.0 / p_dropout;
+    const unsigned long long seed   = 1;
+    const unsigned long long offset = 0;
+
+    bool input_permute  = false;
+    bool output_permute = true;
+
+    if(argc == 1)
+    {
+        // use default case
+    }
+    else if(argc == 4)
+    {
+        do_verification = std::stoi(argv[1]);
+        init_method     = std::stoi(argv[2]);
+        time_kernel     = std::stoi(argv[3]);
+    }
+    else if(argc == 7)
+    {
+        do_verification = std::stoi(argv[1]);
+        init_method     = std::stoi(argv[2]);
+        time_kernel     = std::stoi(argv[3]);
+
+        alpha = std::stof(argv[4]);
+
+        input_permute  = std::stoi(argv[5]);
+        output_permute = std::stoi(argv[6]);
+    }
+    else
+    {
+        printf("arg1: verification (0=no, 1=yes)\n");
+        printf("arg2: initialization (0=no init, 1=integer value, 2=decimal value)\n");
+        printf("arg3: time kernel (0=no, 1=yes)\n");
+        printf("arg4 to 11: M, N, K, O, G0, G1\n");
+        printf("arg10: scale (alpha)\n");
+        printf("arg11 to 12: input / output permute\n");
+        exit(0);
+    }
+
+    auto gemm    = DeviceGemmInstance{};
+    auto invoker = gemm.MakeInvoker();
+    std::vector<DeviceGemmInstance::ProblemDesc> problem_descs;
+
+    using DeviceMemPtr = std::unique_ptr<DeviceMem>;
+    std::vector<const void*> p_q;
+    std::vector<const void*> p_k;
+    std::vector<void*> p_z;
+    std::vector<const void*> p_v;
+    std::vector<const void*> p_y;
+    std::vector<const void*> p_lse;
+    std::vector<void*> p_qgrad;
+    std::vector<void*> p_kgrad;
+    std::vector<void*> p_vgrad;
+    std::vector<const void*> p_ygrad;
+
+    std::vector<Tensor<DataType>> q_g_m_ks;
+    std::vector<Tensor<DataType>> k_g_n_ks;
+    std::vector<Tensor<ZDataType>> z_g_m_ns;
+    std::vector<Tensor<DataType>> v_g_n_os;
+    std::vector<Tensor<AccDataType>> s_g_m_ns;
+    std::vector<Tensor<DataType>> p_g_m_ns;
+    std::vector<Tensor<DataType>> y_g_m_os;
+    std::vector<Tensor<LSEDataType>> lse_g_ms;
+    std::vector<Tensor<DataType>> p_drop_g_m_ns;
+
+    std::vector<Tensor<DataType>> q_tensors;
+    std::vector<Tensor<DataType>> k_tensors;
+    std::vector<Tensor<DataType>> v_tensors;
+    std::vector<Tensor<DataType>> y_tensors;
+    std::vector<Tensor<ZDataType>> z_tensors;
+    std::vector<Tensor<LSEDataType>> lse_tensors;
+    std::vector<Tensor<DataType>> qgrad_tensors;
+    std::vector<Tensor<DataType>> kgrad_tensors;
+    std::vector<Tensor<DataType>> vgrad_tensors;
+    std::vector<Tensor<DataType>> ygrad_tensors;
+
+    std::vector<DeviceMemPtr> q_tensors_device;
+    std::vector<DeviceMemPtr> k_tensors_device;
+    std::vector<DeviceMemPtr> z_tensors_device;
+    std::vector<DeviceMemPtr> v_tensors_device;
+    std::vector<DeviceMemPtr> y_tensors_device;
+    std::vector<DeviceMemPtr> lse_tensors_device;
+    std::vector<DeviceMemPtr> qgrad_tensors_device;
+    std::vector<DeviceMemPtr> ygrad_tensors_device;
+    std::vector<DeviceMemPtr> kgrad_tensors_device;
+    std::vector<DeviceMemPtr> vgrad_tensors_device;
+    std::size_t group_count = 3;
+    std::size_t flop = 0, num_byte = 0;
+    for(std::size_t i = 0; i < group_count; i++)
+    {
+        int M  = 128 * (rand() % 4 + 1);
+        int N  = 128 * (rand() % 4 + 1);
+        int K  = DIM;
+        int O  = DIM;
+        int G0 = rand() % 3 + 1;
+        int G1 = rand() % 2 + 1;
+        std::vector<ck::index_t> q_gs_ms_ks_lengths{G0, G1, M, K};
+        std::vector<ck::index_t> q_gs_ms_ks_strides =
+            input_permute
+                ? std::vector<ck::index_t>{M * G1 * K, K, G1 * K, 1} // Q layout [G0, M, G1, K]
+                : std::vector<ck::index_t>{G1 * M * K, M * K, K, 1}; // Q layout [G0, G1, M, K]
+
+        std::vector<ck::index_t> k_gs_ns_ks_lengths{G0, G1, N, K};
+        std::vector<ck::index_t> k_gs_ns_ks_strides =
+            input_permute
+                ? std::vector<ck::index_t>{N * G1 * K, K, G1 * K, 1} // K layout [G0, N, G1, K]
+                : std::vector<ck::index_t>{G1 * N * K, N * K, K, 1}; // K layout [G0, G1, N, K]
+
+        std::vector<ck::index_t> v_gs_os_ns_lengths{G0, G1, O, N};
+        std::vector<ck::index_t> v_gs_os_ns_strides =
+            input_permute
+                ? std::vector<ck::index_t>{N * G1 * O, O, 1, G1 * O} // V layout [G0, N, G1, O]
+                : std::vector<ck::index_t>{G1 * N * O, N * O, 1, O}; // V layout [G0, G1, N, O]
+
+        std::vector<ck::index_t> y_gs_ms_os_lengths{G0, G1, M, O};
+        std::vector<ck::index_t> y_gs_ms_os_strides =
+            output_permute
+                ? std::vector<ck::index_t>{M * G1 * O, O, G1 * O, 1} // Y layout [G0, M, G1, O]
+                : std::vector<ck::index_t>{G1 * M * O, M * O, O, 1}; // Y layout [G0, G1, M, O]
+
+        std::vector<ck::index_t> z_gs_ms_ns_lengths{G0, G1, M, N};
+        std::vector<ck::index_t> z_gs_ms_ns_strides =
+            input_permute
+                ? std::vector<ck::index_t>{M * G1 * N, N, G1 * N, 1} // Z layout [G0, M, G1, N]
+                : std::vector<ck::index_t>{G1 * M * N, M * N, N, 1}; // Z layout [G0, G1, M, N]
+        // The softmax stat log-sum-exp (LSE) is used to speed up softmax calculation in backward
+        // pass Pi = exp(Si) / sum(exp(S0) + exp(S1) + ...)
+        //    = exp(Si) / exp(log(sum(exp() + ...)))
+        //    = exp(Si - log(sum(exp() + ...)))
+        //               ^^^^^^^^^^^^^^^^^^^^^
+        //                       LSE
+        std::vector<ck::index_t> lse_gs_ms_lengths{G0, G1, M};
+        std::vector<ck::index_t> lse_gs_ms_strides{G1 * M, M, 1}; // LSE layout [G0, G1, M]
+        problem_descs.push_back({
+            q_gs_ms_ks_lengths,
+            q_gs_ms_ks_strides,
+            k_gs_ns_ks_lengths,
+            k_gs_ns_ks_strides,
+            z_gs_ms_ns_lengths,
+            z_gs_ms_ns_strides,
+            v_gs_os_ns_lengths,
+            v_gs_os_ns_strides,
+            y_gs_ms_os_lengths,
+            y_gs_ms_os_strides,
+            lse_gs_ms_lengths,
+            lse_gs_ms_strides,
+            {}, // std::array<std::vector<ck::index_t>, 1>{acc0_biases_gs_ms_ns_lengths},
+            {}, // std::array<std::vector<ck::index_t>, 1>{acc0_biases_gs_ms_ns_strides},
+            {}, // std::array<std::vector<ck::index_t>, 1>{acc1_biases_gs_ms_os_lengths},
+            {}, // std::array<std::vector<ck::index_t>, 1>{acc1_biases_gs_ms_os_strides},
+        });
+
+        int BatchCount = G0 * G1;
+        flop += (size_t(3) * M * N * K + size_t(2) * M * N * O) * 2 * BatchCount;
+        // Q/K/V/Y, dQ/dK/dV/dY, LSE
+        num_byte += (sizeof(DataType) * M * K + sizeof(DataType) * K * N +
+                     sizeof(DataType) * N * O + sizeof(DataType) * M * O) *
+                        size_t(2) * BatchCount +
+                    sizeof(LSEDataType) * M * BatchCount;
+
+        Tensor<DataType> q_gs_ms_ks(q_gs_ms_ks_lengths, q_gs_ms_ks_strides);
+        Tensor<DataType> k_gs_ns_ks(k_gs_ns_ks_lengths, k_gs_ns_ks_strides);
+        Tensor<ZDataType> z_gs_ms_ns(z_gs_ms_ns_lengths, z_gs_ms_ns_strides);
+        Tensor<DataType> v_gs_os_ns(v_gs_os_ns_lengths, v_gs_os_ns_strides);
+        Tensor<DataType> y_gs_ms_os(y_gs_ms_os_lengths, y_gs_ms_os_strides);
+        Tensor<DataType> ygrad_gs_ms_os(y_gs_ms_os_lengths, y_gs_ms_os_strides);
+        Tensor<LSEDataType> lse_gs_ms(lse_gs_ms_lengths, lse_gs_ms_strides);
+        if(i < 4)
+        {
+            std::cout << "q_gs_ms_ks: " << q_gs_ms_ks.mDesc << std::endl;
+            std::cout << "k_gs_ns_ks: " << k_gs_ns_ks.mDesc << std::endl;
+            std::cout << "v_gs_os_ns: " << v_gs_os_ns.mDesc << std::endl;
+            std::cout << "y_gs_ms_os: " << y_gs_ms_os.mDesc << std::endl;
+            std::cout << "lse_gs_ms_os: " << lse_gs_ms.mDesc << std::endl;
+        }
+        z_gs_ms_ns.GenerateTensorValue(GeneratorTensor_1<DataType>{0});
+        switch(init_method)
+        {
+        case 0: break;
+        case 1:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_2<DataType>{-2, 2});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_2<DataType>{-2, 2});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_2<DataType>{-2, 2});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_2<DataType>{-2, 2});
+            break;
+        case 2:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_3<DataType>{0.0, 1.0});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_3<DataType>{0.0, 1.0});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_3<DataType>{-0.5, 0.5});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_3<DataType>{-0.5, 0.5});
+            break;
+        case 3:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_2<DataType>{-5, 5});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            break;
+        case 4:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_1<DataType>{1});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_1<DataType>{2});
+            break;
+        case 5:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_1<DataType>{1});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_Sequential<2>{}); // dy[g0, g1, m, o]
+            // dO dot O = [0; 1; 2; ...]
+            break;
+        case 6:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_1<DataType>{1});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_Sequential<3>{}); // dy[g0, g1, m, o]
+            // assume mnko = 256
+            // P = softmax(QK) = 0.0039 * ones
+            // O = P V = 0.0039 * ones
+            // dP = dO V = [0, 1, 2, ...; 0, 1, 2, ...; ...]
+            // dO dot O = [127.5; ...]
+            // dS = P * (dP - dO dot O)
+            //
+            break;
+        default:
+            q_gs_ms_ks.GenerateTensorValue(GeneratorTensor_1<DataType>{1});
+            k_gs_ns_ks.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            v_gs_os_ns.GenerateTensorValue(GeneratorTensor_Diagonal<DataType>{});
+            ygrad_gs_ms_os.GenerateTensorValue(GeneratorTensor_1<DataType>{1}); // dy[g0, g1, m, o]
+            // assume mnko = 256
+            // P = softmax(QK) = 0.0039 * ones
+            // O = P V = 0.0039 * ones
+            // dP = dO V = ones
+            // dS = P * (dP - (dO dot O))
+            //    = 0.0039 * ones * (ones - 0.0039*256)
+            //    = 0.0039 * ones * (ones - 1)
+            //    = 0
+        }
+        Tensor<DataType> q_g_m_k({BatchCount, M, K});
+        Tensor<DataType> k_g_n_k({BatchCount, N, K});
+        Tensor<ZDataType> z_g_m_n({BatchCount, M, N});
+        Tensor<DataType> v_g_n_o({BatchCount, N, O});
+        Tensor<AccDataType> s_g_m_n({BatchCount, M, N});
+        Tensor<DataType> p_g_m_n({BatchCount, M, N});
+        Tensor<DataType> y_g_m_o({BatchCount, M, O});
+        Tensor<LSEDataType> lse_g_m({BatchCount, M});
+        Tensor<DataType> p_drop_g_m_n({BatchCount, M, N});
+
+        q_gs_ms_ks.ForEach([&](auto& self, auto idx) {
+            q_g_m_k(idx[0] * G1 + idx[1], idx[2], idx[3]) = self(idx);
+        });
+        k_gs_ns_ks.ForEach([&](auto& self, auto idx) {
+            k_g_n_k(idx[0] * G1 + idx[1], idx[2], idx[3]) = self(idx);
+        });
+        z_gs_ms_ns.ForEach([&](auto& self, auto idx) {
+            z_g_m_n(idx[0] * G1 + idx[1], idx[2], idx[3]) = self(idx);
+        });
+        v_gs_os_ns.ForEach([&](auto& self, auto idx) {
+            v_g_n_o(idx[0] * G1 + idx[1], idx[3], idx[2]) = self(idx);
+        });
+        lse_gs_ms.ForEach(
+            [&](auto& self, auto idx) { lse_g_m(idx[0] * G1 + idx[1], idx[2]) = self(idx); });
+
+        run_attention_fwd_host(q_g_m_k,
+                               k_g_n_k,
+                               v_g_n_o,
+                               alpha,
+                               s_g_m_n,
+                               p_g_m_n,
+                               y_g_m_o,
+                               lse_g_m,
+                               p_drop_g_m_n,
+                               z_g_m_n,
+                               p_dropout_in_16bits,
+                               rp_dropout);
+
+        y_gs_ms_os.ForEach([&](auto& self, auto idx) {
+            self(idx) = y_g_m_o(idx[0] * G1 + idx[1], idx[2], idx[3]);
+        });
+        lse_gs_ms.ForEach(
+            [&](auto& self, auto idx) { self(idx) = lse_g_m(idx[0] * G1 + idx[1], idx[2]); });
+
+        q_g_m_ks.push_back(q_g_m_k);
+        k_g_n_ks.push_back(k_g_n_k);
+        z_g_m_ns.push_back(z_g_m_n);
+        v_g_n_os.push_back(v_g_n_o);
+        s_g_m_ns.push_back(s_g_m_n);
+        p_g_m_ns.push_back(p_g_m_n);
+        y_g_m_os.push_back(y_g_m_o);
+        lse_g_ms.push_back(lse_g_m);
+        q_tensors.push_back(q_gs_ms_ks);
+        k_tensors.push_back(k_gs_ns_ks);
+        v_tensors.push_back(v_gs_os_ns);
+        y_tensors.push_back(y_gs_ms_os);
+        z_tensors.push_back(z_gs_ms_ns);
+        lse_tensors.push_back(lse_gs_ms);
+        ygrad_tensors.push_back(ygrad_gs_ms_os);
+        p_drop_g_m_ns.push_back(p_drop_g_m_n);
+        q_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * q_gs_ms_ks.GetElementSpaceSize()));
+        k_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * k_gs_ns_ks.GetElementSpaceSize()));
+        z_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(ZDataType) * z_gs_ms_ns.GetElementSpaceSize()));
+        v_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * v_gs_os_ns.GetElementSpaceSize()));
+        y_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * y_gs_ms_os.GetElementSpaceSize()));
+        lse_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(LSEDataType) * lse_gs_ms.GetElementSpaceSize()));
+        qgrad_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * q_gs_ms_ks.GetElementSpaceSize()));
+        kgrad_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * k_gs_ns_ks.GetElementSpaceSize()));
+        vgrad_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * v_gs_os_ns.GetElementSpaceSize()));
+        ygrad_tensors_device.emplace_back(
+            std::make_unique<DeviceMem>(sizeof(DataType) * y_gs_ms_os.GetElementSpaceSize()));
+        q_tensors_device.back()->ToDevice(q_gs_ms_ks.data());
+        k_tensors_device.back()->ToDevice(k_gs_ns_ks.data());
+        z_tensors_device.back()->ToDevice(z_gs_ms_ns.data());
+        v_tensors_device.back()->ToDevice(v_gs_os_ns.data());
+        y_tensors_device.back()->ToDevice(y_gs_ms_os.data());
+        lse_tensors_device.back()->ToDevice(lse_gs_ms.data());
+        qgrad_tensors_device.back()->SetZero();
+        kgrad_tensors_device.back()->SetZero();
+        vgrad_tensors_device.back()->SetZero();
+        ygrad_tensors_device.back()->ToDevice(ygrad_gs_ms_os.data());
+        p_q.push_back(q_tensors_device.back()->GetDeviceBuffer());
+        p_k.push_back(k_tensors_device.back()->GetDeviceBuffer());
+        p_z.push_back(z_tensors_device.back()->GetDeviceBuffer());
+        p_v.push_back(v_tensors_device.back()->GetDeviceBuffer());
+        p_y.push_back(y_tensors_device.back()->GetDeviceBuffer());
+        p_lse.push_back(lse_tensors_device.back()->GetDeviceBuffer());
+        p_kgrad.push_back(kgrad_tensors_device.back()->GetDeviceBuffer());
+        p_vgrad.push_back(vgrad_tensors_device.back()->GetDeviceBuffer());
+        p_ygrad.push_back(ygrad_tensors_device.back()->GetDeviceBuffer());
+        p_qgrad.push_back(qgrad_tensors_device.back()->GetDeviceBuffer());
+    }
+    auto argument =
+        gemm.MakeArgument(p_q,
+                          p_k,
+                          p_z,
+                          p_v,
+                          p_y,
+                          p_lse,
+                          p_ygrad,
+                          p_qgrad,
+                          p_kgrad,
+                          p_vgrad,
+                          {}, // std::array<void*, 1> p_acc0_biases;
+                          {}, // std::array<void*, 1> p_acc1_biases;
+                          problem_descs,
+                          QKVElementOp{},
+                          QKVElementOp{},
+                          Scale{alpha},
+                          QKVElementOp{},
+                          YElementOp{},
+                          p_drop,
+                          std::tuple<unsigned long long, unsigned long long>(seed, offset));
+
+    DeviceMem problem_desc_workspace(gemm.GetWorkSpaceSize(&argument));
+
+    gemm.SetWorkSpacePointer(&argument, problem_desc_workspace.GetDeviceBuffer());
+
+    if(!gemm.IsSupportedArgument(argument))
+    {
+        std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
+
+        return 0;
+    }
+
+    float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel});
+
+    float tflops = static_cast<float>(flop) / 1.E9 / ave_time;
+
+    float gb_per_sec = num_byte / 1.E6 / ave_time;
+
+    std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, "
+              << gemm.GetTypeString() << std::endl;
+
+    bool pass = true;
+    if(do_verification)
+    {
+
+        for(int i = 0; i < group_count; i++)
+        {
+            int G1 = v_tensors[i].GetLengths()[1];
+            z_tensors_device[i]->FromDevice(z_g_m_ns[i].data());
+            run_attention_fwd_host(q_g_m_ks[i],
+                k_g_n_ks[i],
+                v_g_n_os[i],
+                alpha,
+                s_g_m_ns[i],
+                p_g_m_ns[i],
+                y_g_m_os[i],
+                lse_g_ms[i],
+                p_drop_g_m_ns[i],
+                z_g_m_ns[i],
+                p_dropout_in_16bits,
+                rp_dropout);
+
+            y_tensors[i].ForEach([&](auto& self, auto idx) {
+                self(idx) = y_g_m_os[i](idx[0] * G1 + idx[1], idx[2], idx[3]);
+            });
+            y_tensors_device[i]->ToDevice(y_tensors[i].data());
+            qgrad_tensors_device[i]->SetZero();
+            kgrad_tensors_device[i]->SetZero();
+            vgrad_tensors_device[i]->SetZero();
+        }
+        p_z = std::vector<void*>(p_z.size(), nullptr);
+        argument =
+            gemm.MakeArgument(p_q,
+                              p_k,
+                              p_z,
+                              p_v,
+                              p_y,
+                              p_lse,
+                              p_ygrad,
+                              p_qgrad,
+                              p_kgrad,
+                              p_vgrad,
+                              {}, // std::array<void*, 1> p_acc0_biases;
+                              {}, // std::array<void*, 1> p_acc1_biases;
+                              problem_descs,
+                              QKVElementOp{},
+                              QKVElementOp{},
+                              Scale{alpha},
+                              QKVElementOp{},
+                              YElementOp{},
+                              p_drop,
+                              std::tuple<unsigned long long, unsigned long long>(seed, offset));
+        DeviceMem problem_desc_workspace(gemm.GetWorkSpaceSize(&argument));
+        gemm.SetWorkSpacePointer(&argument, problem_desc_workspace.GetDeviceBuffer());
+        if(!gemm.IsSupportedArgument(argument))
+        {
+            std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
+
+            return 0;
+        }
+        invoker.Run(argument, StreamConfig{nullptr, false});
+        for(std::size_t i = 0; i < group_count; i++)
+        {
+
+            int G0         = v_tensors[i].GetLengths()[0];
+            int G1         = v_tensors[i].GetLengths()[1];
+            int O          = v_tensors[i].GetLengths()[2];
+            int N          = v_tensors[i].GetLengths()[3];
+            int M          = q_tensors[i].GetLengths()[2];
+            int K          = q_tensors[i].GetLengths()[3];
+            int BatchCount = G0 * G1;
+            Tensor<DataType> qgrad_g_m_k({BatchCount, M, K});
+            Tensor<DataType> kgrad_g_n_k({BatchCount, N, K});
+            Tensor<DataType> vgrad_g_n_o({BatchCount, N, O});
+            Tensor<DataType> sgrad_g_m_n({BatchCount, M, N});
+            Tensor<DataType> pgrad_g_m_n({BatchCount, M, N});
+            Tensor<DataType> pgrad_drop_g_m_n({BatchCount, M, N});
+            Tensor<DataType> ygrad_g_m_o({BatchCount, M, O});
+
+            ygrad_tensors[i].ForEach([&](auto& self, auto idx) {
+                ygrad_g_m_o(idx[0] * G1 + idx[1], idx[2], idx[3]) = self(idx);
+            });
+            auto ref_gemm_grad         = ReferenceGemmGradInstance{};
+            auto ref_gemm_grad_invoker = ref_gemm_grad.MakeInvoker();
+            using RefGemmGradArg       = ReferenceGemmGradInstance::Argument;
+            // dP = dY * V^T
+            auto v_g_o_n = v_g_n_os[i].Transpose({0, 2, 1});
+            ref_gemm_grad_invoker.Run(RefGemmGradArg{
+                ygrad_g_m_o, v_g_o_n, pgrad_drop_g_m_n, PassThrough{}, PassThrough{}, Scale{1.f}});
+            auto ref_dropout         = ReferenceDropoutInstance{};
+            auto ref_dropout_invoker = ref_dropout.MakeInvoker();
+            auto ref_dropout_argment = ref_dropout.MakeArgument(
+                z_g_m_ns[i], pgrad_drop_g_m_n, pgrad_g_m_n, p_dropout_in_16bits, rp_dropout);
+            ref_dropout_invoker.Run(ref_dropout_argment);
+
+            sgrad_g_m_n.ForEach([&](auto& self, auto idx_gmn) {
+                float ygrad_dot_y = 0;
+                for(int o = 0; o < O; o++)
+                {
+                    auto idx_gmo = idx_gmn;
+                    idx_gmo[2]   = o;
+                    ygrad_dot_y += ygrad_g_m_o(idx_gmo) * y_g_m_os[i](idx_gmo);
+                }
+                self(idx_gmn) = p_g_m_ns[i](idx_gmn) * (pgrad_g_m_n(idx_gmn) - ygrad_dot_y);
+            });
+
+            auto p_drop_g_n_m = p_drop_g_m_ns[i].Transpose({0, 2, 1});
+            ref_gemm_grad_invoker.Run(RefGemmGradArg{
+                p_drop_g_n_m, ygrad_g_m_o, vgrad_g_n_o, PassThrough{}, PassThrough{}, Scale{1.0f}});
+            ref_gemm_grad_invoker.Run(RefGemmGradArg{
+                sgrad_g_m_n, k_g_n_ks[i], qgrad_g_m_k, PassThrough{}, PassThrough{}, Scale{alpha}});
+            auto sgrad_g_n_m = sgrad_g_m_n.Transpose({0, 2, 1});
+            ref_gemm_grad_invoker.Run(RefGemmGradArg{
+                sgrad_g_n_m, q_g_m_ks[i], kgrad_g_n_k, PassThrough{}, PassThrough{}, Scale{alpha}});
+
+            Tensor<DataType> qgrad_gs_ms_ks_host_result(q_tensors[i].GetLengths(),
+                                                        q_tensors[i].GetStrides());
+            Tensor<DataType> kgrad_gs_ns_ks_host_result(k_tensors[i].GetLengths(),
+                                                        k_tensors[i].GetStrides());
+            Tensor<DataType> vgrad_gs_os_ns_host_result(v_tensors[i].GetLengths(),
+                                                        v_tensors[i].GetStrides());
+
+            Tensor<DataType> qgrad_gs_ms_ks_device_result(q_tensors[i].GetLengths(),
+                                                          q_tensors[i].GetStrides());
+            Tensor<DataType> kgrad_gs_ns_ks_device_result(k_tensors[i].GetLengths(),
+                                                          k_tensors[i].GetStrides());
+            Tensor<DataType> vgrad_gs_os_ns_device_result(v_tensors[i].GetLengths(),
+                                                          v_tensors[i].GetStrides());
+
+            qgrad_tensors_device[i]->FromDevice(qgrad_gs_ms_ks_device_result.data());
+            kgrad_tensors_device[i]->FromDevice(kgrad_gs_ns_ks_device_result.data());
+            vgrad_tensors_device[i]->FromDevice(vgrad_gs_os_ns_device_result.data());
+            // permute
+            qgrad_gs_ms_ks_host_result.ForEach([&](auto& self, auto idx) {
+                const size_t& g0 = idx[0];
+                const size_t& g1 = idx[1];
+
+                const size_t g = g0 * G1 + g1;
+
+                self(idx) = qgrad_g_m_k(g, idx[2], idx[3]);
+            });
+            kgrad_gs_ns_ks_host_result.ForEach([&](auto& self, auto idx) {
+                const size_t& g0 = idx[0];
+                const size_t& g1 = idx[1];
+
+                const size_t g = g0 * G1 + g1;
+
+                self(idx) = kgrad_g_n_k(g, idx[2], idx[3]);
+            });
+            vgrad_gs_os_ns_host_result.ForEach([&](auto& self, auto idx) {
+                const size_t& g0 = idx[0];
+                const size_t& g1 = idx[1];
+
+                const size_t g = g0 * G1 + g1;
+
+                self(idx) = vgrad_g_n_o(g, idx[3], idx[2]);
+            });
+
+            std::cout << "Checking qgrad:\n";
+            pass &= ck::utils::check_err(qgrad_gs_ms_ks_device_result.mData,
+                                         qgrad_gs_ms_ks_host_result.mData,
+                                         "error",
+                                         1e-2,
+                                         1e-2);
+            std::cout << "Checking kgrad:\n";
+            pass &= ck::utils::check_err(kgrad_gs_ns_ks_device_result.mData,
+                                         kgrad_gs_ns_ks_host_result.mData,
+                                         "error",
+                                         1e-2,
+                                         1e-2);
+            std::cout << "Checking vgrad:\n";
+            pass &= ck::utils::check_err(vgrad_gs_os_ns_device_result.mData,
+                                         vgrad_gs_os_ns_host_result.mData,
+                                         "error",
+                                         1e-2,
+                                         1e-2);
+        }
+    }
+
+    return pass ? ((void)(std::cout << "pass\n"), 0) : ((void)(std::cout << "fail\n"), 1);
+}
+
+int main(int argc, char* argv[]) { return run(argc, argv); }

include/ck/tensor_operation/gpu/device/impl/device_grouped_multihead_attention_backward_xdl_cshuffle_pt1.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include <iostream>
+#include <sstream>
+
+#include "ck/utility/common_header.hpp"
+#include "ck/utility/philox_rand.hpp"
+#include "ck/tensor_description/tensor_descriptor.hpp"
+#include "ck/tensor_description/tensor_descriptor_helper.hpp"
+// #include "ck/tensor_operation/gpu/device/device_batched_multihead_attention_backward.hpp" // TODO
+#include "ck/tensor_operation/gpu/device/device_base.hpp"
+#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
+#include "ck/tensor_operation/gpu/device/masking_specialization.hpp"
+#include "ck/tensor_operation/gpu/device/matrix_padder.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_pt1.hpp"
+#include "ck/tensor_operation/operator_transform/transform_contraction_to_gemm.hpp"
+#include "ck/host_utility/device_prop.hpp"
+#include "ck/host_utility/kernel_launch.hpp"
+
+#include "ck/library/utility/host_tensor.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+
+template <typename GridwiseGemm,
+          typename GroupKernelArg,
+          typename AElementwiseOperation,
+          typename BElementwiseOperation,
+          typename AccElementwiseOperation,
+          typename B1ElementwiseOperation,
+          typename CElementwiseOperation,
+          bool HasMainKBlockLoop>
+__global__ void
+#if CK_USE_LAUNCH_BOUNDS
+    __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
+#endif
+        kernel_grouped_multihead_attention_backward_xdl_cshuffle_pt1(
+            const void CK_CONSTANT_ADDRESS_SPACE* group_kernel_args,
+            const index_t group_count,
+            const AElementwiseOperation a_element_op,
+            const BElementwiseOperation b_element_op,
+            const AccElementwiseOperation acc_element_op,
+            const B1ElementwiseOperation b1_element_op,
+            const CElementwiseOperation c_element_op,
+            const float p_dropout,
+            const unsigned long long seed,
+            const unsigned long long offset)
+{
+#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx908__) || defined(__gfx90a__))
+    __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
+    const index_t block_id = get_block_1d_id();
+    const auto arg_ptr     = reinterpret_cast<const GroupKernelArg*>(
+        cast_pointer_to_generic_address_space(group_kernel_args));
+
+    index_t left     = 0;
+    index_t right    = group_count;
+    index_t group_id = index_t((left + right) / 2);
+
+    while(
+        (!(block_id >= arg_ptr[group_id].block_start_ && block_id < arg_ptr[group_id].block_end_)))
+    {
+        if(block_id < arg_ptr[group_id].block_start_)
+        {
+            right = group_id;
+        }
+        else
+        {
+            left = group_id;
+        }
+        group_id = index_t((left + right) / 2);
+    }
+
+    // per-group batch offset
+    const index_t num_blocks_per_batch = arg_ptr[group_id].num_blocks_per_batch_;
+    const index_t g_idx                = __builtin_amdgcn_readfirstlane(
+        (block_id - arg_ptr[group_id].block_start_) / num_blocks_per_batch);
+
+    const long_index_t a_batch_offset = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetABasePtr(g_idx)));
+    const long_index_t b_batch_offset = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetBBasePtr(g_idx)));
+    const long_index_t z_batch_offset = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetZBasePtr(g_idx)));
+    const long_index_t b1_batch_offset = __builtin_amdgcn_readfirstlane(static_cast<long_index_t>(
+        arg_ptr[group_id].compute_base_ptr_of_batch_.GetB1BasePtr(g_idx)));
+    const long_index_t c_batch_offset  = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetCBasePtr(g_idx)));
+    const long_index_t lse_batch_offset = __builtin_amdgcn_readfirstlane(static_cast<long_index_t>(
+        arg_ptr[group_id].compute_base_ptr_of_batch_.GetLSEBasePtr(g_idx)));
+
+    const index_t global_thread_id = get_thread_global_1d_id();
+    ck::philox ph(seed, global_thread_id, offset);
+    unsigned short* z_matrix_ptr =
+        (arg_ptr[group_id].p_z_grid_ == nullptr ? nullptr
+                                                : arg_ptr[group_id].p_z_grid_ + z_batch_offset);
+    
+    GridwiseGemm::template Run<HasMainKBlockLoop>(
+        arg_ptr[group_id].p_a_grid_ + a_batch_offset,
+        arg_ptr[group_id].p_b_grid_ + b_batch_offset,
+        z_matrix_ptr,
+        arg_ptr[group_id].p_b1_grid_ + b1_batch_offset,
+        arg_ptr[group_id].p_c_grid_ + c_batch_offset,
+        arg_ptr[group_id].p_lse_grid_ + lse_batch_offset,
+        arg_ptr[group_id].p_ygrad_grid_ + c_batch_offset,
+        arg_ptr[group_id].p_qgrad_grid_ + a_batch_offset,
+        arg_ptr[group_id].p_kgrad_grid_ + b_batch_offset,
+        arg_ptr[group_id].p_vgrad_grid_ + b1_batch_offset,
+        p_shared,
+        a_element_op,
+        b_element_op,
+        acc_element_op,
+        b1_element_op,
+        c_element_op,
+        arg_ptr[group_id].a_grid_desc_ak0_m_ak1_,
+        arg_ptr[group_id].b_grid_desc_bk0_n_bk1_,
+        arg_ptr[group_id].c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5_,
+        arg_ptr[group_id].b1_grid_desc_bk0_n_bk1_,
+        arg_ptr[group_id].y_grid_desc_mblock_mperblock_oblock_operblock_,
+        arg_ptr[group_id].lse_grid_desc_m_,
+        arg_ptr[group_id].vgrad_grid_desc_n_o_,
+        arg_ptr[group_id].ygrad_grid_desc_o0_m_o1_,
+        arg_ptr[group_id].block_2_ctile_map_,
+        arg_ptr[group_id].c0_matrix_mask_,
+        p_dropout,
+        ph);
+#else
+    ignore = group_kernel_args;
+    ignore = group_count;
+    ignore = a_element_op;
+    ignore = b_element_op;
+    ignore = acc_element_op;
+    ignore = b1_element_op;
+    ignore = c_element_op;
+    ignore = p_dropout;
+    ignore = seed;
+    ignore = offset;
+#endif // end of if (defined(__gfx908__) || defined(__gfx90a__))
+}
+
+// Computes C = A * B0 * B1
+//              ^^^^^^ (Acc0)
+//              ^^^^^^^^^^^ (Acc1)
+template <index_t NumDimG,
+          index_t NumDimM,
+          index_t NumDimN,
+          index_t NumDimK,
+          index_t NumDimO, // NumDimGemm1N
+          typename DataType,
+          typename GemmDataType,
+          typename ZDataType,
+          typename LSEDataType,
+          typename Acc0BiasDataType,
+          typename Acc1BiasDataType,
+          typename GemmAccDataType,
+          typename CShuffleDataType,
+          typename AElementwiseOperation,
+          typename BElementwiseOperation,
+          typename AccElementwiseOperation,
+          typename B1ElementwiseOperation,
+          typename CElementwiseOperation,
+          GemmSpecialization GemmSpec,
+          TensorSpecialization ASpec,
+          TensorSpecialization BSpec,
+          TensorSpecialization B1Spec,
+          TensorSpecialization CSpec,
+          index_t NumGemmKPrefetchStage,
+          index_t BlockSize,
+          index_t MPerBlock,
+          index_t NPerBlock, // Gemm0NPerBlock
+          index_t KPerBlock, // Gemm0KPerBlock
+          index_t Gemm1NPerBlock,
+          index_t Gemm1KPerBlock,
+          index_t AK1,
+          index_t BK1,
+          index_t B1K1,
+          index_t MPerXDL,
+          index_t NPerXDL,
+          index_t MXdlPerWave,
+          index_t NXdlPerWave,
+          index_t Gemm1NXdlPerWave,
+          index_t Gemm2NXdlPerWave,
+          typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
+          typename ABlockTransferThreadClusterArrangeOrder,
+          typename ABlockTransferSrcAccessOrder,
+          index_t ABlockTransferSrcVectorDim,
+          index_t ABlockTransferSrcScalarPerVector,
+          index_t ABlockTransferDstScalarPerVector_AK1,
+          bool ABlockLdsExtraM,
+          typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
+          typename BBlockTransferThreadClusterArrangeOrder,
+          typename BBlockTransferSrcAccessOrder,
+          index_t BBlockTransferSrcVectorDim,
+          index_t BBlockTransferSrcScalarPerVector,
+          index_t BBlockTransferDstScalarPerVector_BK1,
+          bool BBlockLdsExtraN,
+          typename B1BlockTransferThreadClusterLengths_BK0_N_BK1,
+          typename B1BlockTransferThreadClusterArrangeOrder,
+          typename B1BlockTransferSrcAccessOrder,
+          index_t B1BlockTransferSrcVectorDim,
+          index_t B1BlockTransferSrcScalarPerVector,
+          index_t B1BlockTransferDstScalarPerVector_BK1,
+          bool B1BlockLdsExtraN,
+          index_t CShuffleMXdlPerWavePerShuffle,
+          index_t CShuffleNXdlPerWavePerShuffle,
+          typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
+          index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
+          MaskingSpecialization MaskingSpec,
+          LoopScheduler LoopSched = LoopScheduler::Default>
+struct DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_PT1
+    : public BaseOperator // TODO inherit atten bwd op once API stablizes
+{
+    static_assert(NumDimG > 0 && NumDimM > 0 && NumDimN > 0 && NumDimK > 0 && NumDimO > 0,
+                  "Number of dimension must be greater than 0");
+
+    static constexpr index_t NumAcc0Bias = Acc0BiasDataType::Size();
+    static constexpr index_t NumAcc1Bias = Acc1BiasDataType::Size();
+
+    // TODO: implement bias combination
+    static_assert(NumAcc0Bias == 0 && NumAcc0Bias == 0, "Bias addition is unimplemented");
+
+    using DeviceOp = DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_PT1;
+    struct ProblemDesc
+    {
+        std::vector<index_t> a_gs_ms_ks_lengths;
+        std::vector<index_t> a_gs_ms_ks_strides;
+
+        std::vector<index_t> b_gs_ns_ks_lengths;
+        std::vector<index_t> b_gs_ns_ks_strides;
+
+        std::vector<index_t> z_gs_ms_ns_lengths;
+        std::vector<index_t> z_gs_ms_ns_strides;
+
+        std::vector<index_t> b1_gs_gemm1ns_gemm1ks_lengths;
+        std::vector<index_t> b1_gs_gemm1ns_gemm1ks_strides;
+
+        std::vector<index_t> c_gs_ms_gemm1ns_lengths;
+        std::vector<index_t> c_gs_ms_gemm1ns_strides;
+
+        std::vector<index_t> lse_gs_ms_lengths;
+        std::vector<index_t> lse_gs_ms_strides;
+
+        std::vector<std::vector<index_t>> acc0_biases_gs_ms_ns_lengths;
+        std::vector<std::vector<index_t>> acc0_biases_gs_ms_ns_strides;
+
+        std::vector<std::vector<index_t>> acc1_biases_gs_ms_os_lengths;
+        std::vector<std::vector<index_t>> acc1_biases_gs_ms_os_strides;
+    };
+    static constexpr auto I0 = Number<0>{};
+    static constexpr auto I1 = Number<1>{};
+    static constexpr auto I2 = Number<2>{};
+
+    static constexpr index_t Q_K1 = 8;
+    static constexpr index_t K_K1 = 8;
+    static constexpr index_t V_N1 = 2;
+
+    static constexpr index_t Q_M1 = 2;
+    static constexpr index_t K_N1 = 2;
+    static constexpr index_t V_O1 = 8;
+    static constexpr index_t Y_O1 = 8;
+    static constexpr index_t Y_M1 = 2;
+
+    static constexpr auto padder = GemmGemmPadder<GemmSpec,
+                                                  Number<MPerBlock>,
+                                                  Number<NPerBlock>,
+                                                  Number<KPerBlock>,
+                                                  Number<Gemm1NPerBlock>>{};
+
+    using Transform = TransformBatchedContractionContractionToBatchedGemmGemm<
+        Sequence<NumDimG, NumDimM, NumDimN, NumDimK, NumDimO>,
+        Sequence<MPerBlock, NPerBlock, KPerBlock, Gemm1NPerBlock>,
+        GemmSpec,
+        ASpec,
+        BSpec,
+        B1Spec,
+        CSpec>;
+
+    /*
+    Descriptors for inputs:
+
+      Q, K, V, Y, dY, per-row softmax stats
+
+    Descriptors for outputs:
+
+      dQ, dK, dV
+
+    */
+
+    // Q in Gemm A position
+    static auto MakeAGridDescriptor_AK0_M_AK1(const std::vector<index_t>& a_gs_ms_ks_lengths_vec,
+                                              const std::vector<index_t>& a_gs_ms_ks_strides_vec)
+    {
+        return Transform::MakeAGridDescriptor_AK0_M_AK1(
+            Transform::MakeAGridDescriptor_M_K(a_gs_ms_ks_lengths_vec, a_gs_ms_ks_strides_vec),
+            Number<AK1>{});
+    }
+
+    // K in Gemm B0 position
+    static auto MakeBGridDescriptor_BK0_N_BK1(const std::vector<index_t>& b_gs_ns_ks_lengths_vec,
+                                              const std::vector<index_t>& b_gs_ns_ks_strides_vec)
+    {
+        return Transform::MakeB0GridDescriptor_BK0_N_BK1(
+            Transform::MakeB0GridDescriptor_N_K(b_gs_ns_ks_lengths_vec, b_gs_ns_ks_strides_vec),
+            Number<BK1>{});
+    }
+    //
+    // dV = P^T * dY
+    //
+
+    // VGrad in Gemm C position
+    static auto MakeVGradGridDescriptor_N_O(const std::vector<index_t>& v_gs_os_ns_lengths_vec,
+                                            const std::vector<index_t>& v_gs_os_ns_strides_vec)
+    {
+        // v_gs_os_ns -> vgrad_gs_ns_os. O dims last because output is row-major.
+        // Here directly rearrange lengths/strides before constructing tensor descriptor to reduce
+        // transformation overhead
+        // TODO: This will be much easier when inputs are Gs, Ms, Ns, Os. So there's no need to
+        // extract subsequence and shuffle them.
+        const index_t num_dims = NumDimG + NumDimN + NumDimO;
+
+        // 0, 1, .. NumDimG - 1
+        std::vector<index_t> gs_ids(NumDimG);
+        std::iota(gs_ids.begin(), gs_ids.end(), 0);
+
+        // NumDimG, NumDimG + 1, ... NumDimG + NumDimO - 1
+        std::vector<index_t> os_ids(NumDimO);
+        std::iota(os_ids.begin(), os_ids.end(), NumDimG);
+
+        // NumDimG + NumDimO, NumDimG + NumDimO + 1, ... NumDimG + NumDimO + NumDimN - 1
+        std::vector<index_t> ns_ids(NumDimN);
+        std::iota(ns_ids.begin(), ns_ids.end(), NumDimG + NumDimO);
+
+        std::vector<index_t> ids_old2new;
+        ids_old2new.insert(ids_old2new.end(), gs_ids.begin(), gs_ids.end());
+        ids_old2new.insert(ids_old2new.end(), ns_ids.begin(), ns_ids.end());
+        ids_old2new.insert(ids_old2new.end(), os_ids.begin(), os_ids.end());
+
+        std::vector<index_t> v_gs_ns_os_lengths_vec(num_dims), v_gs_ns_os_strides_vec(num_dims);
+        for(int i = 0; i < num_dims; i++)
+        {
+            index_t id_new            = ids_old2new[i];
+            v_gs_ns_os_lengths_vec[i] = v_gs_os_ns_lengths_vec[id_new];
+            v_gs_ns_os_strides_vec[i] = v_gs_os_ns_strides_vec[id_new];
+        }
+
+        const auto vgrad_desc_nraw_oraw =
+            MakeGridDescriptorPair<NumDimG, NumDimN, NumDimO, TensorSpecialization::Default>(
+                v_gs_ns_os_lengths_vec, v_gs_ns_os_strides_vec)
+                .second;
+
+        return PadTensorDescriptor(vgrad_desc_nraw_oraw,
+                                   make_tuple(NPerBlock, Gemm1NPerBlock),
+                                   Sequence<padder.PadN, padder.PadO>{});
+    }
+
+    //
+    // dQ = alpha * dS * K
+    //
+    static auto MakeYGradGridDescriptor_O0_M_O1(const std::vector<index_t>& y_gs_ms_os_lengths_vec,
+                                                const std::vector<index_t>& y_gs_ms_os_strides_vec)
+    {
+        return Transform::MakeAGridDescriptor_AK0_M_AK1(
+            Transform::MakeAGridDescriptor_M_K(y_gs_ms_os_lengths_vec, y_gs_ms_os_strides_vec),
+            Number<Y_O1>{});
+    }
+
+    // V in Gemm B position
+    static auto MakeVGridDescriptor_O0_N_O1(const std::vector<index_t>& v_gs_os_ns_lengths_vec,
+                                            const std::vector<index_t>& v_gs_os_ns_strides_vec)
+    {
+        // v_gs_os_ns -> vgrad_gs_ns_os. O dims last because output is row-major.
+        // Here directly rearrange lengths/strides before constructing tensor descriptor to reduce
+        // transformation overhead
+        // TODO: This will be much easier when inputs are Gs, Ms, Ns, Os. So there's no need to
+        // extract subsequence and shuffle them.
+        const index_t num_dims = NumDimG + NumDimN + NumDimO;
+
+        // 0, 1, .. NumDimG - 1
+        std::vector<index_t> gs_ids(NumDimG);
+        std::iota(gs_ids.begin(), gs_ids.end(), 0);
+
+        // NumDimG, NumDimG + 1, ... NumDimG + NumDimO - 1
+        std::vector<index_t> os_ids(NumDimO);
+        std::iota(os_ids.begin(), os_ids.end(), NumDimG);
+
+        // NumDimG + NumDimO, NumDimG + NumDimO + 1, ... NumDimG + NumDimO + NumDimN - 1
+        std::vector<index_t> ns_ids(NumDimN);
+        std::iota(ns_ids.begin(), ns_ids.end(), NumDimG + NumDimO);
+
+        std::vector<index_t> ids_old2new;
+        ids_old2new.insert(ids_old2new.end(), gs_ids.begin(), gs_ids.end());
+        ids_old2new.insert(ids_old2new.end(), ns_ids.begin(), ns_ids.end());
+        ids_old2new.insert(ids_old2new.end(), os_ids.begin(), os_ids.end());
+
+        std::vector<index_t> v_gs_ns_os_lengths_vec(num_dims), v_gs_ns_os_strides_vec(num_dims);
+        for(int i = 0; i < num_dims; i++)
+        {
+            index_t id_new            = ids_old2new[i];
+            v_gs_ns_os_lengths_vec[i] = v_gs_os_ns_lengths_vec[id_new];
+            v_gs_ns_os_strides_vec[i] = v_gs_os_ns_strides_vec[id_new];
+        }
+
+        const auto v_grid_desc_nraw_oraw =
+            MakeGridDescriptorPair<NumDimG, NumDimN, NumDimO, TensorSpecialization::Default>(
+                v_gs_ns_os_lengths_vec, v_gs_ns_os_strides_vec)
+                .second;
+
+        const auto v_grid_desc_n_o = PadTensorDescriptor(v_grid_desc_nraw_oraw,
+                                                         make_tuple(NPerBlock, Gemm1NPerBlock),
+                                                         Sequence<padder.PadN, padder.PadO>{});
+
+        // N_O to O0_N_O1; to refactor
+        return Transform::MakeB0GridDescriptor_BK0_N_BK1(v_grid_desc_n_o, Number<V_O1>{});
+    }
+
+    static auto MakeZGridDescriptor_M_N(const std::vector<index_t>& z_gs_ms_ns_lengths_vec,
+                                        const std::vector<index_t>& z_gs_ms_ns_strides_vec)
+    {
+        return Transform::MakeCGridDescriptor_M_N(z_gs_ms_ns_lengths_vec, z_gs_ms_ns_strides_vec);
+    }
+
+    static auto MakeLSEGridDescriptor_M(index_t MRaw)
+    {
+        const auto lse_grid_desc_mraw = make_naive_tensor_descriptor_packed(make_tuple(MRaw));
+
+        const auto M    = math::integer_divide_ceil(MRaw, MPerBlock) * MPerBlock;
+        const auto MPad = M - MRaw;
+
+        if constexpr(GemmSpec == GemmSpecialization::MPadding ||
+                     GemmSpec == GemmSpecialization::MNPadding ||
+                     GemmSpec == GemmSpecialization::MKPadding ||
+                     GemmSpec == GemmSpecialization::MNKPadding)
+        {
+            // pad M
+            return transform_tensor_descriptor(lse_grid_desc_mraw,
+                                               make_tuple(make_right_pad_transform(MRaw, MPad)),
+                                               make_tuple(Sequence<0>{}),
+                                               make_tuple(Sequence<0>{}));
+        }
+        else
+        {
+            // not pad M
+            return lse_grid_desc_mraw;
+        }
+    }
+
+
+    using AGridDesc_AK0_M_AK1  = decltype(MakeAGridDescriptor_AK0_M_AK1({}, {}));
+    using BGridDesc_BK0_N_BK1  = decltype(MakeBGridDescriptor_BK0_N_BK1({}, {}));
+    using B1GridDesc_BK0_N_BK1 = decltype(MakeBGridDescriptor_BK0_N_BK1({}, {}));
+    using YGridDesc_M_O        = decltype(Transform::MakeCGridDescriptor_M_N({}, {}));
+    using LSEGridDesc_M        = decltype(MakeLSEGridDescriptor_M(1));
+    using AGridDesc_G_M_K      = decltype(Transform::MakeAGridDescriptor_G_M_K({}, {}));
+    using BGridDesc_G_N_K      = decltype(Transform::MakeB0GridDescriptor_G_N_K({}, {}));
+    using B1GridDesc_G_N_K     = decltype(Transform::MakeB1GridDescriptor_G_N_K({}, {}));
+    using CGridDesc_G_M_N      = decltype(Transform::MakeCGridDescriptor_G_M_N({}, {}));
+    using ZGridDesc_G_M_N      = decltype(Transform::MakeCGridDescriptor_G_M_N({}, {}));
+
+    using VGradGridDesc_N_O     = decltype(MakeVGradGridDescriptor_N_O({}, {}));
+    using YGradGridDesc_O0_M_O1 = decltype(MakeYGradGridDescriptor_O0_M_O1({}, {}));
+    using ZGridDesc_M_N         = decltype(MakeZGridDescriptor_M_N({}, {}));
+
+    constexpr static auto make_MaskOutPredicate()
+    {
+        if constexpr(MaskingSpec == MaskingSpecialization::MaskDisabled)
+        {
+            return MaskDisabledPredicate{};
+        }
+        else if constexpr(MaskingSpec == MaskingSpecialization::MaskOutUpperTriangle)
+        {
+            return MaskOutUpperTrianglePredicate{};
+        }
+    }
+    using C0MatrixMask = C0MatrixMask_impl<decltype(make_MaskOutPredicate())>;
+
+    struct ComputeBasePtrOfStridedBatch
+    {
+        ComputeBasePtrOfStridedBatch(const AGridDesc_G_M_K& a_grid_desc_g_m_k,
+                                     const BGridDesc_G_N_K& b_grid_desc_g_n_k,
+                                     const ZGridDesc_G_M_N& z_grid_desc_g_m_n,
+                                     const B1GridDesc_G_N_K& b1_grid_desc_g_n_k,
+                                     const CGridDesc_G_M_N& c_grid_desc_g_m_n,
+                                     index_t batch_stride_lse)
+            : a_grid_desc_g_m_k_(a_grid_desc_g_m_k),
+              b_grid_desc_g_n_k_(b_grid_desc_g_n_k),
+              z_grid_desc_g_m_n_(z_grid_desc_g_m_n),
+              b1_grid_desc_g_n_k_(b1_grid_desc_g_n_k),
+              c_grid_desc_g_m_n_(c_grid_desc_g_m_n),
+              batch_stride_lse_(batch_stride_lse)
+        {
+        }
+
+        __host__ __device__ constexpr long_index_t GetABasePtr(index_t g_idx) const
+        {
+            return a_grid_desc_g_m_k_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetBBasePtr(index_t g_idx) const
+        {
+            return b_grid_desc_g_n_k_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetZBasePtr(index_t g_idx) const
+        {
+            return z_grid_desc_g_m_n_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetB1BasePtr(index_t g_idx) const
+        {
+            return b1_grid_desc_g_n_k_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetCBasePtr(index_t g_idx) const
+        {
+            return c_grid_desc_g_m_n_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetLSEBasePtr(index_t g_idx) const
+        {
+            return g_idx * static_cast<long_index_t>(batch_stride_lse_);
+        }
+
+        private:
+        AGridDesc_G_M_K a_grid_desc_g_m_k_;
+        BGridDesc_G_N_K b_grid_desc_g_n_k_;
+        ZGridDesc_G_M_N z_grid_desc_g_m_n_;
+        B1GridDesc_G_N_K b1_grid_desc_g_n_k_;
+        CGridDesc_G_M_N c_grid_desc_g_m_n_;
+        index_t batch_stride_lse_;
+    };
+
+    // GridwiseGemm
+    using GridwiseGemm = GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_PT1<
+        DataType, // TODO: distinguish A/B datatype
+        GemmDataType,
+        GemmAccDataType,
+        CShuffleDataType,
+        LSEDataType,
+        AElementwiseOperation,
+        BElementwiseOperation,
+        AccElementwiseOperation,
+        B1ElementwiseOperation,
+        CElementwiseOperation,
+        InMemoryDataOperationEnum::Set,
+        AGridDesc_AK0_M_AK1,
+        BGridDesc_BK0_N_BK1,
+        ZGridDesc_M_N,
+        B1GridDesc_BK0_N_BK1,
+        YGridDesc_M_O,
+        LSEGridDesc_M,
+        NumGemmKPrefetchStage,
+        BlockSize,
+        MPerBlock,
+        NPerBlock,
+        KPerBlock,
+        Gemm1NPerBlock,
+        Gemm1KPerBlock,
+        AK1,
+        BK1,
+        B1K1,
+        MPerXDL,
+        NPerXDL,
+        MXdlPerWave,
+        NXdlPerWave,
+        Gemm1NXdlPerWave,
+        Gemm2NXdlPerWave,
+        ABlockTransferThreadClusterLengths_AK0_M_AK1,
+        ABlockTransferThreadClusterArrangeOrder,
+        ABlockTransferSrcAccessOrder,
+        ABlockTransferSrcVectorDim,
+        ABlockTransferSrcScalarPerVector,
+        ABlockTransferDstScalarPerVector_AK1,
+        true,
+        ABlockLdsExtraM,
+        BBlockTransferThreadClusterLengths_BK0_N_BK1,
+        BBlockTransferThreadClusterArrangeOrder,
+        BBlockTransferSrcAccessOrder,
+        BBlockTransferSrcVectorDim,
+        BBlockTransferSrcScalarPerVector,
+        BBlockTransferDstScalarPerVector_BK1,
+        true,
+        BBlockLdsExtraN,
+        B1BlockTransferThreadClusterLengths_BK0_N_BK1,
+        B1BlockTransferThreadClusterArrangeOrder,
+        B1BlockTransferSrcAccessOrder,
+        B1BlockTransferSrcVectorDim,
+        B1BlockTransferSrcScalarPerVector,
+        B1BlockTransferDstScalarPerVector_BK1,
+        false,
+        B1BlockLdsExtraN,
+        CShuffleMXdlPerWavePerShuffle,
+        CShuffleNXdlPerWavePerShuffle,
+        CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
+        CShuffleBlockTransferScalarPerVector_NPerBlock,
+        LoopSched,
+        Transform::matrix_padder.PadN,
+        MaskingSpec == MaskingSpecialization::MaskOutUpperTriangle>;
+
+    using Block2CTileMap = OffsettedBlockToCTileMap<typename GridwiseGemm::DefaultBlock2CTileMap>;
+
+    struct GroupKernelArg
+    {
+        // pointers
+        const DataType* p_a_grid_;
+        const DataType* p_b_grid_;
+        ZDataType* p_z_grid_;
+        const DataType* p_b1_grid_;
+        const DataType* p_c_grid_;
+        const LSEDataType* p_lse_grid_;
+        const DataType* p_ygrad_grid_;
+        DataType* p_qgrad_grid_;
+        DataType* p_kgrad_grid_;
+        DataType* p_vgrad_grid_;
+
+        // tensor descriptors for block/thread-wise copy
+        AGridDesc_AK0_M_AK1 a_grid_desc_ak0_m_ak1_;
+        BGridDesc_BK0_N_BK1 b_grid_desc_bk0_n_bk1_;
+        ZGridDesc_M_N z_grid_desc_m_n_;
+        B1GridDesc_BK0_N_BK1 b1_grid_desc_bk0_n_bk1_;
+        YGridDesc_M_O y_grid_desc_m_o_;
+
+        typename GridwiseGemm::YGridDescriptor_MBlock_MPerBlock_OBlock_OPerBlock
+            y_grid_desc_mblock_mperblock_oblock_operblock_;
+        typename GridwiseGemm::ZGridDescriptor_M0_N0_M1_N1_M2_N2_M3_N3_N4_N5
+            c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5_;
+        LSEGridDesc_M lse_grid_desc_m_;
+        VGradGridDesc_N_O vgrad_grid_desc_n_o_;
+        YGradGridDesc_O0_M_O1 ygrad_grid_desc_o0_m_o1_;
+        // block-to-c-tile map
+        Block2CTileMap block_2_ctile_map_;
+        index_t num_blocks_per_batch_;
+        ComputeBasePtrOfStridedBatch compute_base_ptr_of_batch_;
+
+        // check C0 masking and padding
+        C0MatrixMask c0_matrix_mask_;
+        index_t block_start_, block_end_;
+    };
+
+    struct GroupDeviceArg
+    {
+        // lengths for the last dimensions of overall problem for sanity check of vector load/store
+        std::vector<index_t> raw_lengths_mz_nz_kz_gemm1nz_;
+
+        // strides for the last dimensions of each tensor for sanity check of vector load/store
+        std::vector<index_t> a_mz_kz_strides_;
+        std::vector<index_t> b_nz_kz_strides_;
+        std::vector<index_t> b1_nz_kz_strides_;
+        std::vector<index_t> c_mz_gemm1nz_strides_;
+
+        // for gridwise gemm check
+        CGridDesc_G_M_N c_grid_desc_g_m_n_;
+
+        index_t batch_count_;
+    };
+    // Argument
+    struct Argument : public BaseArgument
+    {
+        Argument(const std::vector<const void*>& p_As,
+                 const std::vector<const void*>& p_Bs,
+                 const std::vector<void*>& p_Zs,
+                 const std::vector<const void*>& p_B1s,
+                 const std::vector<const void*>& p_Cs, // for dS
+                 const std::vector<const void*>& p_LSEs,
+                 const std::vector<const void*>& p_Ygrads,
+                 std::vector<void*>& p_Qgrads,
+                 std::vector<void*>& p_Kgrads,
+                 std::vector<void*>& p_Vgrads,
+                 const std::array<void*, NumAcc0Bias>& p_acc0_biases,
+                 const std::array<void*, NumAcc1Bias>& p_acc1_biases,
+                 const std::vector<ProblemDesc>& problem_desc_vec,
+                 AElementwiseOperation a_element_op,
+                 BElementwiseOperation b_element_op,
+                 AccElementwiseOperation acc_element_op,
+                 B1ElementwiseOperation b1_element_op,
+                 CElementwiseOperation c_element_op,
+                 float p_drop,
+                 std::tuple<unsigned long long, unsigned long long> seeds)
+            : a_element_op_{a_element_op},
+              b_element_op_{b_element_op},
+              acc_element_op_{acc_element_op},
+              b1_element_op_{b1_element_op},
+              c_element_op_{c_element_op},
+              p_dropout_{p_drop}
+        {
+            seed_   = std::get<0>(seeds);
+            offset_ = std::get<1>(seeds);
+
+            group_count_ = ck::type_convert<ck::index_t>(problem_desc_vec.size());
+
+            if(!(group_count_ == ck::type_convert<ck::index_t>(p_As.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Bs.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Zs.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_B1s.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Cs.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Ygrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Qgrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Kgrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Vgrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_LSEs.size())))
+            {
+                throw std::runtime_error("wrong! group_count_ != p_As/b/b1/c.size");
+            }
+
+            if(!(p_acc0_biases.size() == p_acc1_biases.size()))
+            {
+                throw std::runtime_error("wrong! acc0_bias_vec.size != acc1_bias_vec.size");
+            }
+
+            grid_size_ = 0;
+            for(std::size_t i = 0; i < group_count_; i++)
+            {
+                const auto p_a_grid     = static_cast<const DataType*>(p_As[i]);
+                const auto p_b_grid     = static_cast<const DataType*>(p_Bs[i]);
+                auto p_z_grid           = static_cast<ZDataType*>(p_Zs[i]);
+                const auto p_b1_grid    = static_cast<const DataType*>(p_B1s[i]);
+                const auto p_c_grid     = static_cast<const DataType*>(p_Cs[i]);
+                const auto p_lse_grid   = static_cast<const LSEDataType*>(p_LSEs[i]);
+                const auto p_ygrad_grid = static_cast<const DataType*>(p_Ygrads[i]);
+                auto p_qgrad_grid       = static_cast<DataType*>(p_Qgrads[i]);
+                auto p_kgrad_grid       = static_cast<DataType*>(p_Kgrads[i]);
+                auto p_vgrad_grid       = static_cast<DataType*>(p_Vgrads[i]);
+
+                const auto& problem_desc = problem_desc_vec[i];
+
+                const auto a_grid_desc_ak0_m_ak1 = DeviceOp::MakeAGridDescriptor_AK0_M_AK1(
+                    problem_desc.a_gs_ms_ks_lengths, problem_desc.a_gs_ms_ks_strides);
+                const auto b_grid_desc_bk0_n_bk1 = DeviceOp::MakeBGridDescriptor_BK0_N_BK1(
+                    problem_desc.b_gs_ns_ks_lengths, problem_desc.b_gs_ns_ks_strides);
+                const auto z_grid_desc_m_n = DeviceOp::MakeZGridDescriptor_M_N(
+                    problem_desc.z_gs_ms_ns_lengths, problem_desc.z_gs_ms_ns_strides);
+                const auto b1_grid_desc_bk0_n_bk1 = DeviceOp::MakeVGridDescriptor_O0_N_O1(
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_lengths,
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_strides);
+                const auto y_grid_desc_m_o = Transform::MakeCGridDescriptor_M_N(
+                    problem_desc.c_gs_ms_gemm1ns_lengths, problem_desc.c_gs_ms_gemm1ns_strides);
+
+                const auto lse_grid_desc_m =
+                    DeviceOp::MakeLSEGridDescriptor_M(problem_desc.lse_gs_ms_lengths[NumDimG]);
+                const auto vgrad_grid_desc_n_o = DeviceOp::MakeVGradGridDescriptor_N_O(
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_lengths,
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_strides);
+                const auto ygrad_grid_desc_o0_m_o1 = DeviceOp::MakeYGradGridDescriptor_O0_M_O1(
+                    problem_desc.c_gs_ms_gemm1ns_lengths, problem_desc.c_gs_ms_gemm1ns_strides);
+
+                const auto a_grid_desc_g_m_k = Transform::MakeAGridDescriptor_G_M_K(
+                    problem_desc.a_gs_ms_ks_lengths, problem_desc.a_gs_ms_ks_strides);
+                const auto b_grid_desc_g_n_k = Transform::MakeB0GridDescriptor_G_N_K(
+                    problem_desc.b_gs_ns_ks_lengths, problem_desc.b_gs_ns_ks_strides);
+                const auto z_grid_desc_g_m_n = Transform::MakeCGridDescriptor_G_M_N(
+                    problem_desc.z_gs_ms_ns_lengths, problem_desc.z_gs_ms_ns_strides);
+                const auto b1_grid_desc_g_n_k = Transform::MakeB1GridDescriptor_G_N_K(
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_lengths,
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_strides);
+                const auto c_grid_desc_g_m_n = Transform::MakeCGridDescriptor_G_M_N(
+                    problem_desc.c_gs_ms_gemm1ns_lengths, problem_desc.c_gs_ms_gemm1ns_strides);
+                typename GridwiseGemm::YGridDescriptor_MBlock_MPerBlock_OBlock_OPerBlock
+                    y_grid_desc_mblock_mperblock_oblock_operblock;
+                typename GridwiseGemm::ZGridDescriptor_M0_N0_M1_N1_M2_N2_M3_N3_N4_N5
+                    c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5;
+                const index_t BlockStart     = grid_size_;
+                const auto block_2_ctile_map = Block2CTileMap(y_grid_desc_m_o, BlockStart);
+                if(GridwiseGemm::CheckValidity(a_grid_desc_ak0_m_ak1,
+                                               b_grid_desc_bk0_n_bk1,
+                                               b1_grid_desc_bk0_n_bk1,
+                                               y_grid_desc_m_o,
+                                               block_2_ctile_map))
+                {
+                    y_grid_desc_mblock_mperblock_oblock_operblock =
+                        GridwiseGemm::MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(
+                            y_grid_desc_m_o);
+                }
+
+                c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5 =
+                    GridwiseGemm::MakeCGridDescriptor_M0_N0_M1_N1_M2_N2_M3_N3_N4_N5(
+                        z_grid_desc_m_n);
+
+                const index_t batch_count = c_grid_desc_g_m_n.GetLength(I0);
+                const index_t grid_size_grp =
+                    block_2_ctile_map.CalculateGridSize(y_grid_desc_m_o) * batch_count;
+                const index_t BlockEnd = grid_size_ + grid_size_grp;
+                // batch stride
+                const auto compute_base_ptr_of_batch = ComputeBasePtrOfStridedBatch(
+                    a_grid_desc_g_m_k,
+                    b_grid_desc_g_n_k,
+                    z_grid_desc_g_m_n,
+                    b1_grid_desc_g_n_k,
+                    c_grid_desc_g_m_n,
+                    type_convert<index_t>(lse_grid_desc_m.GetElementSpaceSize()));
+
+                // C0 mask
+                const auto c0_matrix_mask = C0MatrixMask(b_grid_desc_g_n_k.GetLength(I1));
+
+                grid_size_ += grid_size_grp;
+
+                // for each group, make sure acc0_biases_gs_ms_ns_lengths.size() == NumAcc0Bias and
+                // so on
+                if(!(problem_desc.acc0_biases_gs_ms_ns_lengths.size() == NumAcc0Bias &&
+                     problem_desc.acc0_biases_gs_ms_ns_strides.size() == NumAcc0Bias &&
+                     problem_desc.acc1_biases_gs_ms_os_lengths.size() == NumAcc1Bias &&
+                     problem_desc.acc1_biases_gs_ms_os_strides.size() == NumAcc1Bias))
+                {
+                    throw std::runtime_error(
+                        "wrong! number of biases in function argument does not "
+                        "match that in template argument");
+                }
+
+                group_kernel_args_.push_back({p_a_grid,
+                                              p_b_grid,
+                                              p_z_grid,
+                                              p_b1_grid,
+                                              p_c_grid,
+                                              p_lse_grid,
+                                              p_ygrad_grid,
+                                              p_qgrad_grid,
+                                              p_kgrad_grid,
+                                              p_vgrad_grid,
+                                              a_grid_desc_ak0_m_ak1,
+                                              b_grid_desc_bk0_n_bk1,
+                                              z_grid_desc_m_n,
+                                              b1_grid_desc_bk0_n_bk1,
+                                              y_grid_desc_m_o,
+                                              y_grid_desc_mblock_mperblock_oblock_operblock,
+                                              c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5,
+                                              lse_grid_desc_m,
+                                              vgrad_grid_desc_n_o,
+                                              ygrad_grid_desc_o0_m_o1,
+                                              block_2_ctile_map,
+                                              block_2_ctile_map.CalculateGridSize(y_grid_desc_m_o),
+                                              compute_base_ptr_of_batch,
+                                              c0_matrix_mask,
+                                              BlockStart,
+                                              BlockEnd});
+
+                group_device_args_.push_back(
+                    {{problem_desc.a_gs_ms_ks_lengths[NumDimG + NumDimM - 1],
+                      problem_desc.b_gs_ns_ks_lengths[NumDimG + NumDimN - 1],
+                      problem_desc.b_gs_ns_ks_lengths[NumDimG + NumDimN + NumDimK - 1],
+                      problem_desc.b1_gs_gemm1ns_gemm1ks_lengths[NumDimG + NumDimO - 1]},
+                     {problem_desc.a_gs_ms_ks_strides[NumDimG + NumDimM - 1],
+                      problem_desc.a_gs_ms_ks_strides[NumDimG + NumDimM + NumDimK - 1]},
+                     {problem_desc.b_gs_ns_ks_strides[NumDimG + NumDimN - 1],
+                      problem_desc.b_gs_ns_ks_strides[NumDimG + NumDimN + NumDimK - 1]},
+                     {problem_desc.b1_gs_gemm1ns_gemm1ks_strides[NumDimG + NumDimO - 1],
+                      problem_desc.b1_gs_gemm1ns_gemm1ks_strides[NumDimG + NumDimO + NumDimN - 1]},
+                     {problem_desc.c_gs_ms_gemm1ns_strides[NumDimG + NumDimM - 1],
+                      problem_desc.c_gs_ms_gemm1ns_strides[NumDimG + NumDimM + NumDimO - 1]},
+                     c_grid_desc_g_m_n,
+                     batch_count});
+            }
+            // TODO: implement bias addition
+            // ignore = p_acc0_biases;
+            // ignore = p_acc1_biases;
+            // ignore = acc0_biases_gs_ms_ns_lengths;
+            // ignore = acc0_biases_gs_ms_ns_strides;
+            // ignore = acc1_biases_gs_ms_gemm1ns_lengths;
+            // ignore = acc1_biases_gs_ms_gemm1ns_strides;
+        }
+
+        // element-wise op
+        AElementwiseOperation a_element_op_;
+        BElementwiseOperation b_element_op_;
+        AccElementwiseOperation acc_element_op_;
+        B1ElementwiseOperation b1_element_op_;
+        CElementwiseOperation c_element_op_;
+
+        float p_dropout_;
+        unsigned long long seed_;
+        unsigned long long offset_;
+
+        index_t grid_size_;
+        index_t group_count_;
+
+        std::vector<GroupKernelArg> group_kernel_args_;
+        std::vector<GroupDeviceArg> group_device_args_;
+    };
+
+    // Invoker
+    struct Invoker : public BaseInvoker
+    {
+        using Argument = DeviceOp::Argument;
+
+        float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{})
+        {
+            if(!DeviceOp::IsSupportedArgument(arg))
+            {
+                throw std::runtime_error("wrong! unsupported argument");
+            }
+
+            bool all_has_main_k_block_loop  = false;
+            bool some_has_main_k_block_loop = false;
+            // for(std::size_t i = 0; i < arg.group_count_; i++)
+            // {
+            //     const auto K =
+            //         arg.group_kernel_args_[i].a_grid_desc_ak0_m_ak1_.GetLength(I0) * arg.group_kernel_args_[i].a_grid_desc_ak0_m_ak1_.GetLength(I2);
+            //     const bool y = GridwiseGemm::CalculateHasMainKBlockLoop(K);
+            //     all_has_main_k_block_loop &= y;
+            //     some_has_main_k_block_loop |= y;
+            // }
+
+            hipGetErrorString(hipMemcpy(arg.p_workspace_,
+                                        arg.group_kernel_args_.data(),
+                                        arg.group_kernel_args_.size() * sizeof(GroupKernelArg),
+                                        hipMemcpyHostToDevice));
+
+            float ave_time = 0;
+
+            auto launch_kernel = [&](auto has_main_k_block_loop_) {
+                const auto kernel = kernel_grouped_multihead_attention_backward_xdl_cshuffle_pt1<
+                    GridwiseGemm,
+                    GroupKernelArg,
+                    AElementwiseOperation,
+                    BElementwiseOperation,
+                    AccElementwiseOperation,
+                    B1ElementwiseOperation,
+                    CElementwiseOperation,
+                    has_main_k_block_loop_>;
+
+                return launch_and_time_kernel(
+                    stream_config,
+                    kernel,
+                    dim3(arg.grid_size_),
+                    dim3(BlockSize),
+                    0,
+                    cast_pointer_to_constant_address_space(arg.p_workspace_),
+                    arg.group_count_,
+                    arg.a_element_op_,
+                    arg.b_element_op_,
+                    arg.acc_element_op_,
+                    arg.b1_element_op_,
+                    arg.c_element_op_,
+                    arg.p_dropout_,
+                    arg.seed_,
+                    arg.offset_);
+            };
+
+            // Gemm1_K is split into Gemm1_K0/K1 where K1 is known at compile time, so we only need
+            // to concern Gemm0's loop
+            if(all_has_main_k_block_loop)
+            {
+                ave_time = launch_kernel(integral_constant<bool, true>{});
+            }
+            else if(!some_has_main_k_block_loop)
+            {
+                ave_time = launch_kernel(integral_constant<bool, false>{});
+            }
+            else
+            {
+                throw std::runtime_error("wrong! all gemm problems have to simultaneously meet "
+                                         "has_main_k_block_loop or no_main_k_block_loop");
+            }
+            return ave_time;
+        }
+
+        // polymorphic
+        float Run(const BaseArgument* p_arg,
+                  const StreamConfig& stream_config = StreamConfig{}) override
+        {
+            return Run(*dynamic_cast<const Argument*>(p_arg), stream_config);
+        }
+    };
+
+    static constexpr bool IsValidCompilationParameter()
+    {
+        // TODO: properly implement this check
+        return true;
+    }
+
+    static bool IsSupportedArgument(const Argument& arg)
+    {
+        if(!(ck::get_device_name() == "gfx908" || ck::get_device_name() == "gfx90a"))
+        {
+            return false;
+        }
+
+        for(std::size_t i = 0; i < arg.group_count_; i++)
+        {
+            // TODO: Check if tensor specialization & strides mismatch
+            const auto& kernel_arg = arg.group_kernel_args_[i];
+            const auto& device_arg = arg.group_device_args_[i];
+            // Check if C permute dimension matches GEMM + GEMM shape
+            const index_t c_g       = device_arg.c_grid_desc_g_m_n_.GetLength(I0); // unpadded
+            const index_t c_m       = kernel_arg.y_grid_desc_m_o_.GetLength(I0);
+            const index_t c_gemm1n  = kernel_arg.y_grid_desc_m_o_.GetLength(I1);
+            const index_t a_m       = kernel_arg.a_grid_desc_ak0_m_ak1_.GetLength(I1);
+            const index_t b1_gemm1n = kernel_arg.b1_grid_desc_bk0_n_bk1_.GetLength(I0) * kernel_arg.b1_grid_desc_bk0_n_bk1_.GetLength(I2);
+
+            if(!(c_g == device_arg.batch_count_ && c_m == a_m && c_gemm1n == b1_gemm1n))
+            {
+                return false;
+            }
+
+            // Note: we need raw lengths since threadwise copy can not handle vector load when part
+            // of vector is out of bounds Note: need lowest dim in Ms/Ns/Ks/Os, not merged M/N/K/O
+            const auto MzRaw      = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[0];
+            const auto NzRaw      = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[1];
+            const auto KzRaw      = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[2];
+            const auto Gemm1NzRaw = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[3];
+
+            // Check scalar per vector requirement
+            const auto a_extent_lowest  = ABlockTransferSrcVectorDim == 2 ? KzRaw : MzRaw;
+            const auto b_extent_lowest  = BBlockTransferSrcVectorDim == 2 ? KzRaw : NzRaw;
+            const auto b1_extent_lowest = B1BlockTransferSrcVectorDim == 2 ? NzRaw : Gemm1NzRaw;
+            const auto c_extent_lowest  = Gemm1NzRaw;
+
+            if(!(a_extent_lowest % ABlockTransferSrcScalarPerVector == 0 &&
+                 b_extent_lowest % BBlockTransferSrcScalarPerVector == 0 &&
+                 b1_extent_lowest % B1BlockTransferSrcScalarPerVector == 0 &&
+                 c_extent_lowest % CShuffleBlockTransferScalarPerVector_NPerBlock == 0))
+            {
+                return false;
+            }
+
+            // Check vector load/store requirement
+            const auto a_stride_lowest = ABlockTransferSrcVectorDim == 2
+                                             ? device_arg.a_mz_kz_strides_[1]
+                                             : device_arg.a_mz_kz_strides_[0];
+            const auto b_stride_lowest = BBlockTransferSrcVectorDim == 2
+                                             ? device_arg.b_nz_kz_strides_[1]
+                                             : device_arg.b_nz_kz_strides_[0];
+            const auto b1_stride_lowest = B1BlockTransferSrcVectorDim == 2
+                                              ? device_arg.b1_nz_kz_strides_[1]
+                                              : device_arg.b1_nz_kz_strides_[0];
+            const auto c_stride_lowest =
+                device_arg.c_mz_gemm1nz_strides_[1]; // cshuffle assumes lowest dim in Gemm1Ns to be
+                                                     // contiguous
+
+            if(!(a_stride_lowest == 1 || b_stride_lowest == 1 || b1_stride_lowest == 1 ||
+                 c_stride_lowest == 1))
+            {
+                return false;
+            }
+
+            if(!GridwiseGemm::CheckValidity(kernel_arg.a_grid_desc_ak0_m_ak1_,
+                                            kernel_arg.b_grid_desc_bk0_n_bk1_,
+                                            kernel_arg.b1_grid_desc_bk0_n_bk1_,
+                                            kernel_arg.y_grid_desc_m_o_,
+                                            kernel_arg.block_2_ctile_map_))
+            {
+                return false;
+            }
+        }
+        return true;
+    }
+
+    // polymorphic
+    bool IsSupportedArgument(const BaseArgument* p_arg) override
+    {
+        return IsSupportedArgument(*dynamic_cast<const Argument*>(p_arg));
+    }
+
+    size_t GetWorkSpaceSize(const BaseArgument* p_arg) const override
+    {
+        return dynamic_cast<const Argument*>(p_arg)->group_count_ * sizeof(GroupKernelArg);
+    }
+
+    static auto MakeArgument(const std::vector<const void*>& p_As,
+                             const std::vector<const void*>& p_Bs,
+                             const std::vector<void*>& p_Zs,
+                             const std::vector<const void*>& p_B1s,
+                             const std::vector<const void*>& p_Cs, // for dS
+                             const std::vector<const void*>& p_LSEs,
+                             const std::vector<const void*>& p_Ygrads,
+                             std::vector<void*>& p_Qgrads,
+                             std::vector<void*>& p_Kgrads,
+                             std::vector<void*>& p_Vgrads,
+                             const std::array<void*, NumAcc0Bias>& p_acc0_biases,
+                             const std::array<void*, NumAcc1Bias>& p_acc1_biases,
+                             const std::vector<ProblemDesc>& problem_desc_vec,
+                             AElementwiseOperation a_element_op,
+                             BElementwiseOperation b_element_op,
+                             AccElementwiseOperation acc_element_op,
+                             B1ElementwiseOperation b1_element_op,
+                             CElementwiseOperation c_element_op,
+                             float p_drop,
+                             std::tuple<unsigned long long, unsigned long long> seeds)
+    {
+        return Argument{p_As,
+                        p_Bs,
+                        p_Zs,
+                        p_B1s,
+                        p_Cs,
+                        p_LSEs,
+                        p_Ygrads,
+                        p_Qgrads,
+                        p_Kgrads,
+                        p_Vgrads,
+                        p_acc0_biases,
+                        p_acc1_biases,
+                        problem_desc_vec,
+                        a_element_op,
+                        b_element_op,
+                        acc_element_op,
+                        b1_element_op,
+                        c_element_op,
+                        p_drop,
+                        seeds};
+    }
+
+    static auto MakeInvoker() { return Invoker{}; }
+
+    // polymorphic
+    // FIXME: constness
+    std::unique_ptr<BaseArgument>
+    MakeArgumentPointer(const std::vector<const void*>& p_As,
+                        const std::vector<const void*>& p_Bs,
+                        const std::vector<void*>& p_Zs,
+                        const std::vector<const void*>& p_B1s,
+                        const std::vector<const void*>& p_Cs, // for dS
+                        const std::vector<const void*>& p_LSEs,
+                        const std::vector<const void*>& p_Ygrads,
+                        std::vector<void*>& p_Qgrads,
+                        std::vector<void*>& p_Kgrads,
+                        std::vector<void*>& p_Vgrads,
+                        const std::array<void*, NumAcc0Bias>& p_acc0_biases,
+                        const std::array<void*, NumAcc1Bias>& p_acc1_biases,
+                        const std::vector<ProblemDesc>& problem_desc_vec,
+                        AElementwiseOperation a_element_op,
+                        BElementwiseOperation b_element_op,
+                        AccElementwiseOperation acc_element_op,
+                        B1ElementwiseOperation b1_element_op,
+                        CElementwiseOperation c_element_op,
+                        float p_drop,
+                        std::tuple<unsigned long long, unsigned long long> seeds) // override
+    {
+        return std::make_unique<Argument>(p_As,
+                                          p_Bs,
+                                          p_Zs,
+                                          p_B1s,
+                                          p_Cs,
+                                          p_LSEs,
+                                          p_Ygrads,
+                                          p_Qgrads,
+                                          p_Kgrads,
+                                          p_Vgrads,
+                                          p_acc0_biases, // cast in struct Argument
+                                          p_acc1_biases, // cast in struct Argument
+                                          problem_desc_vec,
+                                          a_element_op,
+                                          b_element_op,
+                                          acc_element_op,
+                                          b1_element_op,
+                                          c_element_op,
+                                          p_drop,
+                                          seeds);
+    }
+
+    // polymorphic
+    std::unique_ptr<BaseInvoker> MakeInvokerPointer() // override
+    {
+        return std::make_unique<Invoker>(Invoker{});
+    }
+
+    // polymorphic
+    std::string GetTypeString() const override
+    {
+        auto str = std::stringstream();
+
+        // clang-format off
+        str << "DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_PT1"
+            << "<"
+            << BlockSize << ", "
+            << MPerBlock << ", "
+            << NPerBlock << ", "
+            << KPerBlock << ", "
+            << AK1 << ", "
+            << BK1 << ", "
+            << MPerBlock << ", "
+            << Gemm1NPerBlock << ", "
+            << Gemm1KPerBlock << ", "
+            << B1K1 << ", "
+            << getGemmSpecializationString(GemmSpec) << ", "
+            << "ASpec" << getTensorSpecializationString(ASpec) << ", "
+            << "B0Spec" << getTensorSpecializationString(BSpec) << ", "
+            << "B1Spec" << getTensorSpecializationString(B1Spec) << ", "
+            << "CSpec" << getTensorSpecializationString(CSpec) << ", "
+            << getMaskingSpecializationString(MaskingSpec) << ">";
+        // clang-format on
+
+        return str.str();
+    }
+};
+
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

include/ck/tensor_operation/gpu/device/impl/device_grouped_multihead_attention_backward_xdl_cshuffle_v2.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include <iostream>
+#include <sstream>
+
+#include "ck/utility/common_header.hpp"
+#include "ck/utility/philox_rand.hpp"
+#include "ck/tensor_description/tensor_descriptor.hpp"
+#include "ck/tensor_description/tensor_descriptor_helper.hpp"
+// #include "ck/tensor_operation/gpu/device/device_batched_multihead_attention_backward.hpp" // TODO
+#include "ck/tensor_operation/gpu/device/device_base.hpp"
+#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
+#include "ck/tensor_operation/gpu/device/masking_specialization.hpp"
+#include "ck/tensor_operation/gpu/device/matrix_padder.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/tensor_operation/gpu/grid/gridwise_batched_multihead_attention_backward_xdl_cshuffle_v2.hpp"
+#include "ck/tensor_operation/operator_transform/transform_contraction_to_gemm.hpp"
+#include "ck/host_utility/device_prop.hpp"
+#include "ck/host_utility/kernel_launch.hpp"
+
+#include "ck/library/utility/host_tensor.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+
+template <typename GridwiseGemm,
+          typename GroupKernelArg,
+          typename AElementwiseOperation,
+          typename BElementwiseOperation,
+          typename AccElementwiseOperation,
+          typename B1ElementwiseOperation,
+          typename CElementwiseOperation,
+          bool HasMainKBlockLoop>
+__global__ void
+#if CK_USE_LAUNCH_BOUNDS
+    __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
+#endif
+        kernel_grouped_multihead_attention_backward_xdl_cshuffle_v2(
+            const void CK_CONSTANT_ADDRESS_SPACE* group_kernel_args,
+            const index_t group_count,
+            const AElementwiseOperation a_element_op,
+            const BElementwiseOperation b_element_op,
+            const AccElementwiseOperation acc_element_op,
+            const B1ElementwiseOperation b1_element_op,
+            const CElementwiseOperation c_element_op,
+            const float p_dropout,
+            const unsigned long long seed,
+            const unsigned long long offset)
+{
+#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx908__) || defined(__gfx90a__))
+    __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
+    const index_t block_id = get_block_1d_id();
+    const auto arg_ptr     = reinterpret_cast<const GroupKernelArg*>(
+        cast_pointer_to_generic_address_space(group_kernel_args));
+
+    index_t left     = 0;
+    index_t right    = group_count;
+    index_t group_id = index_t((left + right) / 2);
+
+    while(
+        (!(block_id >= arg_ptr[group_id].block_start_ && block_id < arg_ptr[group_id].block_end_)))
+    {
+        if(block_id < arg_ptr[group_id].block_start_)
+        {
+            right = group_id;
+        }
+        else
+        {
+            left = group_id;
+        }
+        group_id = index_t((left + right) / 2);
+    }
+
+    // per-group batch offset
+    const index_t num_blocks_per_batch = arg_ptr[group_id].num_blocks_per_batch_;
+    const index_t g_idx                = __builtin_amdgcn_readfirstlane(
+        (block_id - arg_ptr[group_id].block_start_) / num_blocks_per_batch);
+
+    const long_index_t a_batch_offset = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetABasePtr(g_idx)));
+    const long_index_t b_batch_offset = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetBBasePtr(g_idx)));
+    const long_index_t z_batch_offset = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetZBasePtr(g_idx)));
+    const long_index_t b1_batch_offset = __builtin_amdgcn_readfirstlane(static_cast<long_index_t>(
+        arg_ptr[group_id].compute_base_ptr_of_batch_.GetB1BasePtr(g_idx)));
+    const long_index_t c_batch_offset  = __builtin_amdgcn_readfirstlane(
+        static_cast<long_index_t>(arg_ptr[group_id].compute_base_ptr_of_batch_.GetCBasePtr(g_idx)));
+    const long_index_t lse_batch_offset = __builtin_amdgcn_readfirstlane(static_cast<long_index_t>(
+        arg_ptr[group_id].compute_base_ptr_of_batch_.GetLSEBasePtr(g_idx)));
+
+    const index_t global_thread_id = get_thread_global_1d_id();
+    ck::philox ph(seed, global_thread_id, offset);
+    unsigned short* z_matrix_ptr =
+        (arg_ptr[group_id].p_z_grid_ == nullptr ? nullptr
+                                                : arg_ptr[group_id].p_z_grid_ + z_batch_offset);
+    
+    GridwiseGemm::template Run<HasMainKBlockLoop>(
+        arg_ptr[group_id].p_a_grid_ + a_batch_offset,
+        arg_ptr[group_id].p_b_grid_ + b_batch_offset,
+        z_matrix_ptr,
+        arg_ptr[group_id].p_b1_grid_ + b1_batch_offset,
+        arg_ptr[group_id].p_c_grid_ + c_batch_offset,
+        arg_ptr[group_id].p_lse_grid_ + lse_batch_offset,
+        arg_ptr[group_id].p_ygrad_grid_ + c_batch_offset,
+        arg_ptr[group_id].p_qgrad_grid_ + a_batch_offset,
+        arg_ptr[group_id].p_kgrad_grid_ + b_batch_offset,
+        arg_ptr[group_id].p_vgrad_grid_ + b1_batch_offset,
+        p_shared,
+        a_element_op,
+        b_element_op,
+        acc_element_op,
+        b1_element_op,
+        c_element_op,
+        arg_ptr[group_id].a_grid_desc_ak0_m_ak1_,
+        arg_ptr[group_id].b_grid_desc_bk0_n_bk1_,
+        arg_ptr[group_id].c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5_,
+        arg_ptr[group_id].b1_grid_desc_bk0_n_bk1_,
+        arg_ptr[group_id].y_grid_desc_mblock_mperblock_oblock_operblock_,
+        arg_ptr[group_id].lse_grid_desc_m_,
+        arg_ptr[group_id].vgrad_grid_desc_n_o_,
+        arg_ptr[group_id].ygrad_grid_desc_m0_o_m1_,
+        arg_ptr[group_id].block_2_ctile_map_,
+        arg_ptr[group_id].c0_matrix_mask_,
+        p_dropout,
+        ph);
+#else
+    ignore = group_kernel_args;
+    ignore = group_count;
+    ignore = a_element_op;
+    ignore = b_element_op;
+    ignore = acc_element_op;
+    ignore = b1_element_op;
+    ignore = c_element_op;
+    ignore = p_dropout;
+    ignore = seed;
+    ignore = offset;
+#endif // end of if (defined(__gfx908__) || defined(__gfx90a__))
+}
+
+// Computes C = A * B0 * B1
+//              ^^^^^^ (Acc0)
+//              ^^^^^^^^^^^ (Acc1)
+template <index_t NumDimG,
+          index_t NumDimM,
+          index_t NumDimN,
+          index_t NumDimK,
+          index_t NumDimO, // NumDimGemm1N
+          typename DataType,
+          typename GemmDataType,
+          typename ZDataType,
+          typename LSEDataType,
+          typename Acc0BiasDataType,
+          typename Acc1BiasDataType,
+          typename GemmAccDataType,
+          typename CShuffleDataType,
+          typename AElementwiseOperation,
+          typename BElementwiseOperation,
+          typename AccElementwiseOperation,
+          typename B1ElementwiseOperation,
+          typename CElementwiseOperation,
+          GemmSpecialization GemmSpec,
+          TensorSpecialization ASpec,
+          TensorSpecialization BSpec,
+          TensorSpecialization B1Spec,
+          TensorSpecialization CSpec,
+          index_t NumGemmKPrefetchStage,
+          index_t BlockSize,
+          index_t MPerBlock,
+          index_t NPerBlock, // Gemm0NPerBlock
+          index_t KPerBlock, // Gemm0KPerBlock
+          index_t Gemm1NPerBlock,
+          index_t Gemm1KPerBlock,
+          index_t AK1,
+          index_t BK1,
+          index_t B1K1,
+          index_t MPerXDL,
+          index_t NPerXDL,
+          index_t MXdlPerWave,
+          index_t NXdlPerWave,
+          index_t Gemm1NXdlPerWave,
+          index_t Gemm2NXdlPerWave,
+          typename ABlockTransferThreadClusterLengths_AK0_M_AK1,
+          typename ABlockTransferThreadClusterArrangeOrder,
+          typename ABlockTransferSrcAccessOrder,
+          index_t ABlockTransferSrcVectorDim,
+          index_t ABlockTransferSrcScalarPerVector,
+          index_t ABlockTransferDstScalarPerVector_AK1,
+          bool ABlockLdsExtraM,
+          typename BBlockTransferThreadClusterLengths_BK0_N_BK1,
+          typename BBlockTransferThreadClusterArrangeOrder,
+          typename BBlockTransferSrcAccessOrder,
+          index_t BBlockTransferSrcVectorDim,
+          index_t BBlockTransferSrcScalarPerVector,
+          index_t BBlockTransferDstScalarPerVector_BK1,
+          bool BBlockLdsExtraN,
+          typename B1BlockTransferThreadClusterLengths_BK0_N_BK1,
+          typename B1BlockTransferThreadClusterArrangeOrder,
+          typename B1BlockTransferSrcAccessOrder,
+          index_t B1BlockTransferSrcVectorDim,
+          index_t B1BlockTransferSrcScalarPerVector,
+          index_t B1BlockTransferDstScalarPerVector_BK1,
+          bool B1BlockLdsExtraN,
+          index_t CShuffleMXdlPerWavePerShuffle,
+          index_t CShuffleNXdlPerWavePerShuffle,
+          typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
+          index_t CShuffleBlockTransferScalarPerVector_NPerBlock,
+          MaskingSpecialization MaskingSpec,
+          LoopScheduler LoopSched = LoopScheduler::Default>
+struct DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_V2
+    : public BaseOperator // TODO inherit atten bwd op once API stablizes
+{
+    static_assert(NumDimG > 0 && NumDimM > 0 && NumDimN > 0 && NumDimK > 0 && NumDimO > 0,
+                  "Number of dimension must be greater than 0");
+
+    static constexpr index_t NumAcc0Bias = Acc0BiasDataType::Size();
+    static constexpr index_t NumAcc1Bias = Acc1BiasDataType::Size();
+
+    // TODO: implement bias combination
+    static_assert(NumAcc0Bias == 0 && NumAcc0Bias == 0, "Bias addition is unimplemented");
+
+    using DeviceOp = DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_V2;
+    struct ProblemDesc
+    {
+        std::vector<index_t> a_gs_ms_ks_lengths;
+        std::vector<index_t> a_gs_ms_ks_strides;
+
+        std::vector<index_t> b_gs_ns_ks_lengths;
+        std::vector<index_t> b_gs_ns_ks_strides;
+
+        std::vector<index_t> z_gs_ms_ns_lengths;
+        std::vector<index_t> z_gs_ms_ns_strides;
+
+        std::vector<index_t> b1_gs_gemm1ns_gemm1ks_lengths;
+        std::vector<index_t> b1_gs_gemm1ns_gemm1ks_strides;
+
+        std::vector<index_t> c_gs_ms_gemm1ns_lengths;
+        std::vector<index_t> c_gs_ms_gemm1ns_strides;
+
+        std::vector<index_t> lse_gs_ms_lengths;
+        std::vector<index_t> lse_gs_ms_strides;
+
+        std::vector<std::vector<index_t>> acc0_biases_gs_ms_ns_lengths;
+        std::vector<std::vector<index_t>> acc0_biases_gs_ms_ns_strides;
+
+        std::vector<std::vector<index_t>> acc1_biases_gs_ms_os_lengths;
+        std::vector<std::vector<index_t>> acc1_biases_gs_ms_os_strides;
+    };
+    static constexpr auto I0 = Number<0>{};
+    static constexpr auto I1 = Number<1>{};
+    static constexpr auto I2 = Number<2>{};
+
+    static constexpr index_t Q_K1 = 8;
+    static constexpr index_t K_K1 = 8;
+    static constexpr index_t V_N1 = 2;
+
+    static constexpr index_t Q_M1 = 2;
+    static constexpr index_t K_N1 = 2;
+    static constexpr index_t V_O1 = 8;
+    static constexpr index_t Y_O1 = 8;
+    static constexpr index_t Y_M1 = 2;
+
+    static constexpr auto padder = GemmGemmPadder<GemmSpec,
+                                                  Number<MPerBlock>,
+                                                  Number<NPerBlock>,
+                                                  Number<KPerBlock>,
+                                                  Number<Gemm1NPerBlock>>{};
+
+    using Transform = TransformBatchedContractionContractionToBatchedGemmGemm<
+        Sequence<NumDimG, NumDimM, NumDimN, NumDimK, NumDimO>,
+        Sequence<MPerBlock, NPerBlock, KPerBlock, Gemm1NPerBlock>,
+        GemmSpec,
+        ASpec,
+        BSpec,
+        B1Spec,
+        CSpec>;
+
+    /*
+    Descriptors for inputs:
+
+      Q, K, V, Y, dY, per-row softmax stats
+
+    Descriptors for outputs:
+
+      dQ, dK, dV
+
+    */
+
+    // Q in Gemm A position
+    static auto MakeAGridDescriptor_AK0_M_AK1(const std::vector<index_t>& a_gs_ms_ks_lengths_vec,
+                                              const std::vector<index_t>& a_gs_ms_ks_strides_vec)
+    {
+        return Transform::MakeAGridDescriptor_AK0_M_AK1(
+            Transform::MakeAGridDescriptor_M_K(a_gs_ms_ks_lengths_vec, a_gs_ms_ks_strides_vec),
+            Number<AK1>{});
+    }
+
+    // K in Gemm B0 position
+    static auto MakeBGridDescriptor_BK0_N_BK1(const std::vector<index_t>& b_gs_ns_ks_lengths_vec,
+                                              const std::vector<index_t>& b_gs_ns_ks_strides_vec)
+    {
+        return Transform::MakeB0GridDescriptor_BK0_N_BK1(
+            Transform::MakeB0GridDescriptor_N_K(b_gs_ns_ks_lengths_vec, b_gs_ns_ks_strides_vec),
+            Number<BK1>{});
+    }
+
+    // V in Gemm B1 position
+    static auto
+    MakeB1GridDescriptor_BK0_N_BK1(const std::vector<index_t>& b1_gs_gemm1ns_gemm1ks_lengths_vec,
+                                   const std::vector<index_t>& b1_gs_gemm1ns_gemm1ks_strides_vec)
+    {
+        return Transform::MakeB1GridDescriptor_BK0_N_BK1(
+            Transform::MakeB1GridDescriptor_N_K(b1_gs_gemm1ns_gemm1ks_lengths_vec,
+                                                b1_gs_gemm1ns_gemm1ks_strides_vec),
+            Number<B1K1>{});
+    }
+
+    //
+    // dV = P^T * dY
+    //
+
+    // VGrad in Gemm C position
+    static auto MakeVGradGridDescriptor_N_O(const std::vector<index_t>& v_gs_os_ns_lengths_vec,
+                                            const std::vector<index_t>& v_gs_os_ns_strides_vec)
+    {
+        // v_gs_os_ns -> vgrad_gs_ns_os. O dims last because output is row-major.
+        // Here directly rearrange lengths/strides before constructing tensor descriptor to reduce
+        // transformation overhead
+        // TODO: This will be much easier when inputs are Gs, Ms, Ns, Os. So there's no need to
+        // extract subsequence and shuffle them.
+        const index_t num_dims = NumDimG + NumDimN + NumDimO;
+
+        // 0, 1, .. NumDimG - 1
+        std::vector<index_t> gs_ids(NumDimG);
+        std::iota(gs_ids.begin(), gs_ids.end(), 0);
+
+        // NumDimG, NumDimG + 1, ... NumDimG + NumDimO - 1
+        std::vector<index_t> os_ids(NumDimO);
+        std::iota(os_ids.begin(), os_ids.end(), NumDimG);
+
+        // NumDimG + NumDimO, NumDimG + NumDimO + 1, ... NumDimG + NumDimO + NumDimN - 1
+        std::vector<index_t> ns_ids(NumDimN);
+        std::iota(ns_ids.begin(), ns_ids.end(), NumDimG + NumDimO);
+
+        std::vector<index_t> ids_old2new;
+        ids_old2new.insert(ids_old2new.end(), gs_ids.begin(), gs_ids.end());
+        ids_old2new.insert(ids_old2new.end(), ns_ids.begin(), ns_ids.end());
+        ids_old2new.insert(ids_old2new.end(), os_ids.begin(), os_ids.end());
+
+        std::vector<index_t> v_gs_ns_os_lengths_vec(num_dims), v_gs_ns_os_strides_vec(num_dims);
+        for(int i = 0; i < num_dims; i++)
+        {
+            index_t id_new            = ids_old2new[i];
+            v_gs_ns_os_lengths_vec[i] = v_gs_os_ns_lengths_vec[id_new];
+            v_gs_ns_os_strides_vec[i] = v_gs_os_ns_strides_vec[id_new];
+        }
+
+        const auto vgrad_desc_nraw_oraw =
+            MakeGridDescriptorPair<NumDimG, NumDimN, NumDimO, TensorSpecialization::Default>(
+                v_gs_ns_os_lengths_vec, v_gs_ns_os_strides_vec)
+                .second;
+
+        return PadTensorDescriptor(vgrad_desc_nraw_oraw,
+                                   make_tuple(NPerBlock, Gemm1NPerBlock),
+                                   Sequence<padder.PadN, padder.PadO>{});
+    }
+
+    template <typename YGridDesc_M_O>
+    static auto MakeYGradGridDescriptor_M0_O_M1(const YGridDesc_M_O& ygrad_grid_desc_m_o)
+    {
+        const auto M = ygrad_grid_desc_m_o.GetLength(I0);
+        const auto O = ygrad_grid_desc_m_o.GetLength(I1);
+
+        const auto Y_M0 = M / Y_M1;
+
+        return transform_tensor_descriptor(
+            ygrad_grid_desc_m_o,
+            make_tuple(make_unmerge_transform(make_tuple(Y_M0, Y_M1)),
+                       make_pass_through_transform(O)),
+            make_tuple(Sequence<0>{}, Sequence<1>{}),
+            make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
+    }
+    //
+    // dS_i_j = P_i_j .* (dP_i_j - dY_i dot Y_i)
+    //
+
+    //
+    // dQ = alpha * dS * K
+    //
+
+    // QGrad in Gemm C position
+    static auto MakeQGradGridDescriptor_M_K(const std::vector<index_t>& q_gs_ms_ks_lengths_vec,
+                                            const std::vector<index_t>& q_gs_ms_ks_strides_vec)
+    {
+        return Transform::MakeCGridDescriptor_M_N(q_gs_ms_ks_lengths_vec, q_gs_ms_ks_strides_vec);
+    }
+
+    //
+    // dK = alpha * dS^T * Q
+    //
+
+    // KGrad in Gemm C position
+    static auto MakeKGradGridDescriptor_N_K(const std::vector<index_t>& k_gs_ns_ks_lengths_vec,
+                                            const std::vector<index_t>& k_gs_ns_ks_strides_vec)
+    {
+        return Transform::MakeCGridDescriptor_M_N(k_gs_ns_ks_lengths_vec, k_gs_ns_ks_strides_vec);
+    }
+
+    static auto MakeZGridDescriptor_M_N(const std::vector<index_t>& z_gs_ms_ns_lengths_vec,
+                                        const std::vector<index_t>& z_gs_ms_ns_strides_vec)
+    {
+        return Transform::MakeCGridDescriptor_M_N(z_gs_ms_ns_lengths_vec, z_gs_ms_ns_strides_vec);
+    }
+
+    static auto MakeLSEGridDescriptor_M(index_t MRaw)
+    {
+        const auto lse_grid_desc_mraw = make_naive_tensor_descriptor_packed(make_tuple(MRaw));
+
+        const auto M    = math::integer_divide_ceil(MRaw, MPerBlock) * MPerBlock;
+        const auto MPad = M - MRaw;
+
+        if constexpr(GemmSpec == GemmSpecialization::MPadding ||
+                     GemmSpec == GemmSpecialization::MNPadding ||
+                     GemmSpec == GemmSpecialization::MKPadding ||
+                     GemmSpec == GemmSpecialization::MNKPadding)
+        {
+            // pad M
+            return transform_tensor_descriptor(lse_grid_desc_mraw,
+                                               make_tuple(make_right_pad_transform(MRaw, MPad)),
+                                               make_tuple(Sequence<0>{}),
+                                               make_tuple(Sequence<0>{}));
+        }
+        else
+        {
+            // not pad M
+            return lse_grid_desc_mraw;
+        }
+    }
+
+    using AGridDesc_AK0_M_AK1  = decltype(MakeAGridDescriptor_AK0_M_AK1({}, {}));
+    using BGridDesc_BK0_N_BK1  = decltype(MakeBGridDescriptor_BK0_N_BK1({}, {}));
+    using B1GridDesc_BK0_N_BK1 = decltype(MakeB1GridDescriptor_BK0_N_BK1({}, {}));
+    using YGridDesc_M_O        = decltype(Transform::MakeCGridDescriptor_M_N({}, {}));
+    using LSEGridDesc_M        = decltype(MakeLSEGridDescriptor_M(1));
+    using AGridDesc_G_M_K      = decltype(Transform::MakeAGridDescriptor_G_M_K({}, {}));
+    using BGridDesc_G_N_K      = decltype(Transform::MakeB0GridDescriptor_G_N_K({}, {}));
+    using B1GridDesc_G_N_K     = decltype(Transform::MakeB1GridDescriptor_G_N_K({}, {}));
+    using CGridDesc_G_M_N      = decltype(Transform::MakeCGridDescriptor_G_M_N({}, {}));
+    using ZGridDesc_G_M_N      = decltype(Transform::MakeCGridDescriptor_G_M_N({}, {}));
+
+    using VGradGridDesc_N_O     = decltype(MakeVGradGridDescriptor_N_O({}, {}));
+    using YGradGridDesc_M0_O_M1 = decltype(MakeYGradGridDescriptor_M0_O_M1(YGridDesc_M_O{}));
+    using ZGridDesc_M_N         = decltype(MakeZGridDescriptor_M_N({}, {}));
+
+    constexpr static auto make_MaskOutPredicate()
+    {
+        if constexpr(MaskingSpec == MaskingSpecialization::MaskDisabled)
+        {
+            return MaskDisabledPredicate{};
+        }
+        else if constexpr(MaskingSpec == MaskingSpecialization::MaskOutUpperTriangle)
+        {
+            return MaskOutUpperTrianglePredicate{};
+        }
+    }
+    using C0MatrixMask = C0MatrixMask_impl<decltype(make_MaskOutPredicate())>;
+
+    struct ComputeBasePtrOfStridedBatch
+    {
+        ComputeBasePtrOfStridedBatch(const AGridDesc_G_M_K& a_grid_desc_g_m_k,
+                                     const BGridDesc_G_N_K& b_grid_desc_g_n_k,
+                                     const ZGridDesc_G_M_N& z_grid_desc_g_m_n,
+                                     const B1GridDesc_G_N_K& b1_grid_desc_g_n_k,
+                                     const CGridDesc_G_M_N& c_grid_desc_g_m_n,
+                                     index_t BatchStrideLSE)
+            : a_grid_desc_g_m_k_(a_grid_desc_g_m_k),
+              b_grid_desc_g_n_k_(b_grid_desc_g_n_k),
+              z_grid_desc_g_m_n_(z_grid_desc_g_m_n),
+              b1_grid_desc_g_n_k_(b1_grid_desc_g_n_k),
+              c_grid_desc_g_m_n_(c_grid_desc_g_m_n),
+              BatchStrideLSE_(BatchStrideLSE)
+        {
+        }
+
+        __host__ __device__ constexpr long_index_t GetABasePtr(index_t g_idx) const
+        {
+            return a_grid_desc_g_m_k_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetBBasePtr(index_t g_idx) const
+        {
+            return b_grid_desc_g_n_k_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetZBasePtr(index_t g_idx) const
+        {
+            return z_grid_desc_g_m_n_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetB1BasePtr(index_t g_idx) const
+        {
+            return b1_grid_desc_g_n_k_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetCBasePtr(index_t g_idx) const
+        {
+            return c_grid_desc_g_m_n_.CalculateOffset(make_multi_index(g_idx, 0, 0));
+        }
+
+        __host__ __device__ constexpr long_index_t GetLSEBasePtr(index_t g_idx) const
+        {
+            return g_idx * static_cast<long_index_t>(BatchStrideLSE_);
+        }
+
+        private:
+        AGridDesc_G_M_K a_grid_desc_g_m_k_;
+        BGridDesc_G_N_K b_grid_desc_g_n_k_;
+        ZGridDesc_G_M_N z_grid_desc_g_m_n_;
+        B1GridDesc_G_N_K b1_grid_desc_g_n_k_;
+        CGridDesc_G_M_N c_grid_desc_g_m_n_;
+        index_t BatchStrideLSE_;
+    };
+
+    // GridwiseGemm
+    using GridwiseGemm = GridwiseBatchedMultiheadAttentionBackward_Xdl_CShuffle_V2<
+        DataType, // TODO: distinguish A/B datatype
+        GemmDataType,
+        GemmAccDataType,
+        CShuffleDataType,
+        LSEDataType,
+        AElementwiseOperation,
+        BElementwiseOperation,
+        AccElementwiseOperation,
+        B1ElementwiseOperation,
+        CElementwiseOperation,
+        InMemoryDataOperationEnum::Set,
+        AGridDesc_AK0_M_AK1,
+        BGridDesc_BK0_N_BK1,
+        ZGridDesc_M_N,
+        B1GridDesc_BK0_N_BK1,
+        YGridDesc_M_O,
+        LSEGridDesc_M,
+        NumGemmKPrefetchStage,
+        BlockSize,
+        MPerBlock,
+        NPerBlock,
+        KPerBlock,
+        Gemm1NPerBlock,
+        Gemm1KPerBlock,
+        AK1,
+        BK1,
+        B1K1,
+        MPerXDL,
+        NPerXDL,
+        MXdlPerWave,
+        NXdlPerWave,
+        Gemm1NXdlPerWave,
+        Gemm2NXdlPerWave,
+        ABlockTransferThreadClusterLengths_AK0_M_AK1,
+        ABlockTransferThreadClusterArrangeOrder,
+        ABlockTransferSrcAccessOrder,
+        ABlockTransferSrcVectorDim,
+        ABlockTransferSrcScalarPerVector,
+        ABlockTransferDstScalarPerVector_AK1,
+        true,
+        ABlockLdsExtraM,
+        BBlockTransferThreadClusterLengths_BK0_N_BK1,
+        BBlockTransferThreadClusterArrangeOrder,
+        BBlockTransferSrcAccessOrder,
+        BBlockTransferSrcVectorDim,
+        BBlockTransferSrcScalarPerVector,
+        BBlockTransferDstScalarPerVector_BK1,
+        true,
+        BBlockLdsExtraN,
+        B1BlockTransferThreadClusterLengths_BK0_N_BK1,
+        B1BlockTransferThreadClusterArrangeOrder,
+        B1BlockTransferSrcAccessOrder,
+        B1BlockTransferSrcVectorDim,
+        B1BlockTransferSrcScalarPerVector,
+        B1BlockTransferDstScalarPerVector_BK1,
+        false,
+        B1BlockLdsExtraN,
+        CShuffleMXdlPerWavePerShuffle,
+        CShuffleNXdlPerWavePerShuffle,
+        CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,
+        CShuffleBlockTransferScalarPerVector_NPerBlock,
+        LoopSched,
+        Transform::matrix_padder.PadN,
+        MaskingSpec == MaskingSpecialization::MaskOutUpperTriangle>;
+
+    using Block2CTileMap = OffsettedBlockToCTileMap<typename GridwiseGemm::DefaultBlock2CTileMap>;
+
+    struct GroupKernelArg
+    {
+        // pointers
+        const DataType* p_a_grid_;
+        const DataType* p_b_grid_;
+        ZDataType* p_z_grid_;
+        const DataType* p_b1_grid_;
+        const DataType* p_c_grid_;
+        const LSEDataType* p_lse_grid_;
+        const DataType* p_ygrad_grid_;
+        DataType* p_qgrad_grid_;
+        DataType* p_kgrad_grid_;
+        DataType* p_vgrad_grid_;
+
+        // tensor descriptors for block/thread-wise copy
+        AGridDesc_AK0_M_AK1 a_grid_desc_ak0_m_ak1_;
+        BGridDesc_BK0_N_BK1 b_grid_desc_bk0_n_bk1_;
+        ZGridDesc_M_N z_grid_desc_m_n_;
+        B1GridDesc_BK0_N_BK1 b1_grid_desc_bk0_n_bk1_;
+        YGridDesc_M_O y_grid_desc_m_o_;
+
+        typename GridwiseGemm::YGridDescriptor_MBlock_MPerBlock_OBlock_OPerBlock
+            y_grid_desc_mblock_mperblock_oblock_operblock_;
+        typename GridwiseGemm::ZGridDescriptor_M0_N0_M1_N1_M2_N2_M3_N3_N4_N5
+            c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5_;
+        LSEGridDesc_M lse_grid_desc_m_;
+        VGradGridDesc_N_O vgrad_grid_desc_n_o_;
+        YGradGridDesc_M0_O_M1 ygrad_grid_desc_m0_o_m1_;
+        // block-to-c-tile map
+        Block2CTileMap block_2_ctile_map_;
+        index_t num_blocks_per_batch_;
+        ComputeBasePtrOfStridedBatch compute_base_ptr_of_batch_;
+
+        // check C0 masking and padding
+        C0MatrixMask c0_matrix_mask_;
+        index_t block_start_, block_end_;
+    };
+
+    struct GroupDeviceArg
+    {
+        // lengths for the last dimensions of overall problem for sanity check of vector load/store
+        std::vector<index_t> raw_lengths_mz_nz_kz_gemm1nz_;
+
+        // strides for the last dimensions of each tensor for sanity check of vector load/store
+        std::vector<index_t> a_mz_kz_strides_;
+        std::vector<index_t> b_nz_kz_strides_;
+        std::vector<index_t> b1_nz_kz_strides_;
+        std::vector<index_t> c_mz_gemm1nz_strides_;
+
+        // for gridwise gemm check
+        CGridDesc_G_M_N c_grid_desc_g_m_n_;
+
+        index_t batch_count_;
+    };
+    // Argument
+    struct Argument : public BaseArgument
+    {
+        Argument(const std::vector<const void*>& p_As,
+                 const std::vector<const void*>& p_Bs,
+                 const std::vector<void*>& p_Zs,
+                 const std::vector<const void*>& p_B1s,
+                 const std::vector<const void*>& p_Cs, // for dS
+                 const std::vector<const void*>& p_LSEs,
+                 const std::vector<const void*>& p_Ygrads,
+                 std::vector<void*>& p_Qgrads,
+                 std::vector<void*>& p_Kgrads,
+                 std::vector<void*>& p_Vgrads,
+                 const std::array<void*, NumAcc0Bias>& p_acc0_biases,
+                 const std::array<void*, NumAcc1Bias>& p_acc1_biases,
+                 const std::vector<ProblemDesc>& problem_desc_vec,
+                 AElementwiseOperation a_element_op,
+                 BElementwiseOperation b_element_op,
+                 AccElementwiseOperation acc_element_op,
+                 B1ElementwiseOperation b1_element_op,
+                 CElementwiseOperation c_element_op,
+                 float p_drop,
+                 std::tuple<unsigned long long, unsigned long long> seeds)
+            : a_element_op_{a_element_op},
+              b_element_op_{b_element_op},
+              acc_element_op_{acc_element_op},
+              b1_element_op_{b1_element_op},
+              c_element_op_{c_element_op},
+              p_dropout_{p_drop}
+        {
+            seed_   = std::get<0>(seeds);
+            offset_ = std::get<1>(seeds);
+
+            group_count_ = ck::type_convert<ck::index_t>(problem_desc_vec.size());
+
+            if(!(group_count_ == ck::type_convert<ck::index_t>(p_As.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Bs.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Zs.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_B1s.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Cs.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Ygrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Qgrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Kgrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_Vgrads.size()) &&
+                 group_count_ == ck::type_convert<ck::index_t>(p_LSEs.size())))
+            {
+                throw std::runtime_error("wrong! group_count_ != p_As/b/b1/c.size");
+            }
+
+            if(!(p_acc0_biases.size() == p_acc1_biases.size()))
+            {
+                throw std::runtime_error("wrong! acc0_bias_vec.size != acc1_bias_vec.size");
+            }
+
+            grid_size_ = 0;
+            for(std::size_t i = 0; i < group_count_; i++)
+            {
+                const auto p_a_grid     = static_cast<const DataType*>(p_As[i]);
+                const auto p_b_grid     = static_cast<const DataType*>(p_Bs[i]);
+                auto p_z_grid           = static_cast<ZDataType*>(p_Zs[i]);
+                const auto p_b1_grid    = static_cast<const DataType*>(p_B1s[i]);
+                const auto p_c_grid     = static_cast<const DataType*>(p_Cs[i]);
+                const auto p_lse_grid   = static_cast<const LSEDataType*>(p_LSEs[i]);
+                const auto p_ygrad_grid = static_cast<const DataType*>(p_Ygrads[i]);
+                auto p_qgrad_grid       = static_cast<DataType*>(p_Qgrads[i]);
+                auto p_kgrad_grid       = static_cast<DataType*>(p_Kgrads[i]);
+                auto p_vgrad_grid       = static_cast<DataType*>(p_Vgrads[i]);
+
+                const auto& problem_desc = problem_desc_vec[i];
+
+                const auto a_grid_desc_ak0_m_ak1 = DeviceOp::MakeAGridDescriptor_AK0_M_AK1(
+                    problem_desc.a_gs_ms_ks_lengths, problem_desc.a_gs_ms_ks_strides);
+                const auto b_grid_desc_bk0_n_bk1 = DeviceOp::MakeBGridDescriptor_BK0_N_BK1(
+                    problem_desc.b_gs_ns_ks_lengths, problem_desc.b_gs_ns_ks_strides);
+                const auto z_grid_desc_m_n = DeviceOp::MakeZGridDescriptor_M_N(
+                    problem_desc.z_gs_ms_ns_lengths, problem_desc.z_gs_ms_ns_strides);
+                const auto b1_grid_desc_bk0_n_bk1 = DeviceOp::MakeB1GridDescriptor_BK0_N_BK1(
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_lengths,
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_strides);
+                const auto y_grid_desc_m_o = Transform::MakeCGridDescriptor_M_N(
+                    problem_desc.c_gs_ms_gemm1ns_lengths, problem_desc.c_gs_ms_gemm1ns_strides);
+
+                const auto lse_grid_desc_m =
+                    DeviceOp::MakeLSEGridDescriptor_M(problem_desc.lse_gs_ms_lengths[NumDimG]);
+                const auto vgrad_grid_desc_n_o = DeviceOp::MakeVGradGridDescriptor_N_O(
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_lengths,
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_strides);
+                const auto ygrad_grid_desc_m0_o_m1 =
+                    DeviceOp::MakeYGradGridDescriptor_M0_O_M1(y_grid_desc_m_o);
+
+                const auto a_grid_desc_g_m_k = Transform::MakeAGridDescriptor_G_M_K(
+                    problem_desc.a_gs_ms_ks_lengths, problem_desc.a_gs_ms_ks_strides);
+                const auto b_grid_desc_g_n_k = Transform::MakeB0GridDescriptor_G_N_K(
+                    problem_desc.b_gs_ns_ks_lengths, problem_desc.b_gs_ns_ks_strides);
+                const auto z_grid_desc_g_m_n = Transform::MakeCGridDescriptor_G_M_N(
+                    problem_desc.z_gs_ms_ns_lengths, problem_desc.z_gs_ms_ns_strides);
+                const auto b1_grid_desc_g_n_k = Transform::MakeB1GridDescriptor_G_N_K(
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_lengths,
+                    problem_desc.b1_gs_gemm1ns_gemm1ks_strides);
+                const auto c_grid_desc_g_m_n = Transform::MakeCGridDescriptor_G_M_N(
+                    problem_desc.c_gs_ms_gemm1ns_lengths, problem_desc.c_gs_ms_gemm1ns_strides);
+                typename GridwiseGemm::YGridDescriptor_MBlock_MPerBlock_OBlock_OPerBlock
+                    y_grid_desc_mblock_mperblock_oblock_operblock;
+                typename GridwiseGemm::ZGridDescriptor_M0_N0_M1_N1_M2_N2_M3_N3_N4_N5
+                    c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5;
+                const index_t BlockStart     = grid_size_;
+                const auto block_2_ctile_map = Block2CTileMap(y_grid_desc_m_o, BlockStart);
+                if(GridwiseGemm::CheckValidity(a_grid_desc_ak0_m_ak1,
+                                               b_grid_desc_bk0_n_bk1,
+                                               b1_grid_desc_bk0_n_bk1,
+                                               y_grid_desc_m_o,
+                                               block_2_ctile_map))
+                {
+                    y_grid_desc_mblock_mperblock_oblock_operblock =
+                        GridwiseGemm::MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(
+                            y_grid_desc_m_o);
+                }
+
+                c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5 =
+                    GridwiseGemm::MakeCGridDescriptor_M0_N0_M1_N1_M2_N2_M3_N3_N4_N5(
+                        z_grid_desc_m_n);
+
+                const index_t batch_count = c_grid_desc_g_m_n.GetLength(I0);
+                const index_t grid_size_grp =
+                    block_2_ctile_map.CalculateGridSize(y_grid_desc_m_o) * batch_count;
+                const index_t BlockEnd = grid_size_ + grid_size_grp;
+                // batch stride
+                const auto compute_base_ptr_of_batch = ComputeBasePtrOfStridedBatch(
+                    a_grid_desc_g_m_k,
+                    b_grid_desc_g_n_k,
+                    z_grid_desc_g_m_n,
+                    b1_grid_desc_g_n_k,
+                    c_grid_desc_g_m_n,
+                    type_convert<index_t>(lse_grid_desc_m.GetElementSpaceSize()));
+
+                // C0 mask
+                const auto c0_matrix_mask = C0MatrixMask(b_grid_desc_g_n_k.GetLength(I1));
+
+                grid_size_ += grid_size_grp;
+
+                // for each group, make sure acc0_biases_gs_ms_ns_lengths.size() == NumAcc0Bias and
+                // so on
+                if(!(problem_desc.acc0_biases_gs_ms_ns_lengths.size() == NumAcc0Bias &&
+                     problem_desc.acc0_biases_gs_ms_ns_strides.size() == NumAcc0Bias &&
+                     problem_desc.acc1_biases_gs_ms_os_lengths.size() == NumAcc1Bias &&
+                     problem_desc.acc1_biases_gs_ms_os_strides.size() == NumAcc1Bias))
+                {
+                    throw std::runtime_error(
+                        "wrong! number of biases in function argument does not "
+                        "match that in template argument");
+                }
+
+                group_kernel_args_.push_back({p_a_grid,
+                                              p_b_grid,
+                                              p_z_grid,
+                                              p_b1_grid,
+                                              p_c_grid,
+                                              p_lse_grid,
+                                              p_ygrad_grid,
+                                              p_qgrad_grid,
+                                              p_kgrad_grid,
+                                              p_vgrad_grid,
+                                              a_grid_desc_ak0_m_ak1,
+                                              b_grid_desc_bk0_n_bk1,
+                                              z_grid_desc_m_n,
+                                              b1_grid_desc_bk0_n_bk1,
+                                              y_grid_desc_m_o,
+                                              y_grid_desc_mblock_mperblock_oblock_operblock,
+                                              c_grid_desc_m0_n0_m1_n1_m2_n2_m3_n3_n4_n5,
+                                              lse_grid_desc_m,
+                                              vgrad_grid_desc_n_o,
+                                              ygrad_grid_desc_m0_o_m1,
+                                              block_2_ctile_map,
+                                              block_2_ctile_map.CalculateGridSize(y_grid_desc_m_o),
+                                              compute_base_ptr_of_batch,
+                                              c0_matrix_mask,
+                                              BlockStart,
+                                              BlockEnd});
+
+                group_device_args_.push_back(
+                    {{problem_desc.a_gs_ms_ks_lengths[NumDimG + NumDimM - 1],
+                      problem_desc.b_gs_ns_ks_lengths[NumDimG + NumDimN - 1],
+                      problem_desc.b_gs_ns_ks_lengths[NumDimG + NumDimN + NumDimK - 1],
+                      problem_desc.b1_gs_gemm1ns_gemm1ks_lengths[NumDimG + NumDimO - 1]},
+                     {problem_desc.a_gs_ms_ks_strides[NumDimG + NumDimM - 1],
+                      problem_desc.a_gs_ms_ks_strides[NumDimG + NumDimM + NumDimK - 1]},
+                     {problem_desc.b_gs_ns_ks_strides[NumDimG + NumDimN - 1],
+                      problem_desc.b_gs_ns_ks_strides[NumDimG + NumDimN + NumDimK - 1]},
+                     {problem_desc.b1_gs_gemm1ns_gemm1ks_strides[NumDimG + NumDimO - 1],
+                      problem_desc.b1_gs_gemm1ns_gemm1ks_strides[NumDimG + NumDimO + NumDimN - 1]},
+                     {problem_desc.c_gs_ms_gemm1ns_strides[NumDimG + NumDimM - 1],
+                      problem_desc.c_gs_ms_gemm1ns_strides[NumDimG + NumDimM + NumDimO - 1]},
+                     c_grid_desc_g_m_n,
+                     batch_count});
+            }
+            // TODO: implement bias addition
+            // ignore = p_acc0_biases;
+            // ignore = p_acc1_biases;
+            // ignore = acc0_biases_gs_ms_ns_lengths;
+            // ignore = acc0_biases_gs_ms_ns_strides;
+            // ignore = acc1_biases_gs_ms_gemm1ns_lengths;
+            // ignore = acc1_biases_gs_ms_gemm1ns_strides;
+        }
+
+        // element-wise op
+        AElementwiseOperation a_element_op_;
+        BElementwiseOperation b_element_op_;
+        AccElementwiseOperation acc_element_op_;
+        B1ElementwiseOperation b1_element_op_;
+        CElementwiseOperation c_element_op_;
+
+        float p_dropout_;
+        unsigned long long seed_;
+        unsigned long long offset_;
+
+        index_t grid_size_;
+        index_t group_count_;
+
+        std::vector<GroupKernelArg> group_kernel_args_;
+        std::vector<GroupDeviceArg> group_device_args_;
+    };
+
+    // Invoker
+    struct Invoker : public BaseInvoker
+    {
+        using Argument = DeviceOp::Argument;
+
+        float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{})
+        {
+            if(!DeviceOp::IsSupportedArgument(arg))
+            {
+                throw std::runtime_error("wrong! unsupported argument");
+            }
+
+            bool all_has_main_k_block_loop  = true;
+            bool some_has_main_k_block_loop = false;
+            for(std::size_t i = 0; i < arg.group_count_; i++)
+            {
+                const auto K =
+                    arg.group_kernel_args_[i].a_grid_desc_ak0_m_ak1_.GetLength(I0) * arg.group_kernel_args_[i].a_grid_desc_ak0_m_ak1_.GetLength(I2);
+                const bool y = GridwiseGemm::CalculateHasMainKBlockLoop(K);
+                all_has_main_k_block_loop &= y;
+                some_has_main_k_block_loop |= y;
+            }
+
+            hipGetErrorString(hipMemcpy(arg.p_workspace_,
+                                        arg.group_kernel_args_.data(),
+                                        arg.group_kernel_args_.size() * sizeof(GroupKernelArg),
+                                        hipMemcpyHostToDevice));
+
+            float ave_time = 0;
+
+            auto launch_kernel = [&](auto has_main_k_block_loop_) {
+                const auto kernel = kernel_grouped_multihead_attention_backward_xdl_cshuffle_v2<
+                    GridwiseGemm,
+                    GroupKernelArg,
+                    AElementwiseOperation,
+                    BElementwiseOperation,
+                    AccElementwiseOperation,
+                    B1ElementwiseOperation,
+                    CElementwiseOperation,
+                    has_main_k_block_loop_>;
+
+                return launch_and_time_kernel(
+                    stream_config,
+                    kernel,
+                    dim3(arg.grid_size_),
+                    dim3(BlockSize),
+                    0,
+                    cast_pointer_to_constant_address_space(arg.p_workspace_),
+                    arg.group_count_,
+                    arg.a_element_op_,
+                    arg.b_element_op_,
+                    arg.acc_element_op_,
+                    arg.b1_element_op_,
+                    arg.c_element_op_,
+                    arg.p_dropout_,
+                    arg.seed_,
+                    arg.offset_);
+            };
+
+            // Gemm1_K is split into Gemm1_K0/K1 where K1 is known at compile time, so we only need
+            // to concern Gemm0's loop
+            if(all_has_main_k_block_loop)
+            {
+                ave_time = launch_kernel(integral_constant<bool, true>{});
+            }
+            else if(!some_has_main_k_block_loop)
+            {
+                ave_time = launch_kernel(integral_constant<bool, false>{});
+            }
+            else
+            {
+                throw std::runtime_error("wrong! all gemm problems have to simultaneously meet "
+                                         "has_main_k_block_loop or no_main_k_block_loop");
+            }
+            return ave_time;
+        }
+
+        // polymorphic
+        float Run(const BaseArgument* p_arg,
+                  const StreamConfig& stream_config = StreamConfig{}) override
+        {
+            return Run(*dynamic_cast<const Argument*>(p_arg), stream_config);
+        }
+    };
+
+    static constexpr bool IsValidCompilationParameter()
+    {
+        // TODO: properly implement this check
+        return true;
+    }
+
+    static bool IsSupportedArgument(const Argument& arg)
+    {
+        if(!(ck::get_device_name() == "gfx908" || ck::get_device_name() == "gfx90a"))
+        {
+            return false;
+        }
+
+        for(std::size_t i = 0; i < arg.group_count_; i++)
+        {
+            // TODO: Check if tensor specialization & strides mismatch
+            const auto& kernel_arg = arg.group_kernel_args_[i];
+            const auto& device_arg = arg.group_device_args_[i];
+            // Check if C permute dimension matches GEMM + GEMM shape
+            const index_t c_g       = device_arg.c_grid_desc_g_m_n_.GetLength(I0); // unpadded
+            const index_t c_m       = kernel_arg.y_grid_desc_m_o_.GetLength(I0);
+            const index_t c_gemm1n  = kernel_arg.y_grid_desc_m_o_.GetLength(I1);
+            const index_t a_m       = kernel_arg.a_grid_desc_ak0_m_ak1_.GetLength(I1);
+            const index_t b1_gemm1n = kernel_arg.b1_grid_desc_bk0_n_bk1_.GetLength(I1);
+
+            if(!(c_g == device_arg.batch_count_ && c_m == a_m && c_gemm1n == b1_gemm1n))
+            {
+                return false;
+            }
+
+            // Note: we need raw lengths since threadwise copy can not handle vector load when part
+            // of vector is out of bounds Note: need lowest dim in Ms/Ns/Ks/Os, not merged M/N/K/O
+            const auto MzRaw      = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[0];
+            const auto NzRaw      = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[1];
+            const auto KzRaw      = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[2];
+            const auto Gemm1NzRaw = device_arg.raw_lengths_mz_nz_kz_gemm1nz_[3];
+
+            // Check scalar per vector requirement
+            const auto a_extent_lowest  = ABlockTransferSrcVectorDim == 2 ? KzRaw : MzRaw;
+            const auto b_extent_lowest  = BBlockTransferSrcVectorDim == 2 ? KzRaw : NzRaw;
+            const auto b1_extent_lowest = B1BlockTransferSrcVectorDim == 2 ? NzRaw : Gemm1NzRaw;
+            const auto c_extent_lowest  = Gemm1NzRaw;
+
+            if(!(a_extent_lowest % ABlockTransferSrcScalarPerVector == 0 &&
+                 b_extent_lowest % BBlockTransferSrcScalarPerVector == 0 &&
+                 b1_extent_lowest % B1BlockTransferSrcScalarPerVector == 0 &&
+                 c_extent_lowest % CShuffleBlockTransferScalarPerVector_NPerBlock == 0))
+            {
+                return false;
+            }
+
+            // Check vector load/store requirement
+            const auto a_stride_lowest = ABlockTransferSrcVectorDim == 2
+                                             ? device_arg.a_mz_kz_strides_[1]
+                                             : device_arg.a_mz_kz_strides_[0];
+            const auto b_stride_lowest = BBlockTransferSrcVectorDim == 2
+                                             ? device_arg.b_nz_kz_strides_[1]
+                                             : device_arg.b_nz_kz_strides_[0];
+            const auto b1_stride_lowest = B1BlockTransferSrcVectorDim == 2
+                                              ? device_arg.b1_nz_kz_strides_[1]
+                                              : device_arg.b1_nz_kz_strides_[0];
+            const auto c_stride_lowest =
+                device_arg.c_mz_gemm1nz_strides_[1]; // cshuffle assumes lowest dim in Gemm1Ns to be
+                                                     // contiguous
+
+            if(!(a_stride_lowest == 1 || b_stride_lowest == 1 || b1_stride_lowest == 1 ||
+                 c_stride_lowest == 1))
+            {
+                return false;
+            }
+
+            if(!GridwiseGemm::CheckValidity(kernel_arg.a_grid_desc_ak0_m_ak1_,
+                                            kernel_arg.b_grid_desc_bk0_n_bk1_,
+                                            kernel_arg.b1_grid_desc_bk0_n_bk1_,
+                                            kernel_arg.y_grid_desc_m_o_,
+                                            kernel_arg.block_2_ctile_map_))
+            {
+                return false;
+            }
+        }
+        return true;
+    }
+
+    // polymorphic
+    bool IsSupportedArgument(const BaseArgument* p_arg) override
+    {
+        return IsSupportedArgument(*dynamic_cast<const Argument*>(p_arg));
+    }
+
+    size_t GetWorkSpaceSize(const BaseArgument* p_arg) const override
+    {
+        return dynamic_cast<const Argument*>(p_arg)->group_count_ * sizeof(GroupKernelArg);
+    }
+
+    static auto MakeArgument(const std::vector<const void*>& p_As,
+                             const std::vector<const void*>& p_Bs,
+                             const std::vector<void*>& p_Zs,
+                             const std::vector<const void*>& p_B1s,
+                             const std::vector<const void*>& p_Cs, // for dS
+                             const std::vector<const void*>& p_LSEs,
+                             const std::vector<const void*>& p_Ygrads,
+                             std::vector<void*>& p_Qgrads,
+                             std::vector<void*>& p_Kgrads,
+                             std::vector<void*>& p_Vgrads,
+                             const std::array<void*, NumAcc0Bias>& p_acc0_biases,
+                             const std::array<void*, NumAcc1Bias>& p_acc1_biases,
+                             const std::vector<ProblemDesc>& problem_desc_vec,
+                             AElementwiseOperation a_element_op,
+                             BElementwiseOperation b_element_op,
+                             AccElementwiseOperation acc_element_op,
+                             B1ElementwiseOperation b1_element_op,
+                             CElementwiseOperation c_element_op,
+                             float p_drop,
+                             std::tuple<unsigned long long, unsigned long long> seeds)
+    {
+        return Argument{p_As,
+                        p_Bs,
+                        p_Zs,
+                        p_B1s,
+                        p_Cs,
+                        p_LSEs,
+                        p_Ygrads,
+                        p_Qgrads,
+                        p_Kgrads,
+                        p_Vgrads,
+                        p_acc0_biases,
+                        p_acc1_biases,
+                        problem_desc_vec,
+                        a_element_op,
+                        b_element_op,
+                        acc_element_op,
+                        b1_element_op,
+                        c_element_op,
+                        p_drop,
+                        seeds};
+    }
+
+    static auto MakeInvoker() { return Invoker{}; }
+
+    // polymorphic
+    // FIXME: constness
+    std::unique_ptr<BaseArgument>
+    MakeArgumentPointer(const std::vector<const void*>& p_As,
+                        const std::vector<const void*>& p_Bs,
+                        const std::vector<void*>& p_Zs,
+                        const std::vector<const void*>& p_B1s,
+                        const std::vector<const void*>& p_Cs, // for dS
+                        const std::vector<const void*>& p_LSEs,
+                        const std::vector<const void*>& p_Ygrads,
+                        std::vector<void*>& p_Qgrads,
+                        std::vector<void*>& p_Kgrads,
+                        std::vector<void*>& p_Vgrads,
+                        const std::array<void*, NumAcc0Bias>& p_acc0_biases,
+                        const std::array<void*, NumAcc1Bias>& p_acc1_biases,
+                        const std::vector<ProblemDesc>& problem_desc_vec,
+                        AElementwiseOperation a_element_op,
+                        BElementwiseOperation b_element_op,
+                        AccElementwiseOperation acc_element_op,
+                        B1ElementwiseOperation b1_element_op,
+                        CElementwiseOperation c_element_op,
+                        float p_drop,
+                        std::tuple<unsigned long long, unsigned long long> seeds) // override
+    {
+        return std::make_unique<Argument>(p_As,
+                                          p_Bs,
+                                          p_Zs,
+                                          p_B1s,
+                                          p_Cs,
+                                          p_LSEs,
+                                          p_Ygrads,
+                                          p_Qgrads,
+                                          p_Kgrads,
+                                          p_Vgrads,
+                                          p_acc0_biases, // cast in struct Argument
+                                          p_acc1_biases, // cast in struct Argument
+                                          problem_desc_vec,
+                                          a_element_op,
+                                          b_element_op,
+                                          acc_element_op,
+                                          b1_element_op,
+                                          c_element_op,
+                                          p_drop,
+                                          seeds);
+    }
+
+    // polymorphic
+    std::unique_ptr<BaseInvoker> MakeInvokerPointer() // override
+    {
+        return std::make_unique<Invoker>(Invoker{});
+    }
+
+    // polymorphic
+    std::string GetTypeString() const override
+    {
+        auto str = std::stringstream();
+
+        // clang-format off
+        str << "DeviceGroupedMultiheadAttentionBackward_Xdl_CShuffle_V2"
+            << "<"
+            << BlockSize << ", "
+            << MPerBlock << ", "
+            << NPerBlock << ", "
+            << KPerBlock << ", "
+            << AK1 << ", "
+            << BK1 << ", "
+            << MPerBlock << ", "
+            << Gemm1NPerBlock << ", "
+            << Gemm1KPerBlock << ", "
+            << B1K1 << ", "
+            << getGemmSpecializationString(GemmSpec) << ", "
+            << "ASpec" << getTensorSpecializationString(ASpec) << ", "
+            << "B0Spec" << getTensorSpecializationString(BSpec) << ", "
+            << "B1Spec" << getTensorSpecializationString(B1Spec) << ", "
+            << "CSpec" << getTensorSpecializationString(CSpec) << ", "
+            << getMaskingSpecializationString(MaskingSpec) << ">";
+        // clang-format on
+
+        return str.str();
+    }
+};
+
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck