Addressed some review comments

Signed-off-by: Chris Tsiaousis <chris.tsiaousis@streamhpc.com>

Author: chris-tsiaousis-hpc

projects/composablekernel/include/ck_tile/core.hpp

@@ -28,12 +28,12 @@
 #include "ck_tile/core/arch/mma/sparse/mfma/selector.hpp"
 #include "ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse.hpp"
+#include "ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse_selector.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse_traits.hpp"
 #include "ck_tile/core/arch/mma/sparse/sparse_transforms.hpp"
 #include "ck_tile/core/arch/mma/sparse/wmma/selector.hpp"
 #include "ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp"
-#include "ck_tile/core/arch/mma/sparse_mma.hpp"
 #include "ck_tile/core/arch/mma/utility/tile_distribution_encoding_register_mapper.hpp"
 #include "ck_tile/core/arch/mma/wmma/wmma.hpp"
 #include "ck_tile/core/arch/mma/wmma/wmma_gfx11.hpp"

projects/composablekernel/include/ck_tile/core/arch/mma/mma_pipeline.hpp

@@ -8,21 +8,14 @@
 #include "mma_selector.hpp"
 #include "mma_traits.hpp"
 #include "mma_transforms.hpp"
-#include <netdb.h>
-#include <optional>
-#include <type_traits>
 
 namespace ck_tile::core::arch::mma {
 
 enum struct MmaPipelineOptionFlag
 {
-    NONE                   = 0x0,
-    C_TRANSPOSE            = 0x1,
-    SWIZZLE_A              = 0x2,
-    SWIZZLE_B              = 0x4,
-    DOUBLE_ATTR_NUM_ACCESS = 0x8,
-    QUAD_ATTR_NUM_ACCESS   = 0x10,
-    COMPRESS_A             = 0x20,
+    NONE        = 0x0,
+    C_TRANSPOSE = 0x1,
+    COMPRESS_A  = 0x2,
 };
 
 struct MmaPipelineOptionFlags
@@ -81,12 +74,6 @@ struct MmaPipelineBase
     static constexpr auto Flags = MmaPipelineOptionFlags(Flags_);
     // TODO: Implement those cases
     static_assert(!(Flags & MmaPipelineOptionFlag::C_TRANSPOSE), "Flag not yet implemented");
-    static_assert(!(Flags & MmaPipelineOptionFlag::SWIZZLE_A), "Flag not yet implemented");
-    static_assert(!(Flags & MmaPipelineOptionFlag::SWIZZLE_B), "Flag not yet implemented");
-    static_assert(!(Flags & MmaPipelineOptionFlag::DOUBLE_ATTR_NUM_ACCESS),
-                  "Flag not yet implemented");
-    static_assert(!(Flags & MmaPipelineOptionFlag::QUAD_ATTR_NUM_ACCESS),
-                  "Flag not yet implemented");
 
     private:
     template <typename DstT, typename SrcT>
@@ -106,7 +93,7 @@ struct MmaPipelineBase
 
     protected:
     template <MmaPipelineOptionFlag Flag>
-    CK_TILE_DEVICE static bool hasFlag()
+    constexpr CK_TILE_DEVICE static bool hasFlag()
     {
         return Flags & Flag;
     }

projects/composablekernel/include/ck_tile/core/arch/mma/mma_wavewise.hpp

@@ -86,7 +86,7 @@ struct WaveWiseMma : public MmaPipelineBase<static_cast<int>(MmaPipelineOptionFl
     constexpr static uint32_t FragM = MmaOp::kM;
     constexpr static uint32_t FragN = MmaOp::kN;
     constexpr static uint32_t FragK = MmaOp::kK;
-    
+
     using BlockWiseMmaOp       = MmaOp;
     using BlockWiseMmaOpTraits = MmaOpTraits<BlockWiseMmaOp>;
 
@@ -174,7 +174,7 @@ struct WaveWiseMma : public MmaPipelineBase<static_cast<int>(MmaPipelineOptionFl
                 }
             }
         }
-        else
+        else if constexpr(AccumPolicy == MmaAccumPolicy::COL_MAJOR)
         {
             // "Col-major" accumulation over the M-dimension blocks first.
             // Pseudo code here, but we would basically iterate over the blocks in col-major order
@@ -190,6 +190,10 @@ struct WaveWiseMma : public MmaPipelineBase<static_cast<int>(MmaPipelineOptionFl
                 }
             }
         }
+        else
+        {
+            static_assert(false);
+        }
     }
 };
 

…ude/ck_tile/core/arch/mma/sparse_mma.hpp …/arch/mma/sparse/sparse_mma_pipeline.hppprojects/composablekernel/include/ck_tile/core/arch/mma/sparse_mma.hpp renamed to projects/composablekernel/include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp projects/composablekernel/include/ck_tile/core/arch/mma/sparse_mma.hpp renamed to projects/composablekernel/include/ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp

@@ -41,7 +41,7 @@ static CK_TILE_DEVICE int32_t compress_a_impl(AVec& a_vec)
         int32_t non_zero_pos       = 0;
 
         static_for<0, 3, 1>{}([&](auto j) {
-            if(a_vec[i * 4 + j] != 0.0f)
+            if(static_cast<float>(a_vec[i * 4 + j]) != 0.0f)
             {
                 nonzero_elems[non_zero_pos] = a_vec[i * 4 + j];
                 // clear the two‑bit field for this output and insert j
@@ -68,17 +68,16 @@ struct SparseCompressTransform
     template <typename VecType>
     CK_TILE_DEVICE static decltype(auto) exec(VecType&& v, int32_t& idx)
     {
-        using VecTraits                         = vector_traits<std::decay_t<VecType>>;
+        using VecTraits                         = vector_traits<remove_cvref_t<VecType>>;
         using ScalarT                           = typename VecTraits::scalar_type;
         static constexpr auto VecN              = VecTraits::vector_size;
         static constexpr index_t CompressedSize = VecN / CompressionRatio;
         using VecCompressed                     = ext_vector_t<ScalarT, CompressedSize>;
 
         idx = detail::compress_a_impl<ScalarT, CompressedSize>(v);
 
-        VecCompressed result;
-        __builtin_memcpy(&result, &v, sizeof(VecCompressed));
-        return result;
+        // TODO c++20: Use bit_cast
+        return *std::launder(reinterpret_cast<VecCompressed*>(&v));
     }
 };
 

projects/composablekernel/include/ck_tile/core/arch/mma/sparse/sparse_transforms.hpp

@@ -8,7 +8,7 @@
 #include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/arch/mma/mma_op_family.hpp"
 #include "ck_tile/core/arch/mma/mma_selector.hpp"
-#include "ck_tile/core/arch/mma/sparse_mma.hpp"
+#include "ck_tile/core/arch/mma/sparse/sparse_mma_pipeline.hpp"
 #include <hip/hip_runtime.h>
 #include "ck_tile/core/numeric/integer.hpp"
 #include "ck_tile/host/hip_check_error.hpp"