Drive build of rocm-system packages from a single build file.

amdsmi.spec

@@ -1,22 +1,5 @@
 ## INCLUDE rocm-config
 ### RPM external amdsmi %{rocm_version_num}
-
-Source: %{rocm_systems_source}
 Requires: rocm-core python3
-
-%prep
-%setup -q -n rocm-systems
-
-%build
-
-cmake \
-  -S %{_builddir}/rocm-systems/projects/%{n} \
-  -B %{_builddir}/build \
-  -DCMAKE_INSTALL_PREFIX=%{i} \
-  -DCMAKE_BUILD_TYPE=Release \
-  -DCMAKE_PREFIX_PATH="%{cmake_prefix_path}" \
-  -DBUILD_TESTING=OFF
-
-make -C %{_builddir}/build %{makeprocesses}
-%install
-make -C %{_builddir}/build %{makeprocesses} install
+%define cmake_args -DBUILD_TESTING=OFF
+## INCLUDE rocm-systems-build

aqlprofile.spec

@@ -1,21 +1,4 @@
 ## INCLUDE rocm-config
 ### RPM external aqlprofile %{rocm_version_num}
-
-Source: %{rocm_systems_source}
 Requires: rocm-core rocr-runtime
-%prep
-%setup -q -n rocm-systems
-
-%build
-
-cmake \
-  -S %{_builddir}/rocm-systems/projects/%{n} \
-  -B %{_builddir}/build \
-  -DCMAKE_INSTALL_PREFIX=%{i} \
-  -DCMAKE_BUILD_TYPE=Release \
-  -DCMAKE_PREFIX_PATH=%{cmake_prefix_path}
-
-make -C %{_builddir}/build %{makeprocesses}
-
-%install
-make -C %{_builddir}/build %{makeprocesses} install
+## INCLUDE rocm-systems-build

hipify.spec

@@ -1,15 +1,15 @@
 ## INCLUDE rocm-config 
 ### RPM external hipify %{rocm_version_num}
-Source0: https://github.com/ROCm/HIPIFY/archive/refs/tags/rocm-%{rocm_version_num}.tar.gz
+Source0: git+https://github.com/ROCm/HIPIFY.git?obj=amd-develop/therock-%{rocm_version_num}&export=hipify&export=%{n}&submodules=1&output=/%{n}.tar.gz
 Requires: rocm-llvm
 
 %prep
-%setup -q -n HIPIFY-rocm-%{realversion}
+%setup -q -n %{n}
 
 %build
 CMAKE_ARGS=(
   -B %{_builddir}/build
-  -S %{_builddir}/HIPIFY-rocm-%{realversion}
+  -S %{_builddir}/%{n}
   -DCMAKE_INSTALL_PREFIX=%{i}
   -DCMAKE_BUILD_TYPE=Release \
   -DCMAKE_C_COMPILER=${ROCM_LLVM_ROOT}/lib/llvm/bin/clang \

miopen-ai-heuristic-guard-fix.patch

@@ -0,0 +1,185 @@
+--- projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
++++ projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
+@@ -31,11 +31,9 @@
+ #include <miopen/env.hpp>
+ #include <miopen/generic_search.hpp>
+ #include <miopen/conv/data_invoke_params.hpp>
+-#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+ #include <miopen/conv/heuristics/ai_heuristics.hpp>
+ #include <miopen/conv/heuristics/ai_candidate_selection.hpp>
+ #include <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp>
+-#endif
+ #include <miopen/solver/implicitgemm_ck_util_common.hpp>
+ #include <miopen/solver/ck_impl_lib_loader.hpp>
+ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS)
+--- projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp
++++ projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp
+@@ -32,11 +32,9 @@
+ #include <miopen/generic_search.hpp>
+ #include <miopen/conv/data_invoke_params.hpp>
+ #include <miopen/solver/problem_description_interpreter.hpp>
+-#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+ #include <miopen/conv/heuristics/ai_heuristics.hpp>
+ #include <miopen/conv/heuristics/ai_candidate_selection.hpp>
+ #include <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp>
+-#endif
+ #include <miopen/solver/implicitgemm_ck_util_common.hpp>
+ #include <miopen/solver/ck_impl_lib_loader.hpp>
+ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_GROUP_BWD_XDLOPS)
+--- projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp
++++ projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp
+@@ -33,11 +33,9 @@
+ #include <miopen/conv/wrw_invoke_params.hpp>
+ #include <miopen/solver/problem_description_interpreter.hpp>
+ #include <miopen/solver/ck_impl_lib_loader.hpp>
+-#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+ #include <miopen/conv/heuristics/ai_heuristics.hpp>
+ #include <miopen/conv/heuristics/ai_candidate_selection.hpp>
+ #include <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp>
+-#endif
+ #include <miopen/solver/implicitgemm_ck_util_common.hpp>
+ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS)
+ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_WRW_XDLOPS_AI_HEUR)
+--- projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp
++++ projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp
+@@ -34,11 +34,9 @@
+ #include <miopen/generic_search.hpp>
+ #include <miopen/conv/data_invoke_params.hpp>
+ #include <miopen/solver/problem_description_interpreter.hpp>
+-#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+ #include <miopen/conv/heuristics/ai_heuristics.hpp>
+ #include <miopen/conv/heuristics/ai_candidate_selection.hpp>
+ #include <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp>
+-#endif
+ #include <miopen/solver/implicitgemm_ck_util_common.hpp>
+ #include <miopen/solver/ck_impl_lib_loader.hpp>
+
+--- projects/miopen/src/include/miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp
++++ projects/miopen/src/include/miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp
+@@ -154,6 +154,50 @@
+     }
+ };
+
++/**
++ * @brief Configuration state for heuristic initialization results
++ *
++ * This struct holds the output state from AI heuristics. It's designed to be
++ * passed by reference to the centralized heuristics function, which will
++ * populate it with results.
++ *
++ * @note Kept outside the MIOPEN_ENABLE_AI_KERNEL_TUNING guard below because
++ * solvers use it unconditionally for fallback (non-AI) kernel selection state.
++ */
++struct HeuristicInitState
++{
++    std::vector<std::string>& valid_kernels;
++    int& index;
++    int& split_k;
++    std::string& kernel_id;
++
++    HeuristicInitState(std::vector<std::string>& vk, int& idx, int& sk, std::string& kid)
++        : valid_kernels(vk), index(idx), split_k(sk), kernel_id(kid)
++    {
++    }
++
++    void Reset(bool uses_split_k)
++    {
++        index     = 0;
++        kernel_id = "";
++        split_k   = uses_split_k ? 1 : 0;
++    }
++
++    void SetResult(int idx, int sk, bool uses_split_k)
++    {
++        index   = idx;
++        split_k = sk;
++        if(uses_split_k && idx >= 0 && idx < static_cast<int>(valid_kernels.size()))
++        {
++            kernel_id = valid_kernels[idx] + "+" + std::to_string(sk);
++        }
++        else if(idx >= 0 && idx < static_cast<int>(valid_kernels.size()))
++        {
++            kernel_id = valid_kernels[idx];
++        }
++    }
++};
++
+ } // namespace conv
+ } // namespace solver
+ } // namespace miopen
+@@ -257,47 +301,6 @@
+ // Centralized AI Heuristics Runner
+ // ============================================================================
+
+-/**
+- * @brief Configuration state for heuristic initialization results
+- *
+- * This struct holds the output state from AI heuristics. It's designed to be
+- * passed by reference to the centralized heuristics function, which will
+- * populate it with results.
+- */
+-struct HeuristicInitState
+-{
+-    std::vector<std::string>& valid_kernels;
+-    int& index;
+-    int& split_k;
+-    std::string& kernel_id;
+-
+-    HeuristicInitState(std::vector<std::string>& vk, int& idx, int& sk, std::string& kid)
+-        : valid_kernels(vk), index(idx), split_k(sk), kernel_id(kid)
+-    {
+-    }
+-
+-    void Reset(bool uses_split_k)
+-    {
+-        index     = 0;
+-        kernel_id = "";
+-        split_k   = uses_split_k ? 1 : 0;
+-    }
+-
+-    void SetResult(int idx, int sk, bool uses_split_k)
+-    {
+-        index   = idx;
+-        split_k = sk;
+-        if(uses_split_k && idx >= 0 && idx < static_cast<int>(valid_kernels.size()))
+-        {
+-            kernel_id = valid_kernels[idx] + "+" + std::to_string(sk);
+-        }
+-        else if(idx >= 0 && idx < static_cast<int>(valid_kernels.size()))
+-        {
+-            kernel_id = valid_kernels[idx];
+-        }
+-    }
+-};
+-
+ // ============================================================================
+ // Generic Helper Template for KTN Lambda Functions
+ // ============================================================================
+--- projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp
++++ projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp
+@@ -39,11 +39,9 @@
+ #include <miopen/conv/wrw_invoke_params.hpp>
+ #include <miopen/solver/problem_description_interpreter.hpp>
+ #include <miopen/solver/ck_impl_lib_loader.hpp>
+-#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+ #include <miopen/conv/heuristics/ai_heuristics.hpp>
+ #include <miopen/conv/heuristics/ai_candidate_selection.hpp>
+ #include <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp>
+-#endif
+ #include <miopen/solver/implicitgemm_ck_util_common.hpp>
+
+ namespace miopen {
+--- projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp
++++ projects/miopen/src/solver/conv/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp
+@@ -33,11 +33,9 @@
+ #include <miopen/generic_search.hpp>
+ #include <miopen/conv/data_invoke_params.hpp>
+ #include <miopen/solver/problem_description_interpreter.hpp>
+-#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+ #include <miopen/conv/heuristics/ai_heuristics.hpp>
+ #include <miopen/conv/heuristics/ai_candidate_selection.hpp>
+ #include <miopen/conv/heuristics/ai_conv_nd_kernel_tuning_utils.hpp>
+-#endif
+ #include <miopen/solver/implicitgemm_ck_util_common.hpp>
+ #include <miopen/solver/ck_impl_lib_loader.hpp>
+

miopen-boost-optional-fix.patch

@@ -1,75 +1,75 @@
+--- projects/miopen/src/anyramdb.cpp
++++ projects/miopen/src/anyramdb.cpp
+@@ -99,7 +99,7 @@ std::optional<AnyRamDb::TRecord> AnyRamDb::FindRecordUnsafe(const std::string& p
+     if(it == cache.end())
+         return {};
+
+-    return it->second;
++    return std::make_optional(it->second);
+ }
+
+ void AnyRamDb::UpdateCacheEntryUnsafe(const std::string& key, const TRecord& value)
+--- projects/miopen/src/db.cpp
++++ projects/miopen/src/db.cpp
+@@ -183,7 +183,7 @@ std::optional<DbRecord> PlainTextDb::FindRecordUnsafe(const std::string& key, Re
+             pos->begin = line_begin;
+             pos->end   = next_line_begin;
+         }
+-        return record;
++        return std::make_optional(record);
+     }
+     // Record was not found
+     return {};
+--- projects/miopen/src/find_controls.cpp
++++ projects/miopen/src/find_controls.cpp
+@@ -159,7 +159,7 @@ std::optional<std::vector<solver::Id>> GetEnvFindOnlySolverImpl()
+     if(res.empty())
+         return {};
+     else
+-        return {res};
++        return std::make_optional(res);
+ }
+
+ } // namespace
+--- projects/miopen/src/include/miopen/kern_db.hpp
++++ projects/miopen/src/include/miopen/kern_db.hpp
+@@ -128,7 +128,7 @@ public:
+             auto new_md5 = md5(decompressed_blob);
+             if(new_md5 != md5_hash)
+                 MIOPEN_THROW(miopenStatusInternalError, "Possible database corruption");
+-            return decompressed_blob;
++            return std::make_optional(decompressed_blob);
+         }
+         else if(rc == SQLITE_DONE)
+         {
 --- projects/miopen/src/include/miopen/sqlite_db.hpp
 +++ projects/miopen/src/include/miopen/sqlite_db.hpp
-@@ -463,7 +463,7 @@ class SQLitePerfDb
+@@ -437,7 +437,7 @@ public:
                  ovr_rec.SetValues(vals.at(0), vals.at(1));
              }
              if(success)
 -                return {ovr_rec};
-+                return boost::make_optional(ovr_rec);
++                return std::make_optional(ovr_rec);
          }
          std::string clause;
          std::vector<std::string> values;
-@@ -498,7 +498,7 @@ class SQLitePerfDb
+@@ -472,7 +472,7 @@ public:
          if(rec.GetSize() == 0)
-             return boost::none;
+             return {};
          else
 -            return {rec};
-+            return boost::make_optional(rec);
++            return std::make_optional(rec);
      }
-
+ 
      /// Removes ID with associated VALUES from record with key PROBLEM_CONFIG from db.
---- projects/miopen/src/find_controls.cpp
-+++ projects/miopen/src/find_controls.cpp
-@@ -161,7 +161,7 @@ boost::optional<std::vector<solver::Id>> GetEnvFindOnlySolverImpl()
-     if(res.empty())
-         return boost::none;
-     else
--        return {res};
-+        return boost::make_optional(res);
- }
-
- } // namespace
 --- projects/miopen/src/ramdb.cpp
 +++ projects/miopen/src/ramdb.cpp
-@@ -272,7 +272,7 @@ boost::optional<miopen::DbRecord> RamDb::FindRecordUnsafe(const std::string& pro
-         return boost::none;
+@@ -271,7 +271,7 @@ std::optional<miopen::DbRecord> RamDb::FindRecordUnsafe(const std::string& probl
+         return {};
      }
-
+ 
 -    return record;
-+    return boost::make_optional(record);
++    return std::make_optional(record);
  }
-
+ 
  template <class TFunc>
---- projects/miopen/src/db.cpp
-+++ projects/miopen/src/db.cpp
-@@ -209,7 +209,7 @@ boost::optional<DbRecord> PlainTextDb::FindRecordUnsafe(const std::string& key,
-             pos->begin = line_begin;
-             pos->end   = next_line_begin;
-         }
--        return record;
-+        return boost::make_optional(record);
-     }
-     // Record was not found
-     return boost::none;
---- projects/miopen/src/anyramdb.cpp
-+++ projects/miopen/src/anyramdb.cpp
-@@ -99,7 +99,7 @@ boost::optional<AnyRamDb::TRecord> AnyRamDb::FindRecordUnsafe(const std::string&
-     if(it == cache.end())
-         return boost::none;
-
--    return it->second;
-+    return boost::make_optional(it->second);
- }
-
- void AnyRamDb::UpdateCacheEntryUnsafe(const std::string& key, const TRecord& value)
---- projects/miopen/src/include/miopen/kern_db.hpp
-+++ projects/miopen/src/include/miopen/kern_db.hpp
-@@ -136,7 +136,7 @@ struct KernDb
-             auto new_md5 = md5(decompressed_blob);
-             if(new_md5 != md5_hash)
-                 MIOPEN_THROW(miopenStatusInternalError, "Possible database corruption");
--            return decompressed_blob;
-+            return boost::make_optional(decompressed_blob);
-         }
-         else if(rc == SQLITE_DONE)
-         {

miopen.spec

@@ -2,12 +2,13 @@
 ### RPM external miopen %{rocm_version_num}
 Source1: https://raw.githubusercontent.com/suruoxi/half/refs/heads/master/include/half.hpp
 Patch0: miopen-boost-optional-fix
+Patch1: miopen-ai-heuristic-guard-fix
 BuildRequires: rocm-cmake
 Requires: rocm-hip rocm-core rocr-runtime rocminfo python3 roctracer sqlite hipblaslt hipblas rocblas rocrand bz2lib hipblas
 Requires: json hipblas-common boost zstd opencl rocm-llvm rocm-comgr
 
 %define ROCMPrePrep cp %{_sourcedir}/half.hpp %{_builddir}
-%define ROCMPostPrep %patch0 -p0
+%define ROCMPostPrep patch -p0 -i %{PATCH0} ; patch -p0 -i %{PATCH1}
 %define ROCMPreBuild mkdir -p %{_builddir}/half-include/half && cp %{_sourcedir}/half.hpp %{_builddir}/half-include/half/
 %define ROCMPreCMake printf 'macro(enable_clang_tidy)\\nendmacro()\\nmacro(clang_tidy_check)\\nendmacro()\\n' > %{_builddir}/rocm-libraries/projects/%{n}/cmake/ClangTidy.cmake
 %define ROCMPostCMake sed -i '827,830d' %{_builddir}/rocm-libraries/projects/%{n}/CMakeLists.txt