fixed

host/include/ck/host/device_gemm_multiple_d/operation.hpp

@@ -20,7 +20,7 @@ struct Operation_Xdl_CShuffle
     static std::vector<Operation_Xdl_CShuffle> CreateOperations(const Problem& prob);
     TensorDesc A{};
     TensorDesc B{};
-    DataType acc               = DataType::Half;
+    DataType acc               = DataType::Float;
     DataType cs_type           = DataType::Half;
     std::vector<TensorDesc> Ds = {};
     TensorDesc E{};

host/include/ck/host/device_gemm_multiple_d/problem.hpp

@@ -27,7 +27,7 @@ struct Problem
     std::vector<DataType> DsDataType = {};
     std::string AElementOp           = PassThrough;
     std::string BElementOp           = PassThrough;
-    std::string CDEElementOp         = "ck::Tuple<>";
+    std::string CDEElementOp         = PassThrough;
 
     std::string GetIncludeHeader() const;
 

host/test/gemm_multiple_d.cpp

@@ -47,9 +47,9 @@ extern "C" __global__ void f(const ck::half_t* a, const ck::half_t* b, ck::half_
 TEST_CASE(test_problem_kernel)
 {
     ck::host::device_gemm_multiple_d::Problem prob;
-    prob.M = 256;
-    prob.N = 256;
-    prob.K = 256;
+    prob.M = 1024;
+    prob.N = 1024;
+    prob.K = 1024;
     for(auto solution : prob.GetSolutions("gfx90a"))
     {
         auto src  = ck::host::InterpolateString(gemm_compile_check,

include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp

@@ -309,10 +309,10 @@ struct GridwiseGemmMultipleD_xdl_cshuffle
         }
 
         // check block-to-E-tile
-        if(!block_2_etile_map.CheckValidity(e_grid_desc_m_n))
-        {
-            return false;
-        }
+        //if(!block_2_etile_map.CheckValidity(e_grid_desc_m_n))
+        //{
+            //return false;
+        //}
 
         // TODO: also check validity of all components (blockwise-copy, threadwise-copy, etc)
         // check tensor size: cannot be larger than 2GB each