[CK TILE STREAMK] Add 'dp_persistent' and 'reduction_strategy' in output of CK TILE STREAMK

Author: CongMa13

tile_engine/ops/gemm_streamk/gemm_streamk_benchmark.hpp

@@ -102,6 +102,8 @@ struct PerformanceResult
 struct KernelInstance
 {
     std::string name_;
+    std::string dp_persistent_;
+    std::string reduction_strategy_;
     GemmProblem problem_;
     PerformanceResult perf_result_;
 
@@ -114,6 +116,8 @@ struct KernelInstance
     {
         os << "{\n"
            << " \"name\": \"" << obj.name_ << "\",\n"
+           << " \"dp_persistent\": \"" << obj.dp_persistent_ << "\",\n"
+           << " \"reduction_strategy\": \"" << obj.reduction_strategy_ << "\",\n"
            << " \"problem\": " << obj.problem_ << ",\n"
            << " \"perf_result\": " << obj.perf_result_ << "\n"
            << "}";

tile_engine/ops/gemm_streamk/gemm_streamk_benchmark_single.cpp

@@ -160,7 +160,10 @@ int main(int argc, char* argv[])
     {
         auto [result, parser] = create_args(argc, argv);
         if(!result)
+        {
+            parser.print();
             return EXIT_FAILURE;
+        }
 
         benchmark_gemm_single(parser);
         return EXIT_SUCCESS;

tile_engine/ops/gemm_streamk/gemm_streamk_profiler.hpp

@@ -164,8 +164,15 @@ class GemmProfiler
                         const std::tuple<std::string, float>& kernel_run_result)
     {
         auto [name, avg_time] = kernel_run_result;
+        auto dp_persistent =
+            SelectedKernel::UsePersistentKernel ? "PersistentKernel" : "NonPersistentKernel";
+        auto reduction_strategy =
+            SelectedKernel::reduction_strategy == ck_tile::StreamKReductionStrategy::Atomic
+                ? "Atomic"
+                : "Reduction";
 
-        KernelInstance kernel_instance{name, gemm_problem, {-1.0f, -1.0f, -1.0f}};
+        KernelInstance kernel_instance{
+            name, dp_persistent, reduction_strategy, gemm_problem, {-1.0f, -1.0f, -1.0f}};
 
         // compute performance metric
         std::size_t flop     = std::size_t(2) * gemm_problem.m_ * gemm_problem.n_ * gemm_problem.k_;
@@ -244,21 +251,24 @@ class GemmProfiler
                     file << "rocm_version,device_name,"
                          << "split_k,m,n,k,stride_a,stride_b,stride_c,"
                          << "dtype_a,dtype_b,dtype_acc,dtype_c," << "layout_a,layout_b,layout_c,"
-                         << "structured_sparsity," << "name,"
-                         << "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n";
+                         << "structured_sparsity," << "dp_persistent," << "reduction_strategy,"
+                         << "name," << "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n";
                 }
 
-                const auto& problem = kernel_instance.problem_;
-                const auto& name    = kernel_instance.name_;
-                const auto& perf    = kernel_instance.perf_result_;
+                const auto& problem            = kernel_instance.problem_;
+                const auto& name               = kernel_instance.name_;
+                const auto& dp_persistent      = kernel_instance.dp_persistent_;
+                const auto& reduction_strategy = kernel_instance.reduction_strategy_;
+                const auto& perf               = kernel_instance.perf_result_;
 
                 file << get_rocm_version() << "," << ck_tile::get_device_name() << ","
                      << problem.split_k_ << "," << problem.m_ << "," << problem.n_ << ","
                      << problem.k_ << "," << problem.stride_a_ << "," << problem.stride_b_ << ","
                      << problem.stride_c_ << "," << problem.dtype_a_ << "," << problem.dtype_b_
                      << "," << problem.dtype_acc_ << "," << problem.dtype_c_ << ","
                      << problem.layout_a_ << "," << problem.layout_b_ << "," << problem.layout_c_
-                     << "," << problem.structured_sparsity_ << "," << name << "," << std::fixed
+                     << "," << problem.structured_sparsity_ << "," << dp_persistent << ","
+                     << reduction_strategy << "," << name << "," << std::fixed
                      << std::setprecision(4) << perf.latency_ << "," << std::fixed
                      << std::setprecision(4) << perf.tflops_ << "," << std::fixed
                      << std::setprecision(4) << perf.bandwidth_ << "," << get_metric_name(metric)