review changes: code cleanup

Author: arai713

projects/composablekernel/test/ck_tile/CMakeLists.txt

@@ -66,6 +66,8 @@ add_subdirectory(core)
 add_subdirectory(epilogue)
 add_subdirectory(atomic_add_op)
 add_subdirectory(fmha)
+# TODO: The Universal GEMM tile engine test will be either removed
+# or moved to the appropriate location in future work.
 # add_subdirectory(gemm_tile_engine)
 add_subdirectory(pooling)
 add_subdirectory(grouped_conv)

projects/composablekernel/tile_engine/ops/common/utils.hpp

@@ -2,13 +2,15 @@
 // SPDX-License-Identifier: MIT
 
 #pragma once
+#include <hip/hip_version.h>
 #include <iostream>
 #include <functional>
 #include <tuple>
 #include <exception>
 #include <sstream>
 #include <vector>
 #include <string>
+#include <cstdlib>
 
 #include "ck_tile/core.hpp"
 #include "ck_tile/host.hpp"
@@ -54,17 +56,6 @@ struct PerformanceResult
         default: throw std::invalid_argument("Unsupported metric type");
         }
     }
-
-    friend std::ostream& operator<<(std::ostream& os, const PerformanceResult& result)
-    {
-        os << "{\n"
-           << "   \"latency(ms)\": " << std::fixed << std::setprecision(2) << result.latency_
-           << ",\n"
-           << "   \"tflops(TFlops)\": " << result.tflops_ << ",\n"
-           << "   \"bandwidth(GB/s)\": " << result.bandwidth_ << "\n"
-           << "}";
-        return os;
-    }
 };
 
 template <typename Problem>
@@ -78,42 +69,46 @@ struct KernelInstance
     {
         return PerformanceResult::compare(a.perf_result_, b.perf_result_, m);
     }
-
-    friend std::ostream& operator<<(std::ostream& os, const KernelInstance& obj)
-    {
-        os << "{\n"
-           << " \"name\": \"" << obj.name_ << "\",\n"
-           << " \"problem\": " << obj.problem_ << ",\n"
-           << " \"perf_result\": " << obj.perf_result_ << "\n"
-           << "}";
-        return os;
-    }
 };
 
-struct Setting
+template <typename Problem>
+std::ostream& operator<<(std::ostream& os, const KernelInstance<Problem>& obj)
 {
-    int n_warmup_;
-    int n_repeat_;
-    bool is_gpu_timer_;
-    int verify_;
-    int init_method_;
-    bool log_;
-    std::string csv_filename_;
-    bool flush_cache_;
-    int rotating_count_;
-    bool json_output_;
+    os << "{\n"
+       << " \"name\": \"" << obj.name_ << "\",\n"
+       << " \"problem\": " << obj.problem_ << ",\n"
+       << " \"perf_result\": " << obj.perf_result_ << "\n"
+       << "}";
+    return os;
+}
+
+std::ostream& operator<<(std::ostream& os, const PerformanceResult& result)
+{
+    os << "{\n"
+       << "   \"latency(ms)\": " << std::fixed << std::setprecision(2) << result.latency_ << ",\n"
+       << "   \"tflops(TFlops)\": " << result.tflops_ << ",\n"
+       << "   \"bandwidth(GB/s)\": " << result.bandwidth_ << "\n"
+       << "}";
+    return os;
+}
+
+struct Settings
+{
+    int n_warmup;
+    int n_repeat;
+    bool is_gpu_timer;
+    int verify;
+    int init_method;
+    bool log;
+    std::string csv_filename;
+    bool flush_cache;
+    int rotating_count;
+    bool json_output;
 };
 
 inline std::string get_rocm_version()
 {
-    std::ifstream version_file("/opt/rocm/.info/version");
-    if(version_file.is_open())
-    {
-        std::string version;
-        std::getline(version_file, version);
-        return version;
-    }
-    return "Unknown";
+    return std::to_string(HIP_VERSION_MAJOR) + "." + std::to_string(HIP_VERSION_MINOR);
 }
 
 template <typename ADataType, typename BDataType, typename AccDataType, typename CDataType>

projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp

@@ -14,7 +14,7 @@
 #include "gemm/gemm_benchmark.hpp"
 
 #pragma clang diagnostic push
-#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-seggestions"
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
 
 // Data types and Layouts are defined by the generated kernel headers
 // No hardcoded type definitions here to avoid conflicts

projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py

@@ -1,3 +1,4 @@
+#!/usr/bin/env python3
 # Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
 
@@ -9,7 +10,7 @@
 
 
 def _import_gemm_benchmark():
-    """Import validation utilities from commons directory."""
+    """Import gemm benchmark from parent directory."""
     current_dir = os.path.dirname(os.path.abspath(__file__))
     parent_dir = os.path.dirname(current_dir)
 

projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark_single.cpp

@@ -58,17 +58,17 @@ void benchmark_single(const ck_tile::ArgParser& arg_parser)
                                            layout_d0,
                                            layout_d1};
 
-    // Create Setting struct
-    Setting setting{arg_parser.get_int("warmup"),
-                    arg_parser.get_int("repeat"),
-                    arg_parser.get_bool("timer"),
-                    arg_parser.get_int("verify"),
-                    arg_parser.get_int("init"),
-                    arg_parser.get_bool("log"),
-                    arg_parser.get_str("csv_filename"),
-                    arg_parser.get_bool("flush_cache"),
-                    arg_parser.get_int("rotating_count"),
-                    arg_parser.get_bool("json_output")};
+    // Create Settings struct
+    Settings setting{arg_parser.get_int("warmup"),
+                     arg_parser.get_int("repeat"),
+                     arg_parser.get_bool("timer"),
+                     arg_parser.get_int("verify"),
+                     arg_parser.get_int("init"),
+                     arg_parser.get_bool("log"),
+                     arg_parser.get_str("csv_filename"),
+                     arg_parser.get_bool("flush_cache"),
+                     arg_parser.get_int("rotating_count"),
+                     arg_parser.get_bool("json_output")};
 
     // Get the profiler instance
     auto& profiler = GemmMultiDProfiler::instance(setting);

projects/composablekernel/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_profiler.hpp

@@ -27,7 +27,7 @@ class GemmMultiDProfiler : public GemmProfiler<GemmMultiDProfiler,
                                   ck_tile::GemmMultiDHostArgs<DsDataType::size()>>;
     using BaseGemm::benchmark;
 
-    GemmMultiDProfiler(Setting setting)
+    GemmMultiDProfiler(Settings setting)
         : GemmProfiler<GemmMultiDProfiler,
                        GemmMultiDProblem,
                        ck_tile::GemmMultiDHostArgs<DsDataType::size()>>(setting)
@@ -141,18 +141,23 @@ class GemmMultiDProfiler : public GemmProfiler<GemmMultiDProfiler,
                                             gemm_multi_d_problem.stride_c_,
                                             is_row_major(layout_c)));
 
-        if(setting_.verify_)
+        if(setting_.verify)
         {
             gemm_multi_d_host_reference(
-                setting_.verify_, a_m_k, b_k_n, d0_m_n, d1_m_n, c_m_n_host_result);
+                setting_.verify, a_m_k, b_k_n, d0_m_n, d1_m_n, c_m_n_host_result);
         }
 
         for(auto& callable : callables)
         {
-            auto kernel_run_result =
-                callable(gemm_multi_d_args,
-                         ck_tile::stream_config{
-                             nullptr, true, setting_.log_, setting_.n_warmup_, setting_.n_repeat_});
+            auto kernel_run_result = callable(gemm_multi_d_args,
+                                              ck_tile::stream_config{nullptr,
+                                                                     true,
+                                                                     setting_.log,
+                                                                     setting_.n_warmup,
+                                                                     setting_.n_repeat,
+                                                                     setting_.is_gpu_timer,
+                                                                     setting_.flush_cache,
+                                                                     setting_.rotating_count});
             process_result(gemm_multi_d_problem,
                            c_m_n_dev_buf,
                            c_m_n_host_result,

projects/composablekernel/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py

@@ -1,3 +1,4 @@
+#!/usr/bin/env python3
 # Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
 
@@ -9,7 +10,7 @@
 
 
 def _import_gemm_benchmark():
-    """Import validation utilities from commons directory."""
+    """Import gemm benchmark from parent directory."""
     current_dir = os.path.dirname(os.path.abspath(__file__))
     parent_dir = os.path.dirname(current_dir)
 

projects/composablekernel/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark_single.cpp

@@ -49,17 +49,17 @@ void benchmark_single(const ck_tile::ArgParser& arg_parser)
                              layout_c,
                              arg_parser.get_bool("structured_sparsity")};
 
-    // Create Setting struct
-    Setting setting{arg_parser.get_int("warmup"),
-                    arg_parser.get_int("repeat"),
-                    arg_parser.get_bool("timer"),
-                    arg_parser.get_int("verify"),
-                    arg_parser.get_int("init"),
-                    arg_parser.get_bool("log"),
-                    arg_parser.get_str("csv_filename"),
-                    arg_parser.get_bool("flush_cache"),
-                    arg_parser.get_int("rotating_count"),
-                    arg_parser.get_bool("json_output")};
+    // Create Settings struct
+    Settings setting{arg_parser.get_int("warmup"),
+                     arg_parser.get_int("repeat"),
+                     arg_parser.get_bool("timer"),
+                     arg_parser.get_int("verify"),
+                     arg_parser.get_int("init"),
+                     arg_parser.get_bool("log"),
+                     arg_parser.get_str("csv_filename"),
+                     arg_parser.get_bool("flush_cache"),
+                     arg_parser.get_int("rotating_count"),
+                     arg_parser.get_bool("json_output")};
 
     // Get the profiler instance
     auto& profiler = GemmPreshuffleProfiler::instance(setting);

projects/composablekernel/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_profiler.hpp

@@ -16,7 +16,7 @@ class GemmPreshuffleProfiler
     using BaseGemm = GemmProfiler<GemmPreshuffleProfiler, GemmProblem, ck_tile::GemmHostArgs>;
     using BaseGemm::benchmark;
 
-    GemmPreshuffleProfiler(Setting setting)
+    GemmPreshuffleProfiler(Settings setting)
         : GemmProfiler<GemmPreshuffleProfiler, GemmProblem, ck_tile::GemmHostArgs>(setting)
     {
     }
@@ -43,17 +43,17 @@ class GemmPreshuffleProfiler
         ck_tile::HostTensor<CDataType> c_m_n_dev_result(ck_tile::host_tensor_descriptor(
             gemm_problem.m_, gemm_problem.n_, gemm_problem.stride_c_, is_row_major(layout_c)));
 
-        if(setting_.init_method_ == 0)
+        if(setting_.init_method == 0)
         {
             ck_tile::FillUniformDistribution<ADataType>{-.5f, .5f}(a_m_k);
             ck_tile::FillUniformDistribution<BDataType>{-.5f, .5f}(b_k_n);
         }
-        else if(setting_.init_method_ == 1)
+        else if(setting_.init_method == 1)
         {
             ck_tile::FillMonotonicSeq<ADataType>{}(a_m_k);
             ck_tile::FillMonotonicSeq<BDataType>{}(b_k_n);
         }
-        else if(setting_.init_method_ == 2)
+        else if(setting_.init_method == 2)
         {
             ck_tile::FillUniformDistribution<ADataType>{1.f, 1.f}(a_m_k);
             ck_tile::FillUniformDistribution<BDataType>{1.f, 1.f}(b_k_n);
@@ -73,9 +73,9 @@ class GemmPreshuffleProfiler
             gemm_problem.m_, gemm_problem.n_, gemm_problem.stride_c_, is_row_major(layout_c)));
         c_m_n_ref.SetZero();
 
-        if(setting_.verify_)
+        if(setting_.verify)
         {
-            gemm_host_reference(setting_.verify_,
+            gemm_host_reference(setting_.verify,
                                 a_m_k,
                                 b_k_n,
                                 c_m_n_ref,
@@ -89,7 +89,7 @@ class GemmPreshuffleProfiler
                                 gemm_problem.stride_c_);
         }
 
-        // Kerenl Execution
+        // Kernel Execution
 
         a_m_k_dev_buf.ToDevice(a_m_k.data());
         c_m_n_dev_buf.SetZero();
@@ -126,12 +126,12 @@ class GemmPreshuffleProfiler
             auto kernel_run_result = callable(gemm_args,
                                               ck_tile::stream_config{nullptr,
                                                                      true,
-                                                                     setting_.log_,
-                                                                     setting_.n_warmup_,
-                                                                     setting_.n_repeat_,
-                                                                     setting_.is_gpu_timer_,
-                                                                     setting_.flush_cache_,
-                                                                     setting_.rotating_count_});
+                                                                     setting_.log,
+                                                                     setting_.n_warmup,
+                                                                     setting_.n_repeat,
+                                                                     setting_.is_gpu_timer,
+                                                                     setting_.flush_cache,
+                                                                     setting_.rotating_count});
 
             process_result(
                 gemm_problem, c_m_n_dev_buf, c_m_n_ref, c_m_n_dev_result, kernel_run_result);

projects/composablekernel/tile_engine/ops/gemm/gemm_profiler.hpp

@@ -19,7 +19,7 @@ template <typename Gemm, typename Problem, typename GemmArgs>
 class GemmProfiler
 {
     public:
-    static Gemm& instance(Setting setting)
+    static Gemm& instance(Settings setting)
     {
         static Gemm instance{setting};
         return instance;
@@ -68,7 +68,7 @@ class GemmProfiler
             ck_tile::static_for<0, DDataType::size(), 1>{}([&](auto i) {
                 using DType = ck_tile::remove_cvref_t<std::tuple_element_t<i, DDataType>>;
                 num_byte += sizeof(DType) * gemm_problem.m_ * gemm_problem.n_;
-                flop += sizeof(DType) * gemm_problem.m_ * gemm_problem.n_;
+                flop += gemm_problem.m_ * gemm_problem.n_;
             });
         }
 
@@ -77,7 +77,7 @@ class GemmProfiler
         kernel_instance.perf_result_.tflops_    = static_cast<float>(flop) / 1.E9 / avg_time;
         kernel_instance.perf_result_.bandwidth_ = num_byte / 1.E6 / avg_time;
 
-        if(setting_.log_ > 0 && !setting_.json_output_)
+        if(setting_.log > 0 && !setting_.json_output)
         {
             std::cout << kernel_instance << std::endl;
         }
@@ -90,7 +90,7 @@ class GemmProfiler
             split_k = gemm_problem.split_k_;
         }
         bool verified_correct =
-            !setting_.verify_ ||
+            !setting_.verify ||
             compare<Problem>(name, gemm_problem.k_, split_k, c_m_n_dev_result, c_m_n_host_result);
 
         if(verified_correct)
@@ -119,7 +119,7 @@ class GemmProfiler
                                                          b.perf_result_, a.perf_result_, metric);
                                                  });
 
-        if(setting_.json_output_)
+        if(setting_.json_output)
         {
             // Output clean JSON only
             std::cout << kernel_instance << std::endl;
@@ -132,9 +132,9 @@ class GemmProfiler
             std::cout << "**********************************" << std::endl;
         }
 
-        if(!setting_.csv_filename_.empty())
+        if(!setting_.csv_filename.empty())
         {
-            std::ofstream file(setting_.csv_filename_ + ".csv", std::ios::app);
+            std::ofstream file(setting_.csv_filename + ".csv", std::ios::app);
 
             if(!file.is_open())
             {
@@ -182,9 +182,9 @@ class GemmProfiler
 
     protected:
     virtual ~GemmProfiler() { kernel_instances_.clear(); }
-    GemmProfiler(Setting setting) : setting_(setting) {}
+    GemmProfiler(Settings setting) : setting_(setting) {}
 
-    Setting setting_;
+    Settings setting_;
 
     std::vector<KernelInstance<Problem>> kernel_instances_;
 };