resolved conflicts

include/ck/utility/type_convert.hpp

@@ -14,6 +14,41 @@ namespace ck {
 #define __gfx94__
 #endif
 
+// Declare a template function for bf16 conversion using RTN
+template <typename Y, typename X>
+__host__ __device__ constexpr Y bf16_convert_rtn(X x);
+
+// Convert fp32 to bf16 with RTN if higher precision is needed
+template <>
+inline __host__ __device__ constexpr bhalf_t bf16_convert_rtn<bhalf_t, float>(float x)
+{
+    // Nan check
+    if(x != x)
+    {
+        return uint16_t(0x7FC0);
+    }
+
+    union
+    {
+        float fp32;
+        uint32_t int32;
+    } u = {x};
+
+    const uint32_t first_bf16_mantisa_bit = ((u.int32 >> 16) & 1);
+    constexpr uint32_t rounding_bias      = uint32_t((1 << 15) - 1);
+
+    return uint16_t((u.int32 + first_bf16_mantisa_bit + rounding_bias) >> 16);
+}
+
+// convert fp16 to bfp16 via fp32 with RTN if higher precision is needed
+template <>
+inline __host__ __device__ constexpr bhalf_t bf16_convert_rtn<bhalf_t, half_t>(half_t x)
+{
+    float x_fp32 = static_cast<float>(x);
+
+    return bf16_convert_rtn<bhalf_t>(x_fp32);
+}
+
 // Convert X to Y, both X and Y are non-const data types.
 template <typename Y,
           typename X,
@@ -51,17 +86,15 @@ inline __host__ __device__ constexpr float type_convert<float, bhalf_t>(bhalf_t
     return u.fp32;
 }
 
-// convert fp32 to bfp16
+// convert fp32 to bfp16, round to nearest even
 template <>
 inline __host__ __device__ constexpr bhalf_t type_convert<bhalf_t, float>(float x)
 {
-    union
-    {
-        float fp32;
-        uint32_t int32;
-    } u = {x};
-
+#if CK_USE_RNE_BF16_CONVERSION
+    return bf16_convert_rtn<bhalf_t>(x);
+#else
     return uint16_t(u.int32 >> 16);
+#endif
 }
 
 // convert bfp16 to fp16 via fp32
@@ -615,60 +648,4 @@ inline __host__ __device__ void array_convert(Array<Y, NumElems>& y, const Array
     }
 }
 
-// Declare a template function for bf16 conversion using RTN
-template <typename Y, typename X>
-__host__ __device__ constexpr Y bf16_convert_rtn(X x);
-
-// Convert fp32 to bf16 with RTN if higher precision is needed
-template <>
-inline __host__ __device__ constexpr bhalf_t bf16_convert_rtn<bhalf_t, float>(float x)
-{
-    union
-    {
-        float fp32;
-        uint32_t int32;
-    } u = {x};
-
-    // When the exponent bits are not all 1s, then the value is zero, normal,
-    // or subnormal. We round the bfloat16 mantissa up by adding 0x7FFF, plus
-    // 1 if the least significant bit of the bfloat16 mantissa is 1 (odd).
-    // This causes the bfloat16's mantissa to be incremented by 1 if the 16
-    // least significant bits of the float mantissa are greater than 0x8000,
-    // or if they are equal to 0x8000 and the least significant bit of the
-    // bfloat16 mantissa is 1 (odd). This causes it to be rounded to even when
-    // the lower 16 bits are exactly 0x8000. If the bfloat16 mantissa already
-    // has the value 0x7f, then incrementing it causes it to become 0x00 and
-    // the exponent is incremented by one, which is the next higher FP value
-    // to the unrounded bfloat16 value. When the bfloat16 value is subnormal
-    // with an exponent of 0x00 and a mantissa of 0x7f, it may be rounded up
-    // to a normal value with an exponent of 0x01 and a mantissa of 0x00.
-    // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F,
-    // incrementing it causes it to become an exponent of 0xFF and a mantissa
-    // of 0x00, which is Inf, the next higher value to the unrounded value.
-    bool flag0 = ~u.int32 & 0x7f800000;
-
-    // When all of the exponent bits are 1, the value is Inf or NaN.
-    // Inf is indicated by a zero mantissa. NaN is indicated by any nonzero
-    // mantissa bit. Quiet NaN is indicated by the most significant mantissa
-    // bit being 1. Signaling NaN is indicated by the most significant
-    // mantissa bit being 0 but some other bit(s) being 1. If any of the
-    // lower 16 bits of the mantissa are 1, we set the least significant bit
-    // of the bfloat16 mantissa, in order to preserve signaling NaN in case
-    // the bfloat16's mantissa bits are all 0.
-    bool flag1 = !flag0 && (u.int32 & 0xffff);
-
-    u.int32 += flag0 ? 0x7fff + ((u.int32 >> 16) & 1) : 0; // Round to nearest, round to even
-    u.int32 |= flag1 ? 0x10000 : 0x0;                      // Preserve signaling NaN
-
-    return uint16_t(u.int32 >> 16);
-}
-
-// convert fp16 to bfp16 via fp32 with RTN if higher precision is needed
-template <>
-inline __host__ __device__ constexpr bhalf_t bf16_convert_rtn<bhalf_t, half_t>(half_t x)
-{
-    float x_fp32 = static_cast<float>(x);
-
-    return bf16_convert_rtn<bhalf_t>(x_fp32);
-}
 } // namespace ck

include/ck_tile/core.hpp

@@ -54,7 +54,6 @@
 #include "ck_tile/core/tensor/tile_window_linear.hpp"
 #include "ck_tile/core/tensor/tile_window_utils.hpp"
 #include "ck_tile/core/tensor/update_tile.hpp"
-#include "ck_tile/core/utility/amd_address_space.hpp"
 #include "ck_tile/core/utility/bit_cast.hpp"
 #include "ck_tile/core/utility/functional.hpp"
 #include "ck_tile/core/utility/functional_with_tuple.hpp"

include/ck_tile/core/arch/arch.hpp

 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -12,18 +12,37 @@
 
 namespace ck_tile {
 
-enum struct address_space_enum
+template <typename, bool>
+struct safe_underlying_type;
+
+template <typename T>
+struct safe_underlying_type<T, true>
+{
+    using type = std::underlying_type_t<T>;
+};
+
+template <typename T>
+struct safe_underlying_type<T, false>
+{
+    using type = void;
+};
+
+template <typename T>
+using safe_underlying_type_t = typename safe_underlying_type<T, std::is_enum<T>::value>::type;
+
+enum struct address_space_enum : std::uint16_t
 {
-    generic,
+    generic = 0,
     global,
     lds,
     sgpr,
-    vgpr,
+    constant,
+    vgpr
 };
 
-enum struct memory_operation_enum
+enum struct memory_operation_enum : std::uint16_t
 {
-    set,
+    set = 0,
     atomic_add,
     atomic_max,
     add
@@ -109,4 +128,30 @@ CK_TILE_DEVICE void s_nop(index_t cnt = 0)
 #endif
 }
 
+#define CK_CONSTANT_ADDRESS_SPACE \
+    __attribute__((address_space( \
+        static_cast<safe_underlying_type_t<address_space_enum>>(address_space_enum::constant))))
+
+template <typename T>
+__device__ T* cast_pointer_to_generic_address_space(T CK_CONSTANT_ADDRESS_SPACE* p)
+{
+    // cast a pointer in "Constant" address space (4) to "Generic" address space (0)
+    // only c-style pointer cast seems be able to be compiled
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wold-style-cast"
+    return (T*)(p); // NOLINT(old-style-cast)
+#pragma clang diagnostic pop
+}
+
+template <typename T>
+__host__ __device__ T CK_CONSTANT_ADDRESS_SPACE* cast_pointer_to_constant_address_space(T* p)
+{
+    // cast a pointer in "Generic" address space (0) to "Constant" address space (4)
+    // only c-style pointer cast seems be able to be compiled;
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wold-style-cast"
+    return (T CK_CONSTANT_ADDRESS_SPACE*)p; // NOLINT(old-style-cast)
+#pragma clang diagnostic pop
+}
+
 } // namespace ck_tile

include/ck_tile/core/numeric/bfloat16.hpp

 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/utility/bit_cast.hpp"
@@ -376,6 +376,16 @@ struct numeric<bfloat16_t>
     }
 };
 
+template <typename T>
+struct numeric_traits;
+
+template <>
+struct numeric_traits<bfloat16_t>
+{
+    static constexpr int exp  = 8;
+    static constexpr int mant = 7;
+};
+
 #if CK_TILE_USE_CUSTOM_DATA_TYPE
 CK_TILE_ARITHMETIC_USING_FLOAT(CK_TILE_HOST_DEVICE, bfloat16_t)
 #endif

include/ck_tile/core/utility/amd_address_space.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
-
-#pragma once
-
-#include "ck_tile/core/config.hpp"
-
-// Address Space for AMDGCN
-// https://llvm.org/docs/AMDGPUUsage.html#address-space
-
-namespace ck_tile {
-
-#define CK_CONSTANT_ADDRESS_SPACE __attribute__((address_space(4)))
-
-template <typename T>
-__device__ T* cast_pointer_to_generic_address_space(T CK_CONSTANT_ADDRESS_SPACE* p)
-{
-    // cast a pointer in "Constant" address space (4) to "Generic" address space (0)
-    // only c-style pointer cast seems be able to be compiled
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wold-style-cast"
-    return (T*)p; // NOLINT(old-style-cast)
-#pragma clang diagnostic pop
-}
-
-template <typename T>
-__host__ __device__ T CK_CONSTANT_ADDRESS_SPACE* cast_pointer_to_constant_address_space(T* p)
-{
-    // cast a pointer in "Generic" address space (0) to "Constant" address space (4)
-    // only c-style pointer cast seems be able to be compiled
-#pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wold-style-cast"
-    return (T CK_CONSTANT_ADDRESS_SPACE*)p; // NOLINT(old-style-cast)
-#pragma clang diagnostic pop
-}
-
-} // namespace ck_tile

include/ck_tile/core/utility/type_traits.hpp

@@ -109,4 +109,22 @@ CK_TILE_HOST_DEVICE PY c_style_pointer_cast(PX p_x)
 #pragma clang diagnostic pop
 }
 
+template <typename CompareTo, typename... Rest>
+struct is_any_of : std::false_type
+{
+};
+
+template <typename CompareTo, typename FirstType>
+struct is_any_of<CompareTo, FirstType> : std::is_same<CompareTo, FirstType>
+{
+};
+
+template <typename CompareTo, typename FirstType, typename... Rest>
+struct is_any_of<CompareTo, FirstType, Rest...>
+    : std::integral_constant<bool,
+                             std::is_same<CompareTo, FirstType>::value ||
+                                 is_any_of<CompareTo, Rest...>::value>
+{
+};
+
 } // namespace ck_tile

include/ck_tile/core/utility/unary_element_function.hpp

@@ -51,16 +51,18 @@ struct composes<F>
 template <typename... Ts>
 __host__ __device__ composes(Ts&&...)->composes<remove_cvref_t<Ts>...>;
 
-template <typename To>
+template <typename SaturateType>
 struct saturates
 {
-    template <typename From>
-    CK_TILE_HOST_DEVICE constexpr auto operator()(const From& from) const
-        -> std::enable_if_t<std::is_arithmetic_v<From>, From>
+    // NOTE: this function does not return SaturateType value
+    // it is user's responsiblity to do further cast or not
+    template <typename AccType>
+    CK_TILE_HOST_DEVICE constexpr auto operator()(const AccType& a_) const
+        -> std::enable_if_t<std::is_arithmetic_v<AccType>, AccType>
     {
-        return clamp(from,
-                     type_convert<From>(numeric<To>::lowest()),
-                     type_convert<From>(numeric<To>::max()));
+        return clamp(a_,
+                     type_convert<AccType>(numeric<SaturateType>::lowest()),
+                     type_convert<AccType>(numeric<SaturateType>::max()));
     }
 };
 

include/ck_tile/host/check_err.hpp

 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -18,6 +18,112 @@
 
 namespace ck_tile {
 
+template <typename ComputeDataType, typename OutDataType, typename AccDataType = ComputeDataType>
+double get_relative_threshold(const int number_of_accumulations = 1)
+{
+    using F8   = ck_tile::fp8_t;
+    using F16  = ck_tile::half_t;
+    using BF16 = ck_tile::bf16_t;
+    using F32  = float;
+    using I8   = int8_t;
+    using I32  = int32_t;
+
+    static_assert(is_any_of<ComputeDataType, F8, F16, BF16, F32, I8, I32, int>::value,
+                  "Warning: Unhandled ComputeDataType for setting up the relative threshold!");
+
+    double compute_error = 0;
+    if constexpr(is_any_of<ComputeDataType, I8, I32, int>::value)
+    {
+        return 0;
+    }
+    else
+    {
+        compute_error = std::pow(2, -numeric_traits<ComputeDataType>::mant) * 0.5;
+    }
+
+    static_assert(is_any_of<OutDataType, F8, F16, BF16, F32, I8, I32, int>::value,
+                  "Warning: Unhandled OutDataType for setting up the relative threshold!");
+
+    double output_error = 0;
+    if constexpr(is_any_of<OutDataType, I8, I32, int>::value)
+    {
+        return 0;
+    }
+    else
+    {
+        output_error = std::pow(2, -numeric_traits<OutDataType>::mant) * 0.5;
+    }
+    double midway_error = std::max(compute_error, output_error);
+
+    static_assert(is_any_of<AccDataType, F8, F16, BF16, F32, I8, I32, int>::value,
+                  "Warning: Unhandled AccDataType for setting up the relative threshold!");
+
+    double acc_error = 0;
+    if constexpr(is_any_of<AccDataType, I8, I32, int>::value)
+    {
+        return 0;
+    }
+    else
+    {
+        acc_error = std::pow(2, -numeric_traits<AccDataType>::mant) * 0.5 * number_of_accumulations;
+    }
+    return std::max(acc_error, midway_error);
+}
+
+template <typename ComputeDataType, typename OutDataType, typename AccDataType = ComputeDataType>
+double get_absolute_threshold(const double max_possible_num, const int number_of_accumulations = 1)
+{
+    using F8   = ck_tile::fp8_t;
+    using F16  = ck_tile::half_t;
+    using BF16 = ck_tile::bf16_t;
+    using F32  = float;
+    using I8   = int8_t;
+    using I32  = int32_t;
+
+    static_assert(is_any_of<ComputeDataType, F8, F16, BF16, F32, I8, I32, int>::value,
+                  "Warning: Unhandled ComputeDataType for setting up the absolute threshold!");
+
+    auto expo            = std::log2(std::abs(max_possible_num));
+    double compute_error = 0;
+    if constexpr(is_any_of<ComputeDataType, I8, I32, int>::value)
+    {
+        return 0;
+    }
+    else
+    {
+        compute_error = std::pow(2, expo - numeric_traits<ComputeDataType>::mant) * 0.5;
+    }
+
+    static_assert(is_any_of<OutDataType, F8, F16, BF16, F32, I8, I32, int>::value,
+                  "Warning: Unhandled OutDataType for setting up the absolute threshold!");
+
+    double output_error = 0;
+    if constexpr(is_any_of<OutDataType, I8, I32, int>::value)
+    {
+        return 0;
+    }
+    else
+    {
+        output_error = std::pow(2, expo - numeric_traits<OutDataType>::mant) * 0.5;
+    }
+    double midway_error = std::max(compute_error, output_error);
+
+    static_assert(is_any_of<AccDataType, F8, F16, BF16, F32, I8, I32, int>::value,
+                  "Warning: Unhandled AccDataType for setting up the absolute threshold!");
+
+    double acc_error = 0;
+    if constexpr(is_any_of<AccDataType, I8, I32, int>::value)
+    {
+        return 0;
+    }
+    else
+    {
+        acc_error =
+            std::pow(2, expo - numeric_traits<AccDataType>::mant) * 0.5 * number_of_accumulations;
+    }
+    return std::max(acc_error, midway_error);
+}
+
 template <typename T>
 std::ostream& operator<<(std::ostream& os, const std::vector<T>& v)
 {
@@ -337,7 +443,11 @@ std::enable_if_t<(std::is_same_v<ranges::range_value_t<Range>, ranges::range_val
     }
     if(!res)
     {
-        std::cerr << std::setw(12) << std::setprecision(7) << "max err: " << max_err << std::endl;
+        const float error_percent =
+            static_cast<float>(err_count) / static_cast<float>(out.size()) * 100.f;
+        std::cerr << "max err: " << max_err;
+        std::cerr << ", number of errors: " << err_count;
+        std::cerr << ", " << error_percent << "% wrong values" << std::endl;
     }
     return res;
 }

include/ck_tile/host/convolution_host_tensor_descriptor_helper.hpp

@@ -14,57 +14,41 @@ namespace detail {
 template <typename OldLayout>
 CK_TILE_HOST std::vector<std::size_t> get_layout_transpose_gnchw_to_old()
 {
-    if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNCW> ||
-                 std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GKCX> ||
-                 std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNKW>)
+    using namespace ck_tile::tensor_layout::convolution;
+
+    if constexpr(is_any_of<OldLayout, GNCW, GKCX, GNKW>::value)
     {
         return {0, 1, 2, 3};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNCHW> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GKCYX> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNKHW>)
+    else if constexpr(is_any_of<OldLayout, GNCHW, GKCYX, GNKHW>::value)
     {
         return {0, 1, 2, 3, 4};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNCDHW> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GKCZYX> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNKDHW>)
+    else if constexpr(is_any_of<OldLayout, GNCDHW, GKCZYX, GNKDHW>::value)
     {
         return {0, 1, 2, 3, 4, 5};
     }
-    if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNWC> ||
-                 std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GKXC> ||
-                 std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNWK>)
+    if constexpr(is_any_of<OldLayout, GNWC, GKXC, GNWK>::value)
     {
         return {0, 1, 3, 2};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNHWC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GKYXC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNHWK>)
+    else if constexpr(is_any_of<OldLayout, GNHWC, GKYXC, GNHWK>::value)
     {
         return {0, 1, 4, 2, 3};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNDHWC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GKZYXC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::GNDHWK>)
+    else if constexpr(is_any_of<OldLayout, GNDHWC, GKZYXC, GNDHWK>::value)
     {
         return {0, 1, 5, 2, 3, 4};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::NWGC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::KXGC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::NWGK>)
+    else if constexpr(is_any_of<OldLayout, NWGC, KXGC, NWGK>::value)
     {
         return {2, 0, 3, 1};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::NHWGC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::KYXGC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::NHWGK>)
+    else if constexpr(is_any_of<OldLayout, NHWGC, KYXGC, NHWGK>::value)
     {
         return {3, 0, 4, 1, 2};
     }
-    else if constexpr(std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::NDHWGC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::KZYXGC> ||
-                      std::is_same_v<OldLayout, ck_tile::tensor_layout::convolution::NDHWGK>)
+    else if constexpr(is_any_of<OldLayout, NDHWGC, KZYXGC, NDHWGK>::value)
     {
         return {4, 0, 5, 1, 2, 3};
     }
@@ -83,11 +67,11 @@ template <typename InLayout>
 CK_TILE_HOST HostTensorDescriptor
 make_input_host_tensor_descriptor_g_n_c_wis_packed(const ck_tile::conv::ConvParam& param)
 {
+    using namespace ck_tile::tensor_layout::convolution;
+
     std::vector<std::size_t> physical_lengths;
 
-    if constexpr(std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::GNCW> ||
-                 std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::GNCHW> ||
-                 std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::GNCDHW>)
+    if constexpr(is_any_of<InLayout, GNCW, GNCHW, GNCDHW>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.G_),
                                                     static_cast<std::size_t>(param.N_),
@@ -97,9 +81,7 @@ make_input_host_tensor_descriptor_g_n_c_wis_packed(const ck_tile::conv::ConvPara
                                 param.input_spatial_lengths_.begin(),
                                 param.input_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
-    else if constexpr(std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::GNWC> ||
-                      std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::GNHWC> ||
-                      std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::GNDHWC>)
+    else if constexpr(is_any_of<InLayout, GNWC, GNHWC, GNDHWC>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.G_),
                                                     static_cast<std::size_t>(param.N_),
@@ -109,9 +91,7 @@ make_input_host_tensor_descriptor_g_n_c_wis_packed(const ck_tile::conv::ConvPara
                                 param.input_spatial_lengths_.begin(),
                                 param.input_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
-    else if constexpr(std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::NWGC> ||
-                      std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::NHWGC> ||
-                      std::is_same_v<InLayout, ck_tile::tensor_layout::convolution::NDHWGC>)
+    else if constexpr(is_any_of<InLayout, NWGC, NHWGC, NDHWGC>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.N_),
                                                     static_cast<std::size_t>(param.G_),
@@ -139,11 +119,11 @@ template <typename WeiLayout>
 CK_TILE_HOST HostTensorDescriptor
 make_weight_host_tensor_descriptor_g_k_c_xs_packed(const ck_tile::conv::ConvParam& param)
 {
+    using namespace ck_tile::tensor_layout::convolution;
+
     std::vector<std::size_t> physical_lengths;
 
-    if constexpr(std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::KXC> ||
-                 std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::KYXC> ||
-                 std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::KZYXC>)
+    if constexpr(is_any_of<WeiLayout, KXC, KYXC, KZYXC>::value)
     {
         if(param.G_ != 1)
         {
@@ -157,9 +137,7 @@ make_weight_host_tensor_descriptor_g_k_c_xs_packed(const ck_tile::conv::ConvPara
                                 param.filter_spatial_lengths_.begin(),
                                 param.filter_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
-    else if constexpr(std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::GKCX> ||
-                      std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::GKCYX> ||
-                      std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::GKCZYX>)
+    else if constexpr(is_any_of<WeiLayout, GKCX, GKCYX, GKCZYX>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.G_),
                                                     static_cast<std::size_t>(param.K_),
@@ -169,9 +147,7 @@ make_weight_host_tensor_descriptor_g_k_c_xs_packed(const ck_tile::conv::ConvPara
                                 param.filter_spatial_lengths_.begin(),
                                 param.filter_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
-    else if constexpr(std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::GKXC> ||
-                      std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::GKYXC> ||
-                      std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::GKZYXC>)
+    else if constexpr(is_any_of<WeiLayout, GKXC, GKYXC, GKZYXC>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.G_),
                                                     static_cast<std::size_t>(param.K_),
@@ -181,9 +157,7 @@ make_weight_host_tensor_descriptor_g_k_c_xs_packed(const ck_tile::conv::ConvPara
                                 param.filter_spatial_lengths_.begin(),
                                 param.filter_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
-    else if constexpr(std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::KXGC> ||
-                      std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::KYXGC> ||
-                      std::is_same_v<WeiLayout, ck_tile::tensor_layout::convolution::KZYXGC>)
+    else if constexpr(is_any_of<WeiLayout, KXGC, KYXGC, KZYXGC>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.K_),
                                                     static_cast<std::size_t>(param.G_),
@@ -211,11 +185,11 @@ template <typename OutLayout>
 CK_TILE_HOST HostTensorDescriptor
 make_output_host_tensor_descriptor_g_n_k_wos_packed(const ck_tile::conv::ConvParam& param)
 {
+    using namespace ck_tile::tensor_layout::convolution;
+
     std::vector<std::size_t> physical_lengths;
 
-    if constexpr(std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::GNKW> ||
-                 std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::GNKHW> ||
-                 std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::GNKDHW>)
+    if constexpr(is_any_of<OutLayout, GNKW, GNKHW, GNKDHW>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.G_),
                                                     static_cast<std::size_t>(param.N_),
@@ -226,9 +200,7 @@ make_output_host_tensor_descriptor_g_n_k_wos_packed(const ck_tile::conv::ConvPar
                                 param.output_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
     // separate from legacy code above
-    else if constexpr(std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::GNWK> ||
-                      std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::GNHWK> ||
-                      std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::GNDHWK>)
+    else if constexpr(is_any_of<OutLayout, GNWK, GNHWK, GNDHWK>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.G_),
                                                     static_cast<std::size_t>(param.N_),
@@ -238,9 +210,7 @@ make_output_host_tensor_descriptor_g_n_k_wos_packed(const ck_tile::conv::ConvPar
                                 param.output_spatial_lengths_.begin(),
                                 param.output_spatial_lengths_.begin() + param.num_dim_spatial_);
     }
-    else if constexpr(std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::NWGK> ||
-                      std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::NHWGK> ||
-                      std::is_same_v<OutLayout, ck_tile::tensor_layout::convolution::NDHWGK>)
+    else if constexpr(is_any_of<OutLayout, NWGK, NHWGK, NDHWGK>::value)
     {
         physical_lengths = std::vector<std::size_t>{static_cast<std::size_t>(param.N_),
                                                     static_cast<std::size_t>(param.G_),

include/ck_tile/host/host_tensor.hpp

 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -678,4 +678,43 @@ struct HostTensor
     Descriptor mDesc;
     Data mData;
 };
+
+template <bool is_row_major>
+auto host_tensor_descriptor(std::size_t row,
+                            std::size_t col,
+                            std::size_t stride,
+                            bool_constant<is_row_major>)
+{
+    using namespace ck_tile::literals;
+
+    if constexpr(is_row_major)
+    {
+        return HostTensorDescriptor({row, col}, {stride, 1_uz});
+    }
+    else
+    {
+        return HostTensorDescriptor({row, col}, {1_uz, stride});
+    }
+}
+template <bool is_row_major>
+auto get_default_stride(std::size_t row,
+                        std::size_t col,
+                        std::size_t stride,
+                        bool_constant<is_row_major>)
+{
+    if(stride == 0)
+    {
+        if constexpr(is_row_major)
+        {
+            return col;
+        }
+        else
+        {
+            return row;
+        }
+    }
+    else
+        return stride;
+}
+
 } // namespace ck_tile

include/ck_tile/host/reference/reference_fused_moe.hpp

@@ -73,7 +73,7 @@ void reference_fused_moe(
     ck_tile::index_t tokens,
     ck_tile::index_t experts,
     ck_tile::index_t hidden_size,
-    ck_tile::index_t intermediate_size, // this size is for gate/up
+    ck_tile::index_t intermediate_size, // this size is for gate/up/down
     ck_tile::index_t topk,
     ck_tile::index_t gate_only)
 {
@@ -82,19 +82,8 @@ void reference_fused_moe(
     assert(sorted_expert_ids_host.get_num_of_dimension() == 1);
     assert(num_sorted_tiles_host.get_element_size() == 1);
     ck_tile::index_t num_sorted_tiles    = num_sorted_tiles_host.mData[0] / block_m;
-    ck_tile::index_t intermediate_size_0 = intermediate_size;
-    ck_tile::index_t intermediate_size_1 = intermediate_size / (gate_only ? 1 : 2);
-
-    // TODO: better remove this in the future, or modify the token_id value
-    auto get_topk_id = [&](ck_tile::index_t token_id_, ck_tile::index_t expert_id_) {
-        for(ck_tile::index_t i_ = 0; i_ < topk; i_++)
-        {
-            if(token_ids_host(token_id_, i_) == expert_id_)
-                return i_;
-        }
-        throw std::runtime_error("not correct token/expert pair\n");
-        return -1; // TODO: not correct!!
-    };
+    ck_tile::index_t intermediate_size_0 = intermediate_size * (gate_only ? 1 : 2);
+    ck_tile::index_t intermediate_size_1 = intermediate_size;
 
     ck_tile::HostTensor<AccDataType> out_topk_tokens({tokens, topk, hidden_size});
 
@@ -105,11 +94,31 @@ void reference_fused_moe(
         if(i_tile >= num_sorted_tiles)
             return;
         ck_tile::index_t i_expert = sorted_expert_ids_host.mData[i_tile];
-        ck_tile::index_t i_token  = sorted_token_ids_host.mData[i_flatten];
+
+#if CK_TILE_REFERENCE_MOE_SORTING_MOCK_ID
+        ck_tile::index_t i_token = sorted_token_ids_host.mData[i_flatten];
+        ck_tile::index_t i_topk  = i_token >> 24;
+        i_token &= 0xffffff;
+        if(i_token >= tokens)
+            return;
+        (void)token_ids_host;
+#else
+        // TODO: better remove this in the future, or modify the token_id value
+        auto get_topk_id = [&](ck_tile::index_t token_id_, ck_tile::index_t expert_id_) {
+            for(ck_tile::index_t i_ = 0; i_ < topk; i_++)
+            {
+                if(token_ids_host(token_id_, i_) == expert_id_)
+                    return i_;
+            }
+            throw std::runtime_error("not correct token/expert pair\n");
+            return -1; // TODO: not correct!!
+        };
+        ck_tile::index_t i_token = sorted_token_ids_host.mData[i_flatten];
         if(i_token >= tokens)
             return;
         ck_tile::index_t i_topk = get_topk_id(i_token, i_expert); // TODO: ugly
-        auto weight             = sorted_weight_host.mData[i_flatten];
+#endif
+        auto weight = sorted_weight_host.mData[i_flatten];
 
         ck_tile::HostTensor<AccDataType> acc_0({1, intermediate_size_0});
         // first gemm

include/ck_tile/host/reference/reference_rmsnorm2d_fwd.hpp

@@ -8,16 +8,40 @@
 
 namespace ck_tile {
 
+// Note: for simplicity, each functor only care about single M
+struct reference_rmsnorm2d_default_epilogue
+{
+    template <typename OutDataType, typename AccDataType>
+    void operator()(int m, HostTensor<OutDataType>& o, const HostTensor<AccDataType>& acc)
+    {
+        const int N = acc.mDesc.get_lengths()[1];
+        for(int n = 0; n < N; ++n)
+        {
+            o(m, n) = ck_tile::type_convert<OutDataType>(acc(m, n));
+        }
+    }
+
+    template <typename OutDataType, typename AccDataType>
+    auto operator()(int m, const HostTensor<AccDataType>& acc)
+    {
+        HostTensor<OutDataType> o(acc.get_lengths(), acc.get_strides());
+        operator()(m, o, acc);
+        return o;
+    }
+};
+
 template <typename XDataType,
           typename GammaDataType,
           typename ComputeDataType,
           typename YDataType,
-          typename InvRmsDataType>
+          typename InvRmsDataType,
+          typename Epilogue = reference_rmsnorm2d_default_epilogue>
 void reference_rmsnorm2d_fwd(const HostTensor<XDataType>& x_m_n,
                              const HostTensor<GammaDataType>& gamma_n,
                              HostTensor<YDataType>& y_m_n,
                              HostTensor<InvRmsDataType>& invRms_m,
-                             ComputeDataType epsilon)
+                             ComputeDataType epsilon,
+                             Epilogue epilogue_functor = {})
 {
     auto rmsnorm2d_fwd_func = [&](auto m) {
         const int N = x_m_n.mDesc.get_lengths()[1];
@@ -37,13 +61,15 @@ void reference_rmsnorm2d_fwd(const HostTensor<XDataType>& x_m_n,
         if constexpr(!std::is_same_v<InvRmsDataType, ck_tile::null_type>)
             invRms_m(m) = ck_tile::type_convert<InvRmsDataType>(divisor);
 
+        HostTensor<ComputeDataType> acc(x_m_n.get_lengths(), x_m_n.get_strides());
         for(int n = 0; n < N; ++n)
         {
             ComputeDataType x     = ck_tile::type_convert<ComputeDataType>(x_m_n(m, n));
             ComputeDataType gamma = ck_tile::type_convert<ComputeDataType>(gamma_n(n));
-            auto y                = x * divisor * gamma;
-            y_m_n(m, n)           = ck_tile::type_convert<YDataType>(y);
+            acc(m, n)             = x * divisor * gamma;
         }
+
+        epilogue_functor(m, y_m_n, acc);
     };
 
     make_ParallelTensorFunctor(rmsnorm2d_fwd_func, invRms_m.mDesc.get_lengths()[0])(

include/ck_tile/host/reference/reference_rowwise_quantization2d.hpp

@@ -22,7 +22,7 @@ CK_TILE_HOST void reference_rowwise_quantization2d(const HostTensor<XDataType>&
             // scale = amax / 127 for int8
             auto v_scale = type_convert<XDataType>(scale_m(m));
             auto v_qx    = v_x / v_scale;
-            qx_m_n(m, n) = saturates<QXDataType>{}(v_qx);
+            qx_m_n(m, n) = type_convert<QXDataType>(saturates<QXDataType>{}(v_qx));
         }
     };
 

include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp

@@ -719,7 +719,82 @@ struct Silu
         constexpr T one = type_convert<T>(1);
         y               = x * (one / (one + ck_tile::exp(-x)));
     };
+
+    template <>
+    CK_TILE_HOST_DEVICE void operator()<fp32x2_t>(fp32x2_t& y, const fp32x2_t& x) const
+    {
+        constexpr auto one = type_convert<float>(1);
+        y[0]               = x[0] * __builtin_amdgcn_rcpf(one + ck_tile::exp(-x[0]));
+        y[1]               = x[1] * __builtin_amdgcn_rcpf(one + ck_tile::exp(-x[1]));
+    };
+};
+
+#if 0
+// Silu, the formular is not so good to do inline asm (dependency)
+// we put the code here purposely if in the future ppl want to try
+struct SiluAsm
+{
+    template <typename T>
+    CK_TILE_HOST void operator()(T& y, T& x) const
+    {
+        static_assert(std::is_same_v<T, float>, "Data type is not supported by this operation!");
+        constexpr T one = type_convert<T>(1);
+        y               = x * (one / (one + ck_tile::exp(-x)));
+    };
+
+    template <typename T>
+    CK_TILE_DEVICE void operator()(T& y, T& x) const
+    {
+        static_assert(std::is_same_v<T, float>, "Data type is not supported by this operation!");
+
+        const uint32_t log2e_neg_ = 0x3fb8aa3b | 0x80000000; // log2e_v<float> * -1;
+
+        // NOTE: x/y can't be same register before inline asm
+        // "+v" as y, "v" as x is not enought, x/y stil maybe put to same register
+        T tmp = x;
+        asm volatile("v_mul_f32 %[v_y], %[s_log2e], %[v_x]\n"
+                     "v_exp_f32 %[v_y], %[v_y]\n"
+                     "s_nop 0           ; hazard for exp\n"
+                     "v_add_f32 %[v_y], %[v_y], 1.0\n"
+                     "v_rcp_f32 %[v_y], %[v_y]\n"
+                     "s_nop 0           ; hazard for rcp\n"
+                     "v_mul_f32 %[v_y], %[v_x], %[v_y]\n"
+                     : [v_y] "+v"(y), [v_x] "+v"(tmp)
+                     : [s_log2e] "s"(log2e_neg_)
+                     :);
+    };
+
+    template <>
+    CK_TILE_HOST void operator()<fp32x2_t>(fp32x2_t& y, fp32x2_t& x) const
+    {
+        constexpr auto one = type_convert<float>(1);
+        y[0]               = x[0] * (one / (one + ck_tile::exp(-x[0])));
+        y[1]               = x[1] * (one / (one + ck_tile::exp(-x[1])));
+    };
+
+    template <>
+    CK_TILE_DEVICE void operator()<fp32x2_t>(fp32x2_t& y, fp32x2_t& x) const
+    {
+        const uint32_t log2e_neg_ = 0x3fb8aa3b | 0x80000000; // log2e_v<float> * -1;
+
+        // NOTE: x/y can't be same register before inline asm
+        // float tmp0 = x[0], tmp1 = x[1];
+        asm volatile("v_mul_f32 %[v_y0], %[s_log2e], %[v_x0]\n"
+                     "v_mul_f32 %[v_y1], %[s_log2e], %[v_x1]\n"
+                     "v_exp_f32 %[v_y0], %[v_y0]\n"
+                     "v_exp_f32 %[v_y1], %[v_y1]\n"
+                     "v_add_f32 %[v_y0], %[v_y0], 1.0\n"
+                     "v_add_f32 %[v_y1], %[v_y1], 1.0\n"
+                     "v_rcp_f32 %[v_y0], %[v_y0]\n"
+                     "v_rcp_f32 %[v_y1], %[v_y1]\n"
+                     "v_mul_f32 %[v_y0], %[v_x0], %[v_y0]\n"
+                     "v_mul_f32 %[v_y1], %[v_x1], %[v_y1]\n"
+                     : [v_y0] "+v"(y[0]), [v_y1] "+v"(y[1]), [v_x0] "+v"(x[0]), [v_x1] "+v"(x[1])
+                     : [s_log2e] "s"(log2e_neg_)
+                     :);
+    };
 };
+#endif
 
 struct TanH
 {

include/ck_tile/ops/epilogue/dynamic_quant_epilogue.hpp

 // SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #pragma once
 
@@ -24,19 +24,19 @@ struct DynamicQuantEpilogueTraits
 
 // this epilogue just store out a M*N matrix, row major
 template <typename AccDataType_,
-          typename XScaleDataType_,
+          typename SmoothScaleDataType_,
           typename YScaleDataType_,
           typename ODataType_,
           typename BlockShape_,
           typename Traits_>
 struct DynamicQuantEpilogueProblem
 {
-    using AccDataType    = remove_cvref_t<AccDataType_>;
-    using XScaleDataType = remove_cvref_t<XScaleDataType_>;
-    using YScaleDataType = remove_cvref_t<YScaleDataType_>;
-    using ODataType      = remove_cvref_t<ODataType_>;
-    using BlockShape     = remove_cvref_t<BlockShape_>; // can consum generic 2d shape
-    using Traits         = remove_cvref_t<Traits_>;
+    using AccDataType         = remove_cvref_t<AccDataType_>;
+    using SmoothScaleDataType = remove_cvref_t<SmoothScaleDataType_>;
+    using YScaleDataType      = remove_cvref_t<YScaleDataType_>;
+    using ODataType           = remove_cvref_t<ODataType_>;
+    using BlockShape          = remove_cvref_t<BlockShape_>; // can consum generic 2d shape
+    using Traits              = remove_cvref_t<Traits_>;
 };
 
 // TODO: we should put descriptor creation function into policy
@@ -45,7 +45,7 @@ struct DynamicQuantEpilogue
 {
     using Problem                     = remove_cvref_t<Problem_>;
     using AccDataType                 = remove_cvref_t<typename Problem::AccDataType>;
-    using XScaleDataType              = remove_cvref_t<typename Problem::XScaleDataType>;
+    using SmoothScaleDataType         = remove_cvref_t<typename Problem::SmoothScaleDataType>;
     using YScaleDataType              = remove_cvref_t<typename Problem::YScaleDataType>;
     using ODataType                   = remove_cvref_t<typename Problem::ODataType>;
     using BlockShape                  = remove_cvref_t<typename Problem::BlockShape>;
@@ -78,7 +78,7 @@ struct DynamicQuantEpilogue
 #if 0
         // don't remove this
         // Note that if we set encoding purposely like this, you will result in compile fail
-        // TODO: x_scale create local-scratch to accept arbitrary acc input (with same length)
+        // TODO: sm_scale create local-scratch to accept arbitrary acc input (with same length)
         return make_static_tile_distribution(
             tile_distribution_encoding<
                 sequence<S::Repeat_M, S::WarpPerBlock_M, S::ThreadPerWarp_M>,
@@ -105,34 +105,18 @@ struct DynamicQuantEpilogue
         return reduce_crosswarp_sync.GetSmemSize();
     }
 
-    // TODO: this function assume store out vector size is the same as OAccTile last dimension size
-    //       how do we fix this ?
-    template <typename ODramWindowTmp,
-              typename XScaleWindow,
-              typename YScaleWindow,
-              typename OAccTile>
-    CK_TILE_DEVICE auto operator()(ODramWindowTmp& o_dram_window_tmp,
-                                   const XScaleWindow& x_scale_window_,
-                                   YScaleWindow& y_scale_window,
-                                   const OAccTile& o_acc_tile,
-                                   void* smem)
+    template <typename ODramWindowTmp, typename YScaleWindow, typename OAccTile>
+    CK_TILE_DEVICE auto Impl(ODramWindowTmp& o_dram_window_tmp,
+                             YScaleWindow& y_scale_window,
+                             const OAccTile& o_acc_tile,
+                             void* smem)
     {
         auto reduce                = GetBlockReduce2d();
         auto reduce_sync           = GetBlockReduce2dSync();
         auto reduce_crosswarp_sync = GetBlockReduce2dCrossWarpSync();
-        const auto x_scale_window =
-            make_tile_window(x_scale_window_, MakeSmoothInputScaleTileDistribution());
-
-        auto x_scale = load_tile(x_scale_window);
 
         auto o_acc_tmp = o_acc_tile;
 
-        sweep_tile(o_acc_tmp, [&](auto idx) {
-            constexpr auto j_idx = make_tuple(idx[number<1>{}]);
-            const auto xs_       = type_convert<AccDataType>(x_scale[j_idx]);
-            o_acc_tmp(idx)       = o_acc_tmp(idx) * xs_;
-        });
-
         const auto f_absmax = [](auto acc_, auto v_0_) { return max(acc_, abs(v_0_)); };
 
         auto row_absmax = [&]() {
@@ -184,5 +168,45 @@ struct DynamicQuantEpilogue
             store_tile(o_dram_window_tmp, cast_tile<ODataType>(o_acc_tmp));
         }
     }
+
+    // TODO: this function assume store out vector size is the same as OAccTile last dimension size
+    //       how do we fix this ?
+
+    // Smooth Dynamic Quant
+    template <typename ODramWindowTmp,
+              typename SmoothScaleWindow,
+              typename YScaleWindow,
+              typename OAccTile>
+    CK_TILE_DEVICE auto operator()(ODramWindowTmp& o_dram_window_tmp,
+                                   const SmoothScaleWindow& sm_scale_window_,
+                                   YScaleWindow& y_scale_window,
+                                   const OAccTile& o_acc_tile,
+                                   void* smem)
+    {
+        const auto sm_scale_window =
+            make_tile_window(sm_scale_window_, MakeSmoothInputScaleTileDistribution());
+
+        auto sm_scale = load_tile(sm_scale_window);
+
+        auto o_acc_tmp = o_acc_tile;
+
+        sweep_tile(o_acc_tmp, [&](auto idx) {
+            constexpr auto j_idx = make_tuple(idx[number<1>{}]);
+            const auto xs_       = type_convert<AccDataType>(sm_scale[j_idx]);
+            o_acc_tmp(idx)       = o_acc_tmp(idx) * xs_;
+        });
+
+        Impl(o_dram_window_tmp, y_scale_window, o_acc_tmp, smem);
+    }
+
+    // Dynamic Quant
+    template <typename ODramWindowTmp, typename YScaleWindow, typename OAccTile>
+    CK_TILE_DEVICE auto operator()(ODramWindowTmp& o_dram_window_tmp,
+                                   YScaleWindow& y_scale_window,
+                                   const OAccTile& o_acc_tile,
+                                   void* smem)
+    {
+        Impl(o_dram_window_tmp, y_scale_window, o_acc_tile, smem);
+    }
 };
 } // namespace ck_tile

include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp

@@ -234,10 +234,153 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_Base // for f16/bf16
 
     CK_TILE_HOST_DEVICE static constexpr ck_tile::index_t GetSmemSize()
     {
-        return 32 * (128 + 8) * sizeof(bf16_t);
+        // return 32 * (128 + 8) * sizeof(bf16_t);
+        return MakeLdsLoadDesc_A().get_element_space_size() * sizeof(bf16_t) * 2; // 2 lds buffers
     }
 };
 
+// clang-format off
+#define _EXPAND_ASM_ARGS_OUT_ONE_ACC        \
+            [s_loop_cnt]"+s"(loop_cnt),     \
+                [v_acc_0]"+v"(v_acc[0]),    \
+                [v_acc_1]"+v"(v_acc[1]),    \
+                [v_acc_2]"+v"(v_acc[2]),    \
+                [v_acc_3]"+v"(v_acc[3]),    \
+                [v_acc_4]"+v"(v_acc[4]),    \
+                [v_acc_5]"+v"(v_acc[5]),    \
+                [v_acc_6]"+v"(v_acc[6]),    \
+                [v_acc_7]"+v"(v_acc[7]),    \
+                [v_acc_8]"+v"(v_acc[8]),    \
+                [v_acc_9]"+v"(v_acc[9]),    \
+                [v_acc_10]"+v"(v_acc[10]),    \
+                [v_acc_11]"+v"(v_acc[11]),    \
+                [v_acc_12]"+v"(v_acc[12]),    \
+                [v_acc_13]"+v"(v_acc[13]),    \
+                [v_acc_14]"+v"(v_acc[14]),    \
+                [v_acc_15]"+v"(v_acc[15]),    \
+                [s_mem_]"+r"(smem)
+
+#define _EXPAND_ASM_ARGS_OUT_TWO_ACC        \
+            [s_loop_cnt]"+s"(loop_cnt),     \
+                [v_acc_0]"+v"(v_acc[0]),    \
+                [v_acc_1]"+v"(v_acc[1]),    \
+                [v_acc_2]"+v"(v_acc[2]),    \
+                [v_acc_3]"+v"(v_acc[3]),    \
+                [v_acc_4]"+v"(v_acc[4]),    \
+                [v_acc_5]"+v"(v_acc[5]),    \
+                [v_acc_6]"+v"(v_acc[6]),    \
+                [v_acc_7]"+v"(v_acc[7]),    \
+                [v_acc_8]"+v"(v_acc[8]),    \
+                [v_acc_9]"+v"(v_acc[9]),    \
+                [v_acc_10]"+v"(v_acc[10]),    \
+                [v_acc_11]"+v"(v_acc[11]),    \
+                [v_acc_12]"+v"(v_acc[12]),    \
+                [v_acc_13]"+v"(v_acc[13]),    \
+                [v_acc_14]"+v"(v_acc[14]),    \
+                [v_acc_15]"+v"(v_acc[15]),    \
+                [v_acc_16]"+v"(v_acc[16]),    \
+                [v_acc_17]"+v"(v_acc[17]),    \
+                [v_acc_18]"+v"(v_acc[18]),    \
+                [v_acc_19]"+v"(v_acc[19]),    \
+                [v_acc_20]"+v"(v_acc[20]),    \
+                [v_acc_21]"+v"(v_acc[21]),    \
+                [v_acc_22]"+v"(v_acc[22]),    \
+                [v_acc_23]"+v"(v_acc[23]),    \
+                [v_acc_24]"+v"(v_acc[24]),    \
+                [v_acc_25]"+v"(v_acc[25]),    \
+                [v_acc_26]"+v"(v_acc[26]),    \
+                [v_acc_27]"+v"(v_acc[27]),    \
+                [v_acc_28]"+v"(v_acc[28]),    \
+                [v_acc_29]"+v"(v_acc[29]),    \
+                [v_acc_30]"+v"(v_acc[30]),    \
+                [v_acc_31]"+v"(v_acc[31]),    \
+                [s_mem_]"+r"(smem)
+
+#define _EXPAND_ASM_ARGS_IN     \
+              [s_res_a0]"s"(res_a[0]),    \
+                [s_res_a1]"s"(res_a[1]),    \
+                [s_res_a2]"s"(res_a[2]),    \
+                [s_res_a3]"s"(res_a[3]),    \
+                [s_res_b0]"s"(res_b[0]),    \
+                [s_res_b1]"s"(res_b[1]),    \
+                [s_res_b2]"s"(res_b[2]),    \
+                [s_res_b3]"s"(res_b[3]),    \
+                [v_os_a0]"v"(static_cast<index_t>(cached_coords_a[number<0>{}] * sizeof(ADataType))),    \
+                [v_os_a1]"v"(static_cast<index_t>(cached_coords_a[number<1>{}] * sizeof(ADataType))),    \
+                [v_os_a2]"v"(static_cast<index_t>(cached_coords_a[number<2>{}] * sizeof(ADataType))),    \
+                [v_os_a3]"v"(static_cast<index_t>(cached_coords_a[number<3>{}] * sizeof(ADataType))),    \
+                [v_os_a4]"v"(static_cast<index_t>(cached_coords_a[number<4>{}] * sizeof(ADataType))),    \
+                [v_os_a5]"v"(static_cast<index_t>(cached_coords_a[number<5>{}] * sizeof(ADataType))),    \
+                [v_os_a6]"v"(static_cast<index_t>(cached_coords_a[number<6>{}] * sizeof(ADataType))),    \
+                [v_os_a7]"v"(static_cast<index_t>(cached_coords_a[number<7>{}] * sizeof(ADataType))),    \
+                                                                                                        \
+                [v_os_b0]"v"(static_cast<index_t>(cached_coords_b[number<0>{}] * sizeof(BDataType))),    \
+                [v_os_b1]"v"(static_cast<index_t>(cached_coords_b[number<1>{}] * sizeof(BDataType))),    \
+                [v_os_b2]"v"(static_cast<index_t>(cached_coords_b[number<2>{}] * sizeof(BDataType))),    \
+                [v_os_b3]"v"(static_cast<index_t>(cached_coords_b[number<3>{}] * sizeof(BDataType))),    \
+                [v_os_b4]"v"(static_cast<index_t>(cached_coords_b[number<4>{}] * sizeof(BDataType))),    \
+                [v_os_b5]"v"(static_cast<index_t>(cached_coords_b[number<5>{}] * sizeof(BDataType))),    \
+                [v_os_b6]"v"(static_cast<index_t>(cached_coords_b[number<6>{}] * sizeof(BDataType))),    \
+                [v_os_b7]"v"(static_cast<index_t>(cached_coords_b[number<7>{}] * sizeof(BDataType))),    \
+                                                                                                            \
+                [v_os_slda]"v"(static_cast<index_t>(a_sld.cached_coords_[number<0>{}].get_offset() * sizeof(ADataType))),\
+                [s_m0_init]"s"(m0_init_value),    \
+                [s_size_per_issue]"s"(size_per_issue),    \
+                [smem_sz]"n"(smem_buf_size),   \
+                [sld_os_0]"n"(sld_os[number<0>{}].value),    \
+                [sld_os_1]"n"(sld_os[number<1>{}].value),    \
+                [sld_os_2]"n"(sld_os[number<2>{}].value),    \
+                [sld_os_3]"n"(sld_os[number<3>{}].value),    \
+                [sld_os_4]"n"(sld_os[number<4>{}].value),    \
+                [sld_os_5]"n"(sld_os[number<5>{}].value),    \
+                [sld_os_6]"n"(sld_os[number<6>{}].value),    \
+                [sld_os_7]"n"(sld_os[number<7>{}].value),    \
+                [s_tile_os_a]"s"(tile_offset_a_bytes),    \
+                [s_tile_os_b]"s"(tile_offset_b_bytes)
+
+#define _EXPAND_ASM_ARGS_CLOBBER     \
+          "memory", "a0", "a1", "a2", "a3", "a4", "a5", "a6", "a7", "a8", "a9",    \
+          "a10", "a11", "a12", "a13", "a14", "a15", "a16", "a17", "a18", "a19",    \
+          "a20", "a21", "a22", "a23", "a24", "a25", "a26", "a27", "a28", "a29",    \
+          "a30", "a31", "a32", "a33", "a34", "a35", "a36", "a37", "a38", "a39",    \
+          "a40", "a41", "a42", "a43", "a44", "a45", "a46", "a47", "a48", "a49",    \
+          "a50", "a51", "a52", "a53", "a54", "a55", "a56", "a57", "a58", "a59",    \
+          "a60", "a61", "a62", "a63", "a64", "a65", "a66", "a67", "a68", "a69",    \
+          "a70", "a71", "a72", "a73", "a74", "a75", "a76", "a77", "a78", "a79",    \
+          "a80", "a81", "a82", "a83", "a84", "a85", "a86", "a87", "a88", "a89",    \
+          "a90", "a91", "a92", "a93", "a94", "a95", "a96", "a97", "a98", "a99",    \
+          "a100", "a101", "a102", "a103", "a104", "a105", "a106", "a107",    \
+          "a108", "a109", "a110", "a111", "a112", "a113", "a114", "a115",    \
+          "a116", "a117", "a118", "a119", "a120", "a121", "a122", "a123",    \
+          "a124", "a125", "a126", "a127", "a128", "a129", "a130", "a131",    \
+          "a132", "a133", "a134", "a135", "a136", "a137", "a138", "a139",    \
+          "a140", "a141", "a142", "a143", "a144", "a145", "a146", "a147",    \
+          "a148", "a149", "a150", "a151", "a152", "a153", "a154", "a155",    \
+          "a156", "a157", "a158", "a159", "a160", "a161", "a162", "a163",    \
+          "a164", "a165", "a166", "a167", "a168", "a169", "a170", "a171",    \
+          "a172", "a173", "a174", "a175", "a176", "a177", "a178", "a179",    \
+          "a180", "a181", "a182", "a183", "a184", "a185", "a186", "a187",    \
+          "a188", "a189", "a190", "a191", "a192", "a193", "a194", "a195",    \
+          "a196", "a197", "a198", "a199", "a200", "a201", "a202", "a203",    \
+          "a204", "a205", "a206", "a207", "a208", "a209", "a210", "a211",    \
+          "a212", "a213", "a214", "a215", "a216", "a217", "a218", "a219",    \
+          "a220", "a221", "a222", "a223", "a224", "a225", "a226", "a227",    \
+          "a228", "a229", "a230", "a231", "a232", "a233", "a234", "a235",    \
+          "a236", "a237", "a238", "a239", "a240", "a241", "a242", "a243",    \
+          "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",    \
+          "a252", "a253", "a254", "a255",     \
+          "s16", "s17", "s18", "s19", "s20", "s21", "s22", "s23",    \
+          "s86",                         \
+          "v64", "v65", "v66", "v67", "v68", "v69",                 \
+          "v70", "v71", "v72", "v73", "v74", "v75", "v76", "v77", "v78", "v79",     \
+          "v80", "v81", "v82", "v83", "v84", "v85", "v86", "v87", "v88", "v89",    \
+          "v90", "v91", "v92", "v93", "v94", "v95", "v96", "v97", "v98", "v99",    \
+          "v100", "v101", "v102", "v103", "v104", "v105", "v106", "v107",    \
+          "v108", "v109", "v110", "v111", "v112", "v113", "v114", "v115",    \
+          "v116", "v117", "v118", "v119", "v120", "v121", "v122", "v123",    \
+          "v124", "v125", "v126", "v127"
+// clang-format on
+
 struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16x16x32_Base
 {
     using ADataType = bf16_t;
@@ -245,7 +388,9 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16
 
     // TODO: need paired with tile_window_linear!
     // TODO: need call init_raw() before call this function!
-    template <typename ARes, typename ACoords, typename BRes, typename BCoords>
+    // Is2B: originally for B matrix we have 2 prefetch buffers. If set this to true
+    // we can support A matric serve 2 B matrix, B0/B1, each B0/B1 still have same tile size
+    template <typename ARes, typename ACoords, typename BRes, typename BCoords, bool Is2B = false>
     CK_TILE_DEVICE auto
     operator()(const ARes& res_a,
                const ACoords& cached_coords_a,
@@ -254,7 +399,8 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16
                CK_TILE_LDS_ADDR void* smem,
                index_t k,
                index_t tile_offset_a, // for each tile, the offset to move for each unroll
-               index_t tile_offset_b) // for each tile, the offset to move for each unroll
+               index_t tile_offset_b,
+               bool_constant<Is2B> = {}) // for each tile, the offset to move for each unroll
     {
         static_assert(ACoords::size() == Block_M * Block_K / BlockSize / 2 /*2x per dword*/); // 8
         static_assert(BCoords::size() == Repeat_N);
@@ -299,129 +445,78 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16
 
         index_t loop_cnt = k / Block_K;
 
-        // this is the acc thread buffer
-        fp32x4_t v_acc[16]{.0f};
+        if constexpr(Is2B)
+        {
+            // this is the acc thread buffer
+            fp32x4_t v_acc[32]{.0f};
 
-        // B nr->kr
+            // B nr->kr
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Winline-asm"
-        // clang-format off
-        asm volatile(
+            // clang-format off
+            asm volatile(
 #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16
+#define CK_TILE_FLATMM_UK_2B 1
 #include "uk/flatmm_uk_gfx9_32x512x128_1x1x1_16x16x16.inc"
-#undef CK_TILE_FLATMM_UK_MFMA
-            :   [s_loop_cnt]"+s"(loop_cnt),
-                [v_acc_0]"+v"(v_acc[0]),
-                [v_acc_1]"+v"(v_acc[1]),
-                [v_acc_2]"+v"(v_acc[2]),
-                [v_acc_3]"+v"(v_acc[3]),
-                [v_acc_4]"+v"(v_acc[4]),
-                [v_acc_5]"+v"(v_acc[5]),
-                [v_acc_6]"+v"(v_acc[6]),
-                [v_acc_7]"+v"(v_acc[7]),
-                [v_acc_8]"+v"(v_acc[8]),
-                [v_acc_9]"+v"(v_acc[9]),
-                [v_acc_10]"+v"(v_acc[10]),
-                [v_acc_11]"+v"(v_acc[11]),
-                [v_acc_12]"+v"(v_acc[12]),
-                [v_acc_13]"+v"(v_acc[13]),
-                [v_acc_14]"+v"(v_acc[14]),
-                [v_acc_15]"+v"(v_acc[15]),
-                [s_mem_]"+r"(smem)
-            : [s_res_a0]"s"(res_a[0]),
-                [s_res_a1]"s"(res_a[1]),
-                [s_res_a2]"s"(res_a[2]),
-                [s_res_a3]"s"(res_a[3]),
-                [s_res_b0]"s"(res_b[0]),
-                [s_res_b1]"s"(res_b[1]),
-                [s_res_b2]"s"(res_b[2]),
-                [s_res_b3]"s"(res_b[3]),
-                [v_os_a0]"v"(static_cast<index_t>(cached_coords_a[number<0>{}] * sizeof(ADataType))),
-                [v_os_a1]"v"(static_cast<index_t>(cached_coords_a[number<1>{}] * sizeof(ADataType))),
-                [v_os_a2]"v"(static_cast<index_t>(cached_coords_a[number<2>{}] * sizeof(ADataType))),
-                [v_os_a3]"v"(static_cast<index_t>(cached_coords_a[number<3>{}] * sizeof(ADataType))),
-                [v_os_a4]"v"(static_cast<index_t>(cached_coords_a[number<4>{}] * sizeof(ADataType))),
-                [v_os_a5]"v"(static_cast<index_t>(cached_coords_a[number<5>{}] * sizeof(ADataType))),
-                [v_os_a6]"v"(static_cast<index_t>(cached_coords_a[number<6>{}] * sizeof(ADataType))),
-                [v_os_a7]"v"(static_cast<index_t>(cached_coords_a[number<7>{}] * sizeof(ADataType))),
-
-                [v_os_b0]"v"(static_cast<index_t>(cached_coords_b[number<0>{}] * sizeof(BDataType))),
-                [v_os_b1]"v"(static_cast<index_t>(cached_coords_b[number<1>{}] * sizeof(BDataType))),
-                [v_os_b2]"v"(static_cast<index_t>(cached_coords_b[number<2>{}] * sizeof(BDataType))),
-                [v_os_b3]"v"(static_cast<index_t>(cached_coords_b[number<3>{}] * sizeof(BDataType))),
-                [v_os_b4]"v"(static_cast<index_t>(cached_coords_b[number<4>{}] * sizeof(BDataType))),
-                [v_os_b5]"v"(static_cast<index_t>(cached_coords_b[number<5>{}] * sizeof(BDataType))),
-                [v_os_b6]"v"(static_cast<index_t>(cached_coords_b[number<6>{}] * sizeof(BDataType))),
-                [v_os_b7]"v"(static_cast<index_t>(cached_coords_b[number<7>{}] * sizeof(BDataType))),
-
-                [v_os_slda]"v"(static_cast<index_t>(a_sld.cached_coords_[number<0>{}].get_offset() * sizeof(ADataType))),
-                [s_m0_init]"s"(m0_init_value),
-                [s_size_per_issue]"s"(size_per_issue),
-                [smem_sz]"n"(smem_buf_size),  //(smem_buf_size),
-                [sld_os_0]"n"(sld_os[number<0>{}].value),
-                [sld_os_1]"n"(sld_os[number<1>{}].value),
-                [sld_os_2]"n"(sld_os[number<2>{}].value),
-                [sld_os_3]"n"(sld_os[number<3>{}].value),
-                [sld_os_4]"n"(sld_os[number<4>{}].value),
-                [sld_os_5]"n"(sld_os[number<5>{}].value),
-                [sld_os_6]"n"(sld_os[number<6>{}].value),
-                [sld_os_7]"n"(sld_os[number<7>{}].value),
-                [s_tile_os_a]"s"(tile_offset_a_bytes),
-                [s_tile_os_b]"s"(tile_offset_b_bytes)
-            : "memory", "a0", "a1", "a2", "a3", "a4", "a5", "a6", "a7", "a8", "a9",
-          "a10", "a11", "a12", "a13", "a14", "a15", "a16", "a17", "a18", "a19",
-          "a20", "a21", "a22", "a23", "a24", "a25", "a26", "a27", "a28", "a29",
-          "a30", "a31", "a32", "a33", "a34", "a35", "a36", "a37", "a38", "a39",
-          "a40", "a41", "a42", "a43", "a44", "a45", "a46", "a47", "a48", "a49",
-          "a50", "a51", "a52", "a53", "a54", "a55", "a56", "a57", "a58", "a59",
-          "a60", "a61", "a62", "a63", "a64", "a65", "a66", "a67", "a68", "a69",
-          "a70", "a71", "a72", "a73", "a74", "a75", "a76", "a77", "a78", "a79",
-          "a80", "a81", "a82", "a83", "a84", "a85", "a86", "a87", "a88", "a89",
-          "a90", "a91", "a92", "a93", "a94", "a95", "a96", "a97", "a98", "a99",
-          "a100", "a101", "a102", "a103", "a104", "a105", "a106", "a107",
-          "a108", "a109", "a110", "a111", "a112", "a113", "a114", "a115",
-          "a116", "a117", "a118", "a119", "a120", "a121", "a122", "a123",
-          "a124", "a125", "a126", "a127", "a128", "a129", "a130", "a131",
-          "a132", "a133", "a134", "a135", "a136", "a137", "a138", "a139",
-          "a140", "a141", "a142", "a143", "a144", "a145", "a146", "a147",
-          "a148", "a149", "a150", "a151", "a152", "a153", "a154", "a155",
-          "a156", "a157", "a158", "a159", "a160", "a161", "a162", "a163",
-          "a164", "a165", "a166", "a167", "a168", "a169", "a170", "a171",
-          "a172", "a173", "a174", "a175", "a176", "a177", "a178", "a179",
-          "a180", "a181", "a182", "a183", "a184", "a185", "a186", "a187",
-          "a188", "a189", "a190", "a191", "a192", "a193", "a194", "a195",
-          "a196", "a197", "a198", "a199", "a200", "a201", "a202", "a203",
-          "a204", "a205", "a206", "a207", "a208", "a209", "a210", "a211",
-          "a212", "a213", "a214", "a215", "a216", "a217", "a218", "a219",
-          "a220", "a221", "a222", "a223", "a224", "a225", "a226", "a227",
-          "a228", "a229", "a230", "a231", "a232", "a233", "a234", "a235",
-          "a236", "a237", "a238", "a239", "a240", "a241", "a242", "a243",
-          "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",
-          "a252", "a253", "a254", "a255", 
-          "s16", "s17", "s18", "s19", "s20", "s21", "s22", "s23",
-          "s86",    // s86 as tmp
-          "v64", "v65", "v66", "v67", "v68", "v69",
-          "v70", "v71", "v72", "v73", "v74", "v75", "v76", "v77", "v78", "v79",
-          "v80", "v81", "v82", "v83", "v84", "v85", "v86", "v87", "v88", "v89",
-          "v90", "v91", "v92", "v93", "v94", "v95", "v96", "v97", "v98", "v99",
-          "v100", "v101", "v102", "v103", "v104", "v105", "v106", "v107",
-          "v108", "v109", "v110", "v111", "v112", "v113", "v114", "v115",
-          "v116", "v117", "v118", "v119", "v120", "v121", "v122", "v123",
-          "v124", "v125", "v126", "v127"
-        );
-        // clang-format on
+                : _EXPAND_ASM_ARGS_OUT_TWO_ACC
+                : _EXPAND_ASM_ARGS_IN, 
+                    [s_res_b4]"s"(res_b[4]), 
+                    [s_res_b5]"s"(res_b[5]),
+                    [s_res_b6]"s"(res_b[6]),
+                    [s_res_b7]"s"(res_b[7])
+                : _EXPAND_ASM_ARGS_CLOBBER, "s24", "s25", "s26", "s27"
+            );
+            // clang-format on
 #pragma clang diagnostic pop
 
-        // return local scratch
-        auto c = MakeCBlockTile();
-        for(auto i = 0; i < 16; i++)
+            // return local scratch
+            auto c = make_tuple(MakeCBlockTile(), MakeCBlockTile());
+            for(auto i = 0; i < 16; i++)
+            {
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 0] = v_acc[i].x;
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 1] = v_acc[i].y;
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 2] = v_acc[i].z;
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 3] = v_acc[i].w;
+            }
+            for(auto i = 0; i < 16; i++)
+            {
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 0] = v_acc[16 + i].x;
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 1] = v_acc[16 + i].y;
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 2] = v_acc[16 + i].z;
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 3] = v_acc[16 + i].w;
+            }
+            return c;
+        }
+        else
         {
-            c.get_thread_buffer()[4 * i + 0] = v_acc[i].x;
-            c.get_thread_buffer()[4 * i + 1] = v_acc[i].y;
-            c.get_thread_buffer()[4 * i + 2] = v_acc[i].z;
-            c.get_thread_buffer()[4 * i + 3] = v_acc[i].w;
+            // this is the acc thread buffer
+            fp32x4_t v_acc[16]{.0f};
+
+            // B nr->kr
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Winline-asm"
+            // clang-format off
+            asm volatile(
+#define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16
+#include "uk/flatmm_uk_gfx9_32x512x128_1x1x1_16x16x16.inc"
+                : _EXPAND_ASM_ARGS_OUT_ONE_ACC
+                : _EXPAND_ASM_ARGS_IN
+                : _EXPAND_ASM_ARGS_CLOBBER
+            );
+            // clang-format on
+#pragma clang diagnostic pop
+
+            // return local scratch
+            auto c = MakeCBlockTile();
+            for(auto i = 0; i < 16; i++)
+            {
+                c.get_thread_buffer()[4 * i + 0] = v_acc[i].x;
+                c.get_thread_buffer()[4 * i + 1] = v_acc[i].y;
+                c.get_thread_buffer()[4 * i + 2] = v_acc[i].z;
+                c.get_thread_buffer()[4 * i + 3] = v_acc[i].w;
+            }
+            return c;
         }
-        return c;
     }
 };
 
@@ -432,7 +527,7 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16
 
     // TODO: need paired with tile_window_linear!
     // TODO: need call init_raw() before call this function!
-    template <typename ARes, typename ACoords, typename BRes, typename BCoords>
+    template <typename ARes, typename ACoords, typename BRes, typename BCoords, bool Is2B = false>
     CK_TILE_DEVICE auto
     operator()(const ARes& res_a,
                const ACoords& cached_coords_a,
@@ -441,7 +536,8 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16
                CK_TILE_LDS_ADDR void* smem,
                index_t k,
                index_t tile_offset_a, // for each tile, the offset to move for each unroll
-               index_t tile_offset_b) // for each tile, the offset to move for each unroll
+               index_t tile_offset_b, // for each tile, the offset to move for each unroll
+               bool_constant<Is2B> = {})
     {
         static_assert(ACoords::size() == Block_M * Block_K / BlockSize / 2 /*2x per dword*/); // 8
         static_assert(BCoords::size() == Repeat_N);
@@ -486,130 +582,82 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16
 
         index_t loop_cnt = k / Block_K;
 
-        // this is the acc thread buffer
-        fp32x4_t v_acc[16]{.0f};
+        if constexpr(Is2B)
+        {
+            // this is the acc thread buffer
+            fp32x4_t v_acc[32]{.0f};
 
-        // B nr->kr
+            // B nr->kr
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Winline-asm"
-        // clang-format off
-        asm volatile(
+            // clang-format off
+            asm volatile(
 #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16
+#define CK_TILE_FLATMM_UK_2B 1
 #include "uk/flatmm_uk_gfx9_32x512x128_1x1x1_16x16x16.inc"
-#undef CK_TILE_FLATMM_UK_MFMA
-            :   [s_loop_cnt]"+s"(loop_cnt),
-                [v_acc_0]"+v"(v_acc[0]),
-                [v_acc_1]"+v"(v_acc[1]),
-                [v_acc_2]"+v"(v_acc[2]),
-                [v_acc_3]"+v"(v_acc[3]),
-                [v_acc_4]"+v"(v_acc[4]),
-                [v_acc_5]"+v"(v_acc[5]),
-                [v_acc_6]"+v"(v_acc[6]),
-                [v_acc_7]"+v"(v_acc[7]),
-                [v_acc_8]"+v"(v_acc[8]),
-                [v_acc_9]"+v"(v_acc[9]),
-                [v_acc_10]"+v"(v_acc[10]),
-                [v_acc_11]"+v"(v_acc[11]),
-                [v_acc_12]"+v"(v_acc[12]),
-                [v_acc_13]"+v"(v_acc[13]),
-                [v_acc_14]"+v"(v_acc[14]),
-                [v_acc_15]"+v"(v_acc[15]),
-                [s_mem_]"+r"(smem)
-            : [s_res_a0]"s"(res_a[0]),
-                [s_res_a1]"s"(res_a[1]),
-                [s_res_a2]"s"(res_a[2]),
-                [s_res_a3]"s"(res_a[3]),
-                [s_res_b0]"s"(res_b[0]),
-                [s_res_b1]"s"(res_b[1]),
-                [s_res_b2]"s"(res_b[2]),
-                [s_res_b3]"s"(res_b[3]),
-                [v_os_a0]"v"(static_cast<index_t>(cached_coords_a[number<0>{}] * sizeof(ADataType))),
-                [v_os_a1]"v"(static_cast<index_t>(cached_coords_a[number<1>{}] * sizeof(ADataType))),
-                [v_os_a2]"v"(static_cast<index_t>(cached_coords_a[number<2>{}] * sizeof(ADataType))),
-                [v_os_a3]"v"(static_cast<index_t>(cached_coords_a[number<3>{}] * sizeof(ADataType))),
-                [v_os_a4]"v"(static_cast<index_t>(cached_coords_a[number<4>{}] * sizeof(ADataType))),
-                [v_os_a5]"v"(static_cast<index_t>(cached_coords_a[number<5>{}] * sizeof(ADataType))),
-                [v_os_a6]"v"(static_cast<index_t>(cached_coords_a[number<6>{}] * sizeof(ADataType))),
-                [v_os_a7]"v"(static_cast<index_t>(cached_coords_a[number<7>{}] * sizeof(ADataType))),
-
-                [v_os_b0]"v"(static_cast<index_t>(cached_coords_b[number<0>{}] * sizeof(BDataType))),
-                [v_os_b1]"v"(static_cast<index_t>(cached_coords_b[number<1>{}] * sizeof(BDataType))),
-                [v_os_b2]"v"(static_cast<index_t>(cached_coords_b[number<2>{}] * sizeof(BDataType))),
-                [v_os_b3]"v"(static_cast<index_t>(cached_coords_b[number<3>{}] * sizeof(BDataType))),
-                [v_os_b4]"v"(static_cast<index_t>(cached_coords_b[number<4>{}] * sizeof(BDataType))),
-                [v_os_b5]"v"(static_cast<index_t>(cached_coords_b[number<5>{}] * sizeof(BDataType))),
-                [v_os_b6]"v"(static_cast<index_t>(cached_coords_b[number<6>{}] * sizeof(BDataType))),
-                [v_os_b7]"v"(static_cast<index_t>(cached_coords_b[number<7>{}] * sizeof(BDataType))),
-
-                [v_os_slda]"v"(static_cast<index_t>(a_sld.cached_coords_[number<0>{}].get_offset() * sizeof(ADataType))),
-                [s_m0_init]"s"(m0_init_value),
-                [s_size_per_issue]"s"(size_per_issue),
-                [smem_sz]"n"(smem_buf_size),  //(smem_buf_size),
-                [sld_os_0]"n"(sld_os[number<0>{}].value),
-                [sld_os_1]"n"(sld_os[number<1>{}].value),
-                [sld_os_2]"n"(sld_os[number<2>{}].value),
-                [sld_os_3]"n"(sld_os[number<3>{}].value),
-                [sld_os_4]"n"(sld_os[number<4>{}].value),
-                [sld_os_5]"n"(sld_os[number<5>{}].value),
-                [sld_os_6]"n"(sld_os[number<6>{}].value),
-                [sld_os_7]"n"(sld_os[number<7>{}].value),
-                [s_tile_os_a]"s"(tile_offset_a_bytes),
-                [s_tile_os_b]"s"(tile_offset_b_bytes)
-            : "memory", "a0", "a1", "a2", "a3", "a4", "a5", "a6", "a7", "a8", "a9",
-          "a10", "a11", "a12", "a13", "a14", "a15", "a16", "a17", "a18", "a19",
-          "a20", "a21", "a22", "a23", "a24", "a25", "a26", "a27", "a28", "a29",
-          "a30", "a31", "a32", "a33", "a34", "a35", "a36", "a37", "a38", "a39",
-          "a40", "a41", "a42", "a43", "a44", "a45", "a46", "a47", "a48", "a49",
-          "a50", "a51", "a52", "a53", "a54", "a55", "a56", "a57", "a58", "a59",
-          "a60", "a61", "a62", "a63", "a64", "a65", "a66", "a67", "a68", "a69",
-          "a70", "a71", "a72", "a73", "a74", "a75", "a76", "a77", "a78", "a79",
-          "a80", "a81", "a82", "a83", "a84", "a85", "a86", "a87", "a88", "a89",
-          "a90", "a91", "a92", "a93", "a94", "a95", "a96", "a97", "a98", "a99",
-          "a100", "a101", "a102", "a103", "a104", "a105", "a106", "a107",
-          "a108", "a109", "a110", "a111", "a112", "a113", "a114", "a115",
-          "a116", "a117", "a118", "a119", "a120", "a121", "a122", "a123",
-          "a124", "a125", "a126", "a127", "a128", "a129", "a130", "a131",
-          "a132", "a133", "a134", "a135", "a136", "a137", "a138", "a139",
-          "a140", "a141", "a142", "a143", "a144", "a145", "a146", "a147",
-          "a148", "a149", "a150", "a151", "a152", "a153", "a154", "a155",
-          "a156", "a157", "a158", "a159", "a160", "a161", "a162", "a163",
-          "a164", "a165", "a166", "a167", "a168", "a169", "a170", "a171",
-          "a172", "a173", "a174", "a175", "a176", "a177", "a178", "a179",
-          "a180", "a181", "a182", "a183", "a184", "a185", "a186", "a187",
-          "a188", "a189", "a190", "a191", "a192", "a193", "a194", "a195",
-          "a196", "a197", "a198", "a199", "a200", "a201", "a202", "a203",
-          "a204", "a205", "a206", "a207", "a208", "a209", "a210", "a211",
-          "a212", "a213", "a214", "a215", "a216", "a217", "a218", "a219",
-          "a220", "a221", "a222", "a223", "a224", "a225", "a226", "a227",
-          "a228", "a229", "a230", "a231", "a232", "a233", "a234", "a235",
-          "a236", "a237", "a238", "a239", "a240", "a241", "a242", "a243",
-          "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",
-          "a252", "a253", "a254", "a255", 
-          "s16", "s17", "s18", "s19", "s20", "s21", "s22", "s23",
-          "s86",    // s86 as tmp
-          "v64", "v65", "v66", "v67", "v68", "v69",
-          "v70", "v71", "v72", "v73", "v74", "v75", "v76", "v77", "v78", "v79",
-          "v80", "v81", "v82", "v83", "v84", "v85", "v86", "v87", "v88", "v89",
-          "v90", "v91", "v92", "v93", "v94", "v95", "v96", "v97", "v98", "v99",
-          "v100", "v101", "v102", "v103", "v104", "v105", "v106", "v107",
-          "v108", "v109", "v110", "v111", "v112", "v113", "v114", "v115",
-          "v116", "v117", "v118", "v119", "v120", "v121", "v122", "v123",
-          "v124", "v125", "v126", "v127"
-        );
-        // clang-format on
+                : _EXPAND_ASM_ARGS_OUT_TWO_ACC
+                : _EXPAND_ASM_ARGS_IN, 
+                    [s_res_b4]"s"(res_b[4]), 
+                    [s_res_b5]"s"(res_b[5]),
+                    [s_res_b6]"s"(res_b[6]),
+                    [s_res_b7]"s"(res_b[7])
+                : _EXPAND_ASM_ARGS_CLOBBER, "s24", "s25", "s26", "s27"
+            );
+            // clang-format on
 #pragma clang diagnostic pop
 
-        // return local scratch
-        auto c = MakeCBlockTile();
-        for(auto i = 0; i < 16; i++)
+            // return local scratch
+            auto c = make_tuple(MakeCBlockTile(), MakeCBlockTile());
+            for(auto i = 0; i < 16; i++)
+            {
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 0] = v_acc[i].x;
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 1] = v_acc[i].y;
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 2] = v_acc[i].z;
+                c.at(number<0>{}).get_thread_buffer()[4 * i + 3] = v_acc[i].w;
+            }
+            for(auto i = 0; i < 16; i++)
+            {
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 0] = v_acc[16 + i].x;
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 1] = v_acc[16 + i].y;
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 2] = v_acc[16 + i].z;
+                c.at(number<1>{}).get_thread_buffer()[4 * i + 3] = v_acc[16 + i].w;
+            }
+            return c;
+        }
+        else
         {
-            c.get_thread_buffer()[4 * i + 0] = v_acc[i].x;
-            c.get_thread_buffer()[4 * i + 1] = v_acc[i].y;
-            c.get_thread_buffer()[4 * i + 2] = v_acc[i].z;
-            c.get_thread_buffer()[4 * i + 3] = v_acc[i].w;
+            // this is the acc thread buffer
+            fp32x4_t v_acc[16]{.0f};
+
+            // B nr->kr
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Winline-asm"
+            // clang-format off
+            asm volatile(
+#define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16
+#include "uk/flatmm_uk_gfx9_32x512x128_1x1x1_16x16x16.inc"
+                : _EXPAND_ASM_ARGS_OUT_ONE_ACC
+                : _EXPAND_ASM_ARGS_IN
+                : _EXPAND_ASM_ARGS_CLOBBER
+            );
+            // clang-format on
+#pragma clang diagnostic pop
+
+            // return local scratch
+            auto c = MakeCBlockTile();
+            for(auto i = 0; i < 16; i++)
+            {
+                c.get_thread_buffer()[4 * i + 0] = v_acc[i].x;
+                c.get_thread_buffer()[4 * i + 1] = v_acc[i].y;
+                c.get_thread_buffer()[4 * i + 2] = v_acc[i].z;
+                c.get_thread_buffer()[4 * i + 3] = v_acc[i].w;
+            }
+            return c;
         }
-        return c;
     }
 };
-
+#undef _EXPAND_ASM_ARGS_OUT_ONE_ACC
+#undef _EXPAND_ASM_ARGS_OUT_TWO_ACC
+#undef _EXPAND_ASM_ARGS_IN
+#undef _EXPAND_ASM_ARGS_CLOBBER
 } // namespace ck_tile

include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp

@@ -65,7 +65,8 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_Base
         // in LDS we need store as
         //          M0(2)* N0(2) *  Nl(4) * Nw(4) * (Mw(16)*Nv(4) + 4)
         //             y    y       wave-id  lid/16  lid%16   v
-        return 2 * 2 * 4 * 4 * (16 * 4 + 4) * sizeof(bf16_t);
+        constexpr index_t nbufs = 2;
+        return 2 * 2 * 4 * 4 * (16 * 4 + 4) * sizeof(bf16_t) * nbufs;
     }
 };
 
@@ -173,7 +174,6 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 : public FlatmmSn_32x128x512_1x4x
         asm volatile(
 #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16
 #include "uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16.inc"
-#undef CK_TILE_FLATMM_UK_MFMA
             :[smem_]"+r"(smem),
             [s_loop_cnt]"+s"(loop_cnt),
                 [c0]"+v" (v_c0),
@@ -418,7 +418,6 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_FP16 : public FlatmmSn_32x128x512_1x4x
         asm volatile(
 #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16
 #include "uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16.inc"
-#undef CK_TILE_FLATMM_UK_MFMA
             :[smem_]"+r"(smem),
             [s_loop_cnt]"+s"(loop_cnt),
                 [c0]"+v" (v_c0),

include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp

@@ -477,7 +477,7 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_FP16_itl : public FlatmmSn_32x128x512_
           "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",
           "a252", "a253", "a254", "a255", 
           "s8", "s9", "s12", "s13", "s14", "s15", "s38", "s39", "s52", "s86",
-          "s36", "s37","s59","s80",
+          "s36", "s37", "s56", "s59", "s60", "s80",
           "v10", "v11", "v12", "v13", "v14", "v15", "v16", "v17",
           "v50", "v54", "v55",
           "v64","v65","v66","v67","v68","v69","v70","v71",

include/ck_tile/ops/flatmm/block/uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16.inc

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+// clang-format off
+
+// define the CK_TILE_** macro before include this file to change kernel variation
+// we will undef everything defined in this file
+
 #ifndef CK_TILE_FLATMM_UK_MFMA
 #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16
 #endif
@@ -816,3 +823,5 @@
 #undef _UK_MFMA_
 #undef _UK_PK_CVT_
 #undef _UK_ATOMIC_ADD_
+#undef CK_TILE_FLATMM_UK_MFMA
+// clang-format on

include/ck_tile/ops/flatmm/block/uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16_itl.inc

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
+// clang-format off
+
+// define the CK_TILE_** macro before include this file to change kernel variation
+// we will undef everything defined in this file
+
 #ifndef CK_TILE_FLATMM_UK_MFMA
 #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16
 #endif
@@ -159,199 +166,208 @@
     " ds_write_b64 %[v_sfl_sst], [%[c6],%[c7]] offset:23168  \n"
     " s_mov_b32 s80, 0 \n"
     " s_waitcnt vmcnt(24) \n"
-    "label_0AA6: \n"
+    "coreloop_top_%=: \n"
     " s_waitcnt vmcnt(30) & lgkmcnt(0)  \n"
-    " s_barrier \n" _UK_MFMA_ " [%[c0], %[c1], %[c2], %[c3]], acc[0:1], v[128:129], 0  \n"
+    " s_barrier \n"
+    _UK_MFMA_ " [%[c0], %[c1], %[c2], %[c3]], acc[0:1], v[128:129], 0  \n"
     " ds_read_b32 v10, %[v_sfl_sld] offset:16640  \n"
     " ds_read_b32 v11, %[v_sfl_sld] offset:16672  \n"
-    " ds_write_b64 %[v_sfl_sst], [%[c16],%[c17]] offset:25344  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[2:3], v[130:131], v[64:67]  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c16],%[c17]] offset:25344  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[2:3], v[130:131], v[64:67]  \n"
     " buffer_load_dwordx4 acc[128:131],  %[v_os_b0], s[12:15], 0 offen  \n"
-    " ds_write_b64 %[v_sfl_sst], [%[c18],%[c19]] offset:29696  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[4:5], v[132:133], v[64:67]  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c18],%[c19]] offset:29696  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[4:5], v[132:133], v[64:67]  \n"
     " ds_read_b32 v12, %[v_sfl_sld] offset:16704  \n"
     " ds_read_b32 v13, %[v_sfl_sld] offset:16736  \n"
-    " ds_write_b64 %[v_sfl_sst], [%[c20],%[c21]] offset:27520  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[6:7], v[134:135], v[64:67]  \n"
-    " ds_write_b64 %[v_sfl_sst], [%[c22],%[c23]] offset:31872  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[8:9], v[136:137], v[64:67]  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c20],%[c21]] offset:27520  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[6:7], v[134:135], v[64:67]  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c22],%[c23]] offset:31872  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[8:9], v[136:137], v[64:67]  \n"
     " ds_read_b32 v14, %[v_sfl_sld] offset:20992  \n"
-    " ds_read_b32 v15, %[v_sfl_sld] offset:21024  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[10:11], v[138:139], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[132:135],  %[v_os_b0], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[12:13], v[140:141], v[64:67]  \n"
+    " ds_read_b32 v15, %[v_sfl_sld] offset:21024  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[10:11], v[138:139], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[132:135],  %[v_os_b0], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[12:13], v[140:141], v[64:67]  \n"
     " ds_read_b32 v16, %[v_sfl_sld] offset:21056  \n"
-    " ds_read_b32 v17, %[v_sfl_sld] offset:21088  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[14:15], v[142:143], v[64:67]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[0:1], v[192:193], 0  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[2:3], v[194:195], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[136:139],  %[v_os_b0], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[4:5], v[196:197], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[6:7], v[198:199], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[8:9], v[200:201], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[10:11], v[202:203], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[140:143],  %[v_os_b0], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[12:13], v[204:205], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[14:15], v[206:207], v[68:71]  \n"
+    " ds_read_b32 v17, %[v_sfl_sld] offset:21088  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[14:15], v[142:143], v[64:67]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[0:1], v[192:193], 0  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[2:3], v[194:195], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[136:139],  %[v_os_b0], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[4:5], v[196:197], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[6:7], v[198:199], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[8:9], v[200:201], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[10:11], v[202:203], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[140:143],  %[v_os_b0], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[12:13], v[204:205], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[14:15], v[206:207], v[68:71]  \n"
     " s_waitcnt lgkmcnt(0) \n"
-    " s_mov_b64 exec, %[s_execflag_0] \n" _UK_ATOMIC_ADD_ " %[v_os_o0], v10, s[8:9] \n"
-    "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[16:17], v[128:129], 0  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[18:19], v[130:131], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[144:147],  %[v_os_b1], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[20:21], v[132:133], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[22:23], v[134:135], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[24:25], v[136:137], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[26:27], v[138:139], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[148:151],  %[v_os_b1], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[28:29], v[140:141], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[30:31], v[142:143], v[72:75]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[16:17], v[192:193], 0  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[18:19], v[194:195], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[152:155],  %[v_os_b1], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[20:21], v[196:197], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[22:23], v[198:199], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[24:25], v[200:201], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[26:27], v[202:203], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[156:159],  %[v_os_b1], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[28:29], v[204:205], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[30:31], v[206:207], v[76:79]  \n"
-    " s_mov_b64 exec, %[s_execflag_1] \n" _UK_ATOMIC_ADD_ " %[v_os_o1], v11, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_0] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o0], v10, s[8:9] \n"
+    "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[16:17], v[128:129], 0  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[18:19], v[130:131], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[144:147],  %[v_os_b1], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[20:21], v[132:133], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[22:23], v[134:135], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[24:25], v[136:137], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[26:27], v[138:139], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[148:151],  %[v_os_b1], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[28:29], v[140:141], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[30:31], v[142:143], v[72:75]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[16:17], v[192:193], 0  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[18:19], v[194:195], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[152:155],  %[v_os_b1], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[20:21], v[196:197], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[22:23], v[198:199], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[24:25], v[200:201], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[26:27], v[202:203], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[156:159],  %[v_os_b1], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[28:29], v[204:205], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[30:31], v[206:207], v[76:79]  \n"
+    " s_mov_b64 exec, %[s_execflag_1] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o1], v11, s[8:9] \n"
     "  s_mov_b64     exec, s[38:39]  \n"
-    " s_waitcnt vmcnt(30)   \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[32:33], v[144:145], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[34:35], v[146:147], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[160:163],  %[v_os_b2], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[36:37], v[148:149], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[38:39], v[150:151], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[40:41], v[152:153], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[42:43], v[154:155], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[164:167],  %[v_os_b2], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[44:45], v[156:157], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[46:47], v[158:159], v[64:67]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[32:33], v[208:209], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[34:35], v[210:211], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[168:171],  %[v_os_b2], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[36:37], v[212:213], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[38:39], v[214:215], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[40:41], v[216:217], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[42:43], v[218:219], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[172:175],  %[v_os_b2], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[44:45], v[220:221], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[46:47], v[222:223], v[68:71]  \n"
-    " s_mov_b64 exec, %[s_execflag_2] \n" _UK_ATOMIC_ADD_ " %[v_os_o2], v12, s[8:9] \n"
-    "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[48:49], v[144:145], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[50:51], v[146:147], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[176:179],  %[v_os_b3], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[52:53], v[148:149], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[54:55], v[150:151], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[56:57], v[152:153], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[58:59], v[154:155], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[180:183],  %[v_os_b3], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[60:61], v[156:157], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[62:63], v[158:159], v[72:75]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[48:49], v[208:209], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[50:51], v[210:211], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[184:187],  %[v_os_b3], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[52:53], v[212:213], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[54:55], v[214:215], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[56:57], v[216:217], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[58:59], v[218:219], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[188:191],  %[v_os_b3], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[60:61], v[220:221], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[62:63], v[222:223], v[76:79]  \n"
-    " s_mov_b64 exec, %[s_execflag_3] \n" _UK_ATOMIC_ADD_ " %[v_os_o3], v13, s[8:9] \n"
+    " s_waitcnt vmcnt(30)   \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[32:33], v[144:145], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[34:35], v[146:147], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[160:163],  %[v_os_b2], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[36:37], v[148:149], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[38:39], v[150:151], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[40:41], v[152:153], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[42:43], v[154:155], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[164:167],  %[v_os_b2], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[44:45], v[156:157], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[46:47], v[158:159], v[64:67]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[32:33], v[208:209], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[34:35], v[210:211], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[168:171],  %[v_os_b2], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[36:37], v[212:213], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[38:39], v[214:215], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[40:41], v[216:217], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[42:43], v[218:219], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[172:175],  %[v_os_b2], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[44:45], v[220:221], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[46:47], v[222:223], v[68:71]  \n"
+    " s_mov_b64 exec, %[s_execflag_2] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o2], v12, s[8:9] \n"
+    "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[48:49], v[144:145], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[50:51], v[146:147], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[176:179],  %[v_os_b3], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[52:53], v[148:149], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[54:55], v[150:151], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[56:57], v[152:153], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[58:59], v[154:155], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[180:183],  %[v_os_b3], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[60:61], v[156:157], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[62:63], v[158:159], v[72:75]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[48:49], v[208:209], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[50:51], v[210:211], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[184:187],  %[v_os_b3], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[52:53], v[212:213], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[54:55], v[214:215], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[56:57], v[216:217], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[58:59], v[218:219], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[188:191],  %[v_os_b3], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[60:61], v[220:221], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[62:63], v[222:223], v[76:79]  \n"
+    " s_mov_b64 exec, %[s_execflag_3] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o3], v13, s[8:9] \n"
     "  s_mov_b64     exec, s[38:39]    \n"
-    " s_waitcnt vmcnt(30)   \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[64:65], v[160:161], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[66:67], v[162:163], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[192:195],  %[v_os_b4], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[68:69], v[164:165], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[70:71], v[166:167], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[72:73], v[168:169], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[74:75], v[170:171], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[196:199],  %[v_os_b4], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[76:77], v[172:173], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[78:79], v[174:175], v[64:67]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[64:65], v[224:225], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[66:67], v[226:227], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[200:203],  %[v_os_b4], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[68:69], v[228:229], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[70:71], v[230:231], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[72:73], v[232:233], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[74:75], v[234:235], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[204:207],  %[v_os_b4], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[76:77], v[236:237], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[78:79], v[238:239], v[68:71]  \n"
-    " s_mov_b64 exec, %[s_execflag_4] \n" _UK_ATOMIC_ADD_ " %[v_os_o4], v14, s[8:9] \n"
-    "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[80:81], v[160:161], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[82:83], v[162:163], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[208:211],  %[v_os_b5], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[84:85], v[164:165], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[86:87], v[166:167], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[88:89], v[168:169], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[90:91], v[170:171], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[212:215],  %[v_os_b5], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[92:93], v[172:173], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[94:95], v[174:175], v[72:75]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[80:81], v[224:225], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[82:83], v[226:227], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[216:219],  %[v_os_b5], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[84:85], v[228:229], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[86:87], v[230:231], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[88:89], v[232:233], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[90:91], v[234:235], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[220:223],  %[v_os_b5], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[92:93], v[236:237], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[94:95], v[238:239], v[76:79]  \n"
-    " s_mov_b64 exec, %[s_execflag_5] \n" _UK_ATOMIC_ADD_ " %[v_os_o5], v15, s[8:9] \n"
+    " s_waitcnt vmcnt(30)   \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[64:65], v[160:161], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[66:67], v[162:163], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[192:195],  %[v_os_b4], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[68:69], v[164:165], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[70:71], v[166:167], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[72:73], v[168:169], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[74:75], v[170:171], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[196:199],  %[v_os_b4], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[76:77], v[172:173], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[78:79], v[174:175], v[64:67]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[64:65], v[224:225], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[66:67], v[226:227], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[200:203],  %[v_os_b4], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[68:69], v[228:229], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[70:71], v[230:231], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[72:73], v[232:233], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[74:75], v[234:235], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[204:207],  %[v_os_b4], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[76:77], v[236:237], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[78:79], v[238:239], v[68:71]  \n"
+    " s_mov_b64 exec, %[s_execflag_4] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o4], v14, s[8:9] \n"
     "  s_mov_b64     exec, s[38:39]                           \n"
-    " s_waitcnt vmcnt(30)   \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[96:97], v[176:177], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[98:99], v[178:179], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[224:227],  %[v_os_b6], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[100:101], v[180:181], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[102:103], v[182:183], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[104:105], v[184:185], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[106:107], v[186:187], v[64:67]  \n"
-    " buffer_load_dwordx4 acc[228:231],  %[v_os_b6], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[108:109], v[188:189], v[64:67]  \n" _UK_MFMA_
-    " [%[c0], %[c1], %[c2], %[c3]], acc[110:111], v[190:191], v[64:67]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[96:97], v[240:241], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[98:99], v[242:243], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[232:235],  %[v_os_b6], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[100:101], v[244:245], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[102:103], v[246:247], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[104:105], v[248:249], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[106:107], v[250:251], v[68:71]  \n"
-    " buffer_load_dwordx4 acc[236:239],  %[v_os_b6], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[108:109], v[252:253], v[68:71]  \n" _UK_MFMA_
-    " [%[c4], %[c5], %[c6], %[c7]], acc[110:111], v[254:255], v[68:71]  \n"
-    " s_mov_b64 exec, %[s_execflag_6] \n" _UK_ATOMIC_ADD_ " %[v_os_o6], v16, s[8:9] \n"
-    "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[112:113], v[176:177], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[114:115], v[178:179], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[240:243],  %[v_os_b7], s[12:15], 0 offen  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[116:117], v[180:181], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[118:119], v[182:183], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[120:121], v[184:185], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[122:123], v[186:187], v[72:75]  \n"
-    " buffer_load_dwordx4 acc[244:247],  %[v_os_b7], s[12:15], 0 offen offset:1024  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[124:125], v[188:189], v[72:75]  \n" _UK_MFMA_
-    " [%[c8], %[c9], %[c10], %[c11]], acc[126:127], v[190:191], v[72:75]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[112:113], v[240:241], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[114:115], v[242:243], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[248:251],  %[v_os_b7], s[12:15], 0 offen offset:2048  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[116:117], v[244:245], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[118:119], v[246:247], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[120:121], v[248:249], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[122:123], v[250:251], v[76:79]  \n"
-    " buffer_load_dwordx4 acc[252:255],  %[v_os_b7], s[12:15], 0 offen offset:3072  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[124:125], v[252:253], v[76:79]  \n" _UK_MFMA_
-    " [%[c12], %[c13], %[c14], %[c15]], acc[126:127], v[254:255], v[76:79]  \n"
-    " s_mov_b64 exec, %[s_execflag_7] \n" _UK_ATOMIC_ADD_ " %[v_os_o7], v17, s[8:9] \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[80:81], v[160:161], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[82:83], v[162:163], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[208:211],  %[v_os_b5], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[84:85], v[164:165], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[86:87], v[166:167], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[88:89], v[168:169], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[90:91], v[170:171], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[212:215],  %[v_os_b5], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[92:93], v[172:173], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[94:95], v[174:175], v[72:75]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[80:81], v[224:225], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[82:83], v[226:227], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[216:219],  %[v_os_b5], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[84:85], v[228:229], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[86:87], v[230:231], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[88:89], v[232:233], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[90:91], v[234:235], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[220:223],  %[v_os_b5], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[92:93], v[236:237], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[94:95], v[238:239], v[76:79]  \n"
+    " s_mov_b64 exec, %[s_execflag_5] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o5], v15, s[8:9] \n"
+    "  s_mov_b64     exec, s[38:39]                           \n"
+    " s_waitcnt vmcnt(30)   \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[96:97], v[176:177], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[98:99], v[178:179], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[224:227],  %[v_os_b6], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[100:101], v[180:181], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[102:103], v[182:183], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[104:105], v[184:185], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[106:107], v[186:187], v[64:67]  \n"
+    " buffer_load_dwordx4 acc[228:231],  %[v_os_b6], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[108:109], v[188:189], v[64:67]  \n"
+    _UK_MFMA_    " [%[c0], %[c1], %[c2], %[c3]], acc[110:111], v[190:191], v[64:67]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[96:97], v[240:241], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[98:99], v[242:243], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[232:235],  %[v_os_b6], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[100:101], v[244:245], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[102:103], v[246:247], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[104:105], v[248:249], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[106:107], v[250:251], v[68:71]  \n"
+    " buffer_load_dwordx4 acc[236:239],  %[v_os_b6], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[108:109], v[252:253], v[68:71]  \n"
+    _UK_MFMA_    " [%[c4], %[c5], %[c6], %[c7]], acc[110:111], v[254:255], v[68:71]  \n"
+    " s_mov_b64 exec, %[s_execflag_6] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o6], v16, s[8:9] \n"
+    "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[112:113], v[176:177], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[114:115], v[178:179], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[240:243],  %[v_os_b7], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[116:117], v[180:181], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[118:119], v[182:183], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[120:121], v[184:185], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[122:123], v[186:187], v[72:75]  \n"
+    " buffer_load_dwordx4 acc[244:247],  %[v_os_b7], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[124:125], v[188:189], v[72:75]  \n"
+    _UK_MFMA_    " [%[c8], %[c9], %[c10], %[c11]], acc[126:127], v[190:191], v[72:75]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[112:113], v[240:241], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[114:115], v[242:243], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[248:251],  %[v_os_b7], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[116:117], v[244:245], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[118:119], v[246:247], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[120:121], v[248:249], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[122:123], v[250:251], v[76:79]  \n"
+    " buffer_load_dwordx4 acc[252:255],  %[v_os_b7], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[124:125], v[252:253], v[76:79]  \n"
+    _UK_MFMA_    " [%[c12], %[c13], %[c14], %[c15]], acc[126:127], v[254:255], v[76:79]  \n"
+    " s_mov_b64 exec, %[s_execflag_7] \n"
+    _UK_ATOMIC_ADD_ " %[v_os_o7], v17, s[8:9] \n"
     "  s_mov_b64     exec, s[38:39]                           \n"
     " s_add_u32 s60, 0x00000100, s80  \n"
     " s_cmp_lt_u32 s60, %[s_loop_cnt]  \n"
@@ -377,393 +393,333 @@
     " v_mul_f32 %[c12], %[scale_1], %[c12] \n"
     " v_mul_f32 %[c13], %[scale_1], %[c13] \n"
     " v_mul_f32 %[c14], %[scale_1], %[c14] \n"
-    " v_mul_f32 %[c15], %[scale_1], %[c15] \n" _UK_PK_CVT_("%[c0]", "%[c1]", "%[c0]") _UK_PK_CVT_(
-        "%[c2]",
-        "%[c3]",
-        "%[c1]") _UK_PK_CVT_("%[c4]",
-                             "%[c5]",
-                             "%[c2]") _UK_PK_CVT_("%[c6]",
-                                                  "%[c7]",
-                                                  "%[c3]") _UK_PK_CVT_("%[c8]",
-                                                                       "%[c9]",
-                                                                       "%[c4]") _UK_PK_CVT_("%["
-                                                                                            "c10]",
-                                                                                            "%["
-                                                                                            "c11]",
-                                                                                            "%[c5]")
-        _UK_PK_CVT_("%[c12]", "%[c13]", "%[c6]") _UK_PK_CVT_(
-            "%[c14]",
-            "%[c15]",
-            "%[c7]") " s_addk_i32 s80, 0x0080  \n"
-                     " s_cmp_lt_i32 s80, %[s_loop_cnt]  \n"
-                     " s_cbranch_scc0 label_0EC1  \n"
-                     " s_waitcnt vmcnt(30) & lgkmcnt(0)  \n"
-                     " s_barrier  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[128:129], v[128:129], 0  \n"
-                     " ds_read_b32 v10, %[v_sfl_sld] offset:25344  \n"
-                     " ds_read_b32 v11, %[v_sfl_sld] offset:25376  \n"
-                     " ds_write_b64 v3, v[64:65] offset:16640  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[130:131], v[130:131], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[0:3],  %[v_os_b0], s[12:15], 0 offen  \n"
-                     " ds_write_b64 v3, v[66:67] offset:20992  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[132:133], v[132:133], v[80:83]  \n"
-                     " ds_read_b32 v12, %[v_sfl_sld] offset:25408  \n"
-                     " ds_read_b32 v13, %[v_sfl_sld] offset:25440  \n"
-                     " ds_write_b64 v3, v[68:69] offset:18816  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[134:135], v[134:135], v[80:83]  \n"
-                     " ds_write_b64 v3, v[70:71] offset:23168  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[136:137], v[136:137], v[80:83]  \n"
-                     " ds_read_b32 v14, %[v_sfl_sld] offset:29696  \n"
-                     " ds_read_b32 v15, %[v_sfl_sld] offset:29728  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[138:139], v[138:139], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[4:7],  %[v_os_b0], s[12:15], 0 offen offset:1024  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[140:141], v[140:141], v[80:83]  \n"
-                     " ds_read_b32 v16, %[v_sfl_sld] offset:29760  \n"
-                     " ds_read_b32 v17, %[v_sfl_sld] offset:29792  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[142:143], v[142:143], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[128:129], v[192:193], 0  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[130:131], v[194:195], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[8:11],  %[v_os_b0], s[12:15], 0 offen offset:2048  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[132:133], v[196:197], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[134:135], v[198:199], v[84:87]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[136:137], v[200:201], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[138:139], v[202:203], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[12:15],  %[v_os_b0], s[12:15], 0 offen offset:3072  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[140:141], v[204:205], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[142:143], v[206:207], v[84:87]  \n"
-                     " s_waitcnt lgkmcnt(0) \n"
-                     " s_mov_b64 exec, %[s_execflag_0] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o0], v10, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[144:145], v[128:129], 0  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[146:147], v[130:131], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[16:19],  %[v_os_b1], s[12:15], 0 offen  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[148:149], v[132:133], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[150:151], v[134:135], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[152:153], v[136:137], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[154:155], v[138:139], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[20:23],  %[v_os_b1], s[12:15], 0 offen offset:1024  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[156:157], v[140:141], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[158:159], v[142:143], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[144:145], v[192:193], 0  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[146:147], v[194:195], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[24:27],  %[v_os_b1], s[12:15], 0 offen offset:2048  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[148:149], v[196:197], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[150:151], v[198:199], v[92:95]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[152:153], v[200:201], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[154:155], v[202:203], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[28:31],  %[v_os_b1], s[12:15], 0 offen offset:3072  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[156:157], v[204:205], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[158:159], v[206:207], v[92:95]  \n"
-                     " s_mov_b64 exec, %[s_execflag_1] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o1], v11, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n"
-                     " s_waitcnt vmcnt(30) \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[160:161], v[144:145], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[162:163], v[146:147], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[32:35],  %[v_os_b2], s[12:15], 0 offen  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[164:165], v[148:149], v[80:83]  "
-                     "\n" _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[166:167], v[150:151], "
-                     "v[80:83]  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[168:169], v[152:153], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[170:171], v[154:155], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[36:39],  %[v_os_b2], s[12:15], 0 offen offset:1024  "
-                     "\n" _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[172:173], v[156:157], "
-                     "v[80:83]  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[174:175], v[158:159], v[80:83]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[160:161], v[208:209], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[162:163], v[210:211], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[40:43],  %[v_os_b2], s[12:15], 0 offen offset:2048  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[164:165], v[212:213], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[166:167], v[214:215], v[84:87]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[168:169], v[216:217], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[170:171], v[218:219], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[44:47],  %[v_os_b2], s[12:15], 0 offen offset:3072  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[172:173], v[220:221], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[174:175], v[222:223], v[84:87]  \n"
-                     " s_mov_b64 exec, %[s_execflag_2] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o2], v12, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[176:177], v[144:145], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[178:179], v[146:147], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[48:51],  %[v_os_b3], s[12:15], 0 offen  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[180:181], v[148:149], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[182:183], v[150:151], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[184:185], v[152:153], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[186:187], v[154:155], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[52:55],  %[v_os_b3], s[12:15], 0 offen offset:1024  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[188:189], v[156:157], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[190:191], v[158:159], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[176:177], v[208:209], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[178:179], v[210:211], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[56:59],  %[v_os_b3], s[12:15], 0 offen offset:2048  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[180:181], v[212:213], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[182:183], v[214:215], v[92:95]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[184:185], v[216:217], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[186:187], v[218:219], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[60:63],  %[v_os_b3], s[12:15], 0 offen offset:3072  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[188:189], v[220:221], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[190:191], v[222:223], v[92:95]  \n"
-                     " s_mov_b64 exec, %[s_execflag_3] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o3], v13, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n"
-                     " s_waitcnt vmcnt(30) \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[192:193], v[160:161], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[194:195], v[162:163], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[64:67],  %[v_os_b4], s[12:15], 0 offen  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[196:197], v[164:165], v[80:83]  "
-                     "\n" _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[198:199], v[166:167], "
-                     "v[80:83]  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[200:201], v[168:169], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[202:203], v[170:171], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[68:71],  %[v_os_b4], s[12:15], 0 offen offset:1024  "
-                     "\n" _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[204:205], v[172:173], "
-                     "v[80:83]  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[206:207], v[174:175], v[80:83]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[192:193], v[224:225], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[194:195], v[226:227], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[72:75],  %[v_os_b4], s[12:15], 0 offen offset:2048  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[196:197], v[228:229], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[198:199], v[230:231], v[84:87]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[200:201], v[232:233], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[202:203], v[234:235], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[76:79],  %[v_os_b4], s[12:15], 0 offen offset:3072  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[204:205], v[236:237], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[206:207], v[238:239], v[84:87]  \n"
-                     " s_mov_b64 exec, %[s_execflag_4] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o4], v14, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[208:209], v[160:161], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[210:211], v[162:163], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[80:83],  %[v_os_b5], s[12:15], 0 offen  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[212:213], v[164:165], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[214:215], v[166:167], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[216:217], v[168:169], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[218:219], v[170:171], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[84:87],  %[v_os_b5], s[12:15], 0 offen offset:1024  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[220:221], v[172:173], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[222:223], v[174:175], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[208:209], v[224:225], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[210:211], v[226:227], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[88:91],  %[v_os_b5], s[12:15], 0 offen offset:2048  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[212:213], v[228:229], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[214:215], v[230:231], v[92:95]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[216:217], v[232:233], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[218:219], v[234:235], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[92:95],  %[v_os_b5], s[12:15], 0 offen offset:3072  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[220:221], v[236:237], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[222:223], v[238:239], v[92:95]  \n"
-                     " s_mov_b64 exec, %[s_execflag_5] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o5], v15, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n"
-                     " s_waitcnt vmcnt(30)  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[224:225], v[176:177], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[226:227], v[178:179], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[96:99],  %[v_os_b6], s[12:15], 0 offen  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[228:229], v[180:181], v[80:83]  "
-                     "\n" _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[230:231], v[182:183], "
-                     "v[80:83]  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[232:233], v[184:185], v[80:83]  "
-                     "\n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[234:235], v[186:187], v[80:83]  \n"
-                     " buffer_load_dwordx4 acc[100:103],  %[v_os_b6], s[12:15], 0 offen "
-                     "offset:1024  \n" _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[236:237], "
-                     "v[188:189], v[80:83]  \n" _UK_MFMA_
-                     " [%[c16], %[c17], %[c18], %[c19]], acc[238:239], v[190:191], v[80:83]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[224:225], v[240:241], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[226:227], v[242:243], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[104:107],  %[v_os_b6], s[12:15], 0 offen "
-                     "offset:2048  \n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[228:229], "
-                     "v[244:245], v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[230:231], v[246:247], v[84:87]  "
-                     "\n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[232:233], v[248:249], "
-                     "v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[234:235], v[250:251], v[84:87]  \n"
-                     " buffer_load_dwordx4 acc[108:111],  %[v_os_b6], s[12:15], 0 offen "
-                     "offset:3072  \n" _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[236:237], "
-                     "v[252:253], v[84:87]  \n" _UK_MFMA_
-                     " [%[c20], %[c21], %[c22], %[c23]], acc[238:239], v[254:255], v[84:87]  \n"
-                     " s_mov_b64 exec, %[s_execflag_6] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o6], v16, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[240:241], v[176:177], v[88:91]  "
-                     "\n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[242:243], v[178:179], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[112:115],  %[v_os_b7], s[12:15], 0 offen  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[244:245], v[180:181], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[246:247], v[182:183], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[248:249], v[184:185], "
-                     "v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[250:251], v[186:187], v[88:91]  \n"
-                     " buffer_load_dwordx4 acc[116:119],  %[v_os_b7], s[12:15], 0 offen "
-                     "offset:1024  \n" _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[252:253], "
-                     "v[188:189], v[88:91]  \n" _UK_MFMA_
-                     " [%[c24], %[c25], %[c26], %[c27]], acc[254:255], v[190:191], v[88:91]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[240:241], v[240:241], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[242:243], v[242:243], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[120:123],  %[v_os_b7], s[12:15], 0 offen "
-                     "offset:2048  \n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[244:245], "
-                     "v[244:245], v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[246:247], v[246:247], v[92:95]  "
-                     "\n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[248:249], v[248:249], "
-                     "v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[250:251], v[250:251], v[92:95]  \n"
-                     " buffer_load_dwordx4 acc[124:127],  %[v_os_b7], s[12:15], 0 offen "
-                     "offset:3072  \n" _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[252:253], "
-                     "v[252:253], v[92:95]  \n" _UK_MFMA_
-                     " [%[c28], %[c29], %[c30], %[c31]], acc[254:255], v[254:255], v[92:95]  \n"
-                     " s_mov_b64 exec, %[s_execflag_7] \n" _UK_ATOMIC_ADD_
-                     " %[v_os_o7], v17, s[8:9] \n"
-                     "  s_mov_b64     exec, s[38:39]                           \n"
-                     " s_add_u32 s60, 0x00000100, s80  \n"
-                     " s_cmp_lt_u32 s60, %[s_loop_cnt]  \n"
-                     " s_cselect_b32 s56, s56, 0  \n"
-                     " s_add_u32 s12, s56, s12  \n"
-                     " s_addc_u32 s13, 0, s13  \n"
-                     " s_cmp_ge_u32 s80, 0x00000100  \n"
-                     " s_cselect_b32 s59, 0x00000100, s59  \n"
-                     " s_add_u32 s8, s59, s8  \n"
-                     " s_addc_u32 s9, 0, s9  \n"
-                     " v_mul_f32 %[c16], %[scale_0], %[c16] \n"
-                     " v_mul_f32 %[c17], %[scale_0], %[c17] \n"
-                     " v_mul_f32 %[c18], %[scale_0], %[c18] \n"
-                     " v_mul_f32 %[c19], %[scale_0], %[c19] \n"
-                     " v_mul_f32 %[c20], %[scale_1], %[c20] \n"
-                     " v_mul_f32 %[c21], %[scale_1], %[c21] \n"
-                     " v_mul_f32 %[c22], %[scale_1], %[c22] \n"
-                     " v_mul_f32 %[c23], %[scale_1], %[c23] \n"
-                     " v_mul_f32 %[c24], %[scale_0], %[c24] \n"
-                     " v_mul_f32 %[c25], %[scale_0], %[c25] \n"
-                     " v_mul_f32 %[c26], %[scale_0], %[c26] \n"
-                     " v_mul_f32 %[c27], %[scale_0], %[c27] \n"
-                     " v_mul_f32 %[c28], %[scale_1], %[c28] \n"
-                     " v_mul_f32 %[c29], %[scale_1], %[c29] \n"
-                     " v_mul_f32 %[c30], %[scale_1], %[c30] \n"
-                     " v_mul_f32 %[c31], %[scale_1], %[c31] \n" _UK_PK_CVT_(
-                         "%[c16]", "%[c17]", "%[c16]") _UK_PK_CVT_("%[c18]", "%[c19]", "%[c17]")
-                         _UK_PK_CVT_("%[c20]", "%[c21]", "%[c18]") _UK_PK_CVT_(
-                             "%[c22]", "%[c23]", "%[c19]") _UK_PK_CVT_("%[c24]", "%[c25]", "%[c20]")
-                             _UK_PK_CVT_("%[c26]", "%[c27]", "%[c21]")
-                                 _UK_PK_CVT_("%[c28]", "%[c29]", "%[c22]") _UK_PK_CVT_(
-                                     "%[c30]",
-                                     "%[c31]",
-                                     "%[c23]") " s_addk_i32 s80, 0x0080  \n"
-                                               " s_cmp_lt_i32 s80, %[s_loop_cnt]  \n"
-                                               " s_cbranch_scc0 label_0EC1  \n"
-                                               " s_branch label_0AA6  \n"
-                                               " label_0EC1: \n"
-                                               " s_waitcnt lgkmcnt(0)  \n"
-                                               " s_barrier  \n"
-                                               " ds_read_b32 v10, %[v_sfl_sld] offset:16640  \n"
-                                               " ds_read_b32 v11, %[v_sfl_sld] offset:16672  \n"
-                                               " ds_read_b32 v12, %[v_sfl_sld] offset:16704  \n"
-                                               " ds_read_b32 v13, %[v_sfl_sld] offset:16736  \n"
-                                               " ds_read_b32 v14, %[v_sfl_sld] offset:20992  \n"
-                                               " ds_read_b32 v15, %[v_sfl_sld] offset:21024  \n"
-                                               " ds_read_b32 v16, %[v_sfl_sld] offset:21056  \n"
-                                               " ds_read_b32 v17, %[v_sfl_sld] offset:21088  \n"
-                                               " s_waitcnt lgkmcnt(0)  \n"
-                                               " s_mov_b64 exec, %[s_execflag_0] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o0], v10, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_1] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o1], v11, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_2] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o2], v12, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_3] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o3], v13, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_4] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o4], v14, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_5] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o5], v15, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_6] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o6], v16, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_7] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o7], v17, s[8:9] \n"
-                                               "  s_mov_b64     exec, s[38:39]                     "
-                                               "      \n"
-                                               " s_add_u32 s8, s59, s8  \n"
-                                               " s_addc_u32 s9, 0, s9  \n"
-                                               " ds_write_b64 %[v_sfl_sst], [%[c16],%[c17]] "
-                                               "offset:25344  \n"
-                                               " ds_write_b64 %[v_sfl_sst], [%[c18],%[c19]] "
-                                               "offset:29696  \n"
-                                               " ds_write_b64 %[v_sfl_sst], [%[c20],%[c21]] "
-                                               "offset:27520  \n"
-                                               " ds_write_b64 %[v_sfl_sst], [%[c22],%[c23]] "
-                                               "offset:31872  \n"
-                                               " s_waitcnt lgkmcnt(0)  \n"
-                                               " s_barrier  \n"
-                                               " ds_read_b32 v10, %[v_sfl_sld] offset:25344  \n"
-                                               " ds_read_b32 v11, %[v_sfl_sld] offset:25376  \n"
-                                               " ds_read_b32 v12, %[v_sfl_sld] offset:25408  \n"
-                                               " ds_read_b32 v13, %[v_sfl_sld] offset:25440  \n"
-                                               " ds_read_b32 v14, %[v_sfl_sld] offset:29696  \n"
-                                               " ds_read_b32 v15, %[v_sfl_sld] offset:29728  \n"
-                                               " ds_read_b32 v16, %[v_sfl_sld] offset:29760  \n"
-                                               " ds_read_b32 v17, %[v_sfl_sld] offset:29792  \n"
-                                               " s_waitcnt lgkmcnt(0)  \n"
-                                               " s_mov_b64 exec, %[s_execflag_0] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o0], v10, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_1] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o1], v11, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_2] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o2], v12, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_3] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o3], v13, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_4] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o4], v14, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_5] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o5], v15, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_6] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o6], v16, s[8:9] \n"
-                                               " s_mov_b64 exec, %[s_execflag_7] \n" _UK_ATOMIC_ADD_
-                                               " %[v_os_o7], v17, s[8:9] \n"
-                                               "  s_mov_b64     exec, s[38:39]  \n"
+    " v_mul_f32 %[c15], %[scale_1], %[c15] \n"
+    _UK_PK_CVT_("%[c0]","%[c1]","%[c0]")
+    _UK_PK_CVT_("%[c2]","%[c3]","%[c1]")
+    _UK_PK_CVT_("%[c4]","%[c5]","%[c2]")
+    _UK_PK_CVT_("%[c6]","%[c7]","%[c3]")
+    _UK_PK_CVT_("%[c8]","%[c9]","%[c4]")
+    _UK_PK_CVT_("%[c10]","%[c11]","%[c5]")
+    _UK_PK_CVT_("%[c12]", "%[c13]", "%[c6]")
+    _UK_PK_CVT_("%[c14]","%[c15]","%[c7]")
+    " s_addk_i32 s80, 0x0080  \n"
+    " s_cmp_lt_i32 s80, %[s_loop_cnt]  \n"
+    " s_cbranch_scc0 loop_atomic_%=  \n"
+    " s_waitcnt vmcnt(30) & lgkmcnt(0)  \n"
+    " s_barrier  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[128:129], v[128:129], 0  \n"
+    " ds_read_b32 v10, %[v_sfl_sld] offset:25344  \n"
+    " ds_read_b32 v11, %[v_sfl_sld] offset:25376  \n"
+    " ds_write_b64 %[v_sfl_sst], v[64:65] offset:16640  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[130:131], v[130:131], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[0:3],  %[v_os_b0], s[12:15], 0 offen  \n"
+     " ds_write_b64 %[v_sfl_sst], v[66:67] offset:20992  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[132:133], v[132:133], v[80:83]  \n"
+     " ds_read_b32 v12, %[v_sfl_sld] offset:25408  \n"
+     " ds_read_b32 v13, %[v_sfl_sld] offset:25440  \n"
+     " ds_write_b64 %[v_sfl_sst], v[68:69] offset:18816  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[134:135], v[134:135], v[80:83]  \n"
+     " ds_write_b64 %[v_sfl_sst], v[70:71] offset:23168  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[136:137], v[136:137], v[80:83]  \n"
+     " ds_read_b32 v14, %[v_sfl_sld] offset:29696  \n"
+     " ds_read_b32 v15, %[v_sfl_sld] offset:29728  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[138:139], v[138:139], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[4:7],  %[v_os_b0], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[140:141], v[140:141], v[80:83]  \n"
+     " ds_read_b32 v16, %[v_sfl_sld] offset:29760  \n"
+     " ds_read_b32 v17, %[v_sfl_sld] offset:29792  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[142:143], v[142:143], v[80:83]  \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[128:129], v[192:193], 0  \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[130:131], v[194:195], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[8:11],  %[v_os_b0], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[132:133], v[196:197], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[134:135], v[198:199], v[84:87]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[136:137], v[200:201], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[138:139], v[202:203], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[12:15],  %[v_os_b0], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[140:141], v[204:205], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[142:143], v[206:207], v[84:87]  \n"
+     " s_waitcnt lgkmcnt(0) \n"
+     " s_mov_b64 exec, %[s_execflag_0] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o0], v10, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[144:145], v[128:129], 0  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[146:147], v[130:131], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[16:19],  %[v_os_b1], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[148:149], v[132:133], v[88:91]  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[150:151], v[134:135], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[152:153], v[136:137], v[88:91]  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[154:155], v[138:139], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[20:23],  %[v_os_b1], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[156:157], v[140:141], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[158:159], v[142:143], v[88:91]  \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[144:145], v[192:193], 0  \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[146:147], v[194:195], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[24:27],  %[v_os_b1], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[148:149], v[196:197], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[150:151], v[198:199], v[92:95]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[152:153], v[200:201], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[154:155], v[202:203], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[28:31],  %[v_os_b1], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[156:157], v[204:205], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[158:159], v[206:207], v[92:95]  \n"
+     " s_mov_b64 exec, %[s_execflag_1] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o1], v11, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+     " s_waitcnt vmcnt(30) \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[160:161], v[144:145], v[80:83]  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[162:163], v[146:147], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[32:35],  %[v_os_b2], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[164:165], v[148:149], v[80:83]  \n"
+    _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[166:167], v[150:151], v[80:83]      \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[168:169], v[152:153], v[80:83]  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[170:171], v[154:155], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[36:39],  %[v_os_b2], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[172:173], v[156:157], v[80:83]      \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[174:175], v[158:159], v[80:83]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[160:161], v[208:209], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[162:163], v[210:211], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[40:43],  %[v_os_b2], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[164:165], v[212:213], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[166:167], v[214:215], v[84:87]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[168:169], v[216:217], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[170:171], v[218:219], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[44:47],  %[v_os_b2], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[172:173], v[220:221], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[174:175], v[222:223], v[84:87]  \n"
+     " s_mov_b64 exec, %[s_execflag_2] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o2], v12, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[176:177], v[144:145], v[88:91]  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[178:179], v[146:147], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[48:51],  %[v_os_b3], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[180:181], v[148:149], v[88:91]  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[182:183], v[150:151], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[184:185], v[152:153], v[88:91]  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[186:187], v[154:155], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[52:55],  %[v_os_b3], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[188:189], v[156:157], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[190:191], v[158:159], v[88:91]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[176:177], v[208:209], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[178:179], v[210:211], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[56:59],  %[v_os_b3], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[180:181], v[212:213], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[182:183], v[214:215], v[92:95]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[184:185], v[216:217], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[186:187], v[218:219], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[60:63],  %[v_os_b3], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[188:189], v[220:221], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[190:191], v[222:223], v[92:95]  \n"
+     " s_mov_b64 exec, %[s_execflag_3] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o3], v13, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+     " s_waitcnt vmcnt(30) \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[192:193], v[160:161], v[80:83]  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[194:195], v[162:163], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[64:67],  %[v_os_b4], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[196:197], v[164:165], v[80:83]  \n"
+    _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[198:199], v[166:167], v[80:83]      \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[200:201], v[168:169], v[80:83]  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[202:203], v[170:171], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[68:71],  %[v_os_b4], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[204:205], v[172:173], v[80:83]      \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[206:207], v[174:175], v[80:83]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[192:193], v[224:225], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[194:195], v[226:227], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[72:75],  %[v_os_b4], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[196:197], v[228:229], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[198:199], v[230:231], v[84:87]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[200:201], v[232:233], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[202:203], v[234:235], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[76:79],  %[v_os_b4], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[204:205], v[236:237], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[206:207], v[238:239], v[84:87]  \n"
+     " s_mov_b64 exec, %[s_execflag_4] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o4], v14, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[208:209], v[160:161], v[88:91]  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[210:211], v[162:163], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[80:83],  %[v_os_b5], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[212:213], v[164:165], v[88:91]  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[214:215], v[166:167], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[216:217], v[168:169], v[88:91]  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[218:219], v[170:171], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[84:87],  %[v_os_b5], s[12:15], 0 offen offset:1024  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[220:221], v[172:173], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[222:223], v[174:175], v[88:91]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[208:209], v[224:225], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[210:211], v[226:227], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[88:91],  %[v_os_b5], s[12:15], 0 offen offset:2048  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[212:213], v[228:229], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[214:215], v[230:231], v[92:95]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[216:217], v[232:233], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[218:219], v[234:235], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[92:95],  %[v_os_b5], s[12:15], 0 offen offset:3072  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[220:221], v[236:237], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[222:223], v[238:239], v[92:95]  \n"
+     " s_mov_b64 exec, %[s_execflag_5] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o5], v15, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+     " s_waitcnt vmcnt(30)  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[224:225], v[176:177], v[80:83]  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[226:227], v[178:179], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[96:99],  %[v_os_b6], s[12:15], 0 offen  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[228:229], v[180:181], v[80:83]  \n"
+    _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[230:231], v[182:183], v[80:83]      \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[232:233], v[184:185], v[80:83]  \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[234:235], v[186:187], v[80:83]  \n"
+     " buffer_load_dwordx4 acc[100:103],  %[v_os_b6], s[12:15], 0 offen "
+     "offset:1024  \n"
+    _UK_MFMA_ " [%[c16], %[c17], %[c18], %[c19]], acc[236:237], v[188:189], v[80:83]      \n"
+    _UK_MFMA_    " [%[c16], %[c17], %[c18], %[c19]], acc[238:239], v[190:191], v[80:83]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[224:225], v[240:241], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[226:227], v[242:243], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[104:107],  %[v_os_b6], s[12:15], 0 offen "
+     "offset:2048  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[228:229], v[244:245], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[230:231], v[246:247], v[84:87]  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[232:233], v[248:249], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[234:235], v[250:251], v[84:87]  \n"
+     " buffer_load_dwordx4 acc[108:111],  %[v_os_b6], s[12:15], 0 offen "
+     "offset:3072  \n"
+    _UK_MFMA_ " [%[c20], %[c21], %[c22], %[c23]], acc[236:237], v[252:253], v[84:87]      \n"
+    _UK_MFMA_    " [%[c20], %[c21], %[c22], %[c23]], acc[238:239], v[254:255], v[84:87]  \n"
+     " s_mov_b64 exec, %[s_execflag_6] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o6], v16, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[240:241], v[176:177], v[88:91]  \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[242:243], v[178:179], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[112:115],  %[v_os_b7], s[12:15], 0 offen  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[244:245], v[180:181], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[246:247], v[182:183], v[88:91]  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[248:249], v[184:185], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[250:251], v[186:187], v[88:91]  \n"
+     " buffer_load_dwordx4 acc[116:119],  %[v_os_b7], s[12:15], 0 offen "
+     "offset:1024  \n"
+    _UK_MFMA_ " [%[c24], %[c25], %[c26], %[c27]], acc[252:253], v[188:189], v[88:91]      \n"
+    _UK_MFMA_    " [%[c24], %[c25], %[c26], %[c27]], acc[254:255], v[190:191], v[88:91]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[240:241], v[240:241], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[242:243], v[242:243], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[120:123],  %[v_os_b7], s[12:15], 0 offen "
+     "offset:2048  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[244:245], v[244:245], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[246:247], v[246:247], v[92:95]  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[248:249], v[248:249], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[250:251], v[250:251], v[92:95]  \n"
+     " buffer_load_dwordx4 acc[124:127],  %[v_os_b7], s[12:15], 0 offen "
+     "offset:3072  \n"
+    _UK_MFMA_ " [%[c28], %[c29], %[c30], %[c31]], acc[252:253], v[252:253], v[92:95]      \n"
+    _UK_MFMA_    " [%[c28], %[c29], %[c30], %[c31]], acc[254:255], v[254:255], v[92:95]  \n"
+     " s_mov_b64 exec, %[s_execflag_7] \n"
+    _UK_ATOMIC_ADD_     " %[v_os_o7], v17, s[8:9] \n"
+     "  s_mov_b64     exec, s[38:39]                           \n"
+     " s_add_u32 s60, 0x00000100, s80  \n"
+     " s_cmp_lt_u32 s60, %[s_loop_cnt]  \n"
+     " s_cselect_b32 s56, s56, 0  \n"
+     " s_add_u32 s12, s56, s12  \n"
+     " s_addc_u32 s13, 0, s13  \n"
+     " s_cmp_ge_u32 s80, 0x00000100  \n"
+     " s_cselect_b32 s59, 0x00000100, s59  \n"
+     " s_add_u32 s8, s59, s8  \n"
+     " s_addc_u32 s9, 0, s9  \n"
+     " v_mul_f32 %[c16], %[scale_0], %[c16] \n"
+     " v_mul_f32 %[c17], %[scale_0], %[c17] \n"
+     " v_mul_f32 %[c18], %[scale_0], %[c18] \n"
+     " v_mul_f32 %[c19], %[scale_0], %[c19] \n"
+     " v_mul_f32 %[c20], %[scale_1], %[c20] \n"
+     " v_mul_f32 %[c21], %[scale_1], %[c21] \n"
+     " v_mul_f32 %[c22], %[scale_1], %[c22] \n"
+     " v_mul_f32 %[c23], %[scale_1], %[c23] \n"
+     " v_mul_f32 %[c24], %[scale_0], %[c24] \n"
+     " v_mul_f32 %[c25], %[scale_0], %[c25] \n"
+     " v_mul_f32 %[c26], %[scale_0], %[c26] \n"
+     " v_mul_f32 %[c27], %[scale_0], %[c27] \n"
+     " v_mul_f32 %[c28], %[scale_1], %[c28] \n"
+     " v_mul_f32 %[c29], %[scale_1], %[c29] \n"
+     " v_mul_f32 %[c30], %[scale_1], %[c30] \n"
+     " v_mul_f32 %[c31], %[scale_1], %[c31] \n"
+    _UK_PK_CVT_("%[c16]", "%[c17]", "%[c16]")
+    _UK_PK_CVT_("%[c18]", "%[c19]", "%[c17]")
+    _UK_PK_CVT_("%[c20]", "%[c21]", "%[c18]")
+    _UK_PK_CVT_("%[c22]", "%[c23]", "%[c19]")
+    _UK_PK_CVT_("%[c24]", "%[c25]", "%[c20]")
+    _UK_PK_CVT_("%[c26]", "%[c27]", "%[c21]")
+    _UK_PK_CVT_("%[c28]", "%[c29]", "%[c22]")
+    _UK_PK_CVT_("%[c30]","%[c31]","%[c23]")
+    " s_addk_i32 s80, 0x0080  \n"
+    " s_cmp_lt_i32 s80, %[s_loop_cnt]  \n"
+    " s_cbranch_scc0 loop_atomic_%=  \n"
+    " s_branch coreloop_top_%=  \n"
+    " loop_atomic_%=: \n"
+    " s_waitcnt lgkmcnt(0)  \n"
+    " s_barrier  \n"
+    " ds_read_b32 v10, %[v_sfl_sld] offset:16640  \n"
+    " ds_read_b32 v11, %[v_sfl_sld] offset:16672  \n"
+    " ds_read_b32 v12, %[v_sfl_sld] offset:16704  \n"
+    " ds_read_b32 v13, %[v_sfl_sld] offset:16736  \n"
+    " ds_read_b32 v14, %[v_sfl_sld] offset:20992  \n"
+    " ds_read_b32 v15, %[v_sfl_sld] offset:21024  \n"
+    " ds_read_b32 v16, %[v_sfl_sld] offset:21056  \n"
+    " ds_read_b32 v17, %[v_sfl_sld] offset:21088  \n"
+    " s_waitcnt lgkmcnt(0)  \n"
+    " s_mov_b64 exec, %[s_execflag_0] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o0], v10, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_1] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o1], v11, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_2] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o2], v12, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_3] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o3], v13, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_4] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o4], v14, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_5] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o5], v15, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_6] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o6], v16, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_7] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o7], v17, s[8:9] \n"
+    "  s_mov_b64     exec, s[38:39]              \n"
+    " s_add_u32 s8, s59, s8  \n"
+    " s_addc_u32 s9, 0, s9  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c16],%[c17]] "
+    "offset:25344  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c18],%[c19]] "
+    "offset:29696  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c20],%[c21]] "
+    "offset:27520  \n"
+    " ds_write_b64 %[v_sfl_sst], [%[c22],%[c23]] "
+    "offset:31872  \n"
+    " s_waitcnt lgkmcnt(0)  \n"
+    " s_barrier  \n"
+    " ds_read_b32 v10, %[v_sfl_sld] offset:25344  \n"
+    " ds_read_b32 v11, %[v_sfl_sld] offset:25376  \n"
+    " ds_read_b32 v12, %[v_sfl_sld] offset:25408  \n"
+    " ds_read_b32 v13, %[v_sfl_sld] offset:25440  \n"
+    " ds_read_b32 v14, %[v_sfl_sld] offset:29696  \n"
+    " ds_read_b32 v15, %[v_sfl_sld] offset:29728  \n"
+    " ds_read_b32 v16, %[v_sfl_sld] offset:29760  \n"
+    " ds_read_b32 v17, %[v_sfl_sld] offset:29792  \n"
+    " s_waitcnt lgkmcnt(0)  \n"
+    " s_mov_b64 exec, %[s_execflag_0] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o0], v10, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_1] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o1], v11, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_2] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o2], v12, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_3] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o3], v13, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_4] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o4], v14, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_5] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o5], v15, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_6] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o6], v16, s[8:9] \n"
+    " s_mov_b64 exec, %[s_execflag_7] \n"
+    _UK_ATOMIC_ADD_    " %[v_os_o7], v17, s[8:9] \n"
+    "  s_mov_b64     exec, s[38:39]  \n"
 
 #undef _UK_MFMA_
 #undef _UK_PK_CVT_
 #undef _UK_ATOMIC_ADD_
+#undef CK_TILE_FLATMM_UK_MFMA
+// clang-format on