hipBLASLt: bump submodule to include gfx1151 support (#357)

Bump hipBLASLt submodule to include a revision that has gfx1151 support
and can build rocBLAS.

Adapt TheRock hipBLASLt patches to match.

* Disables rocRoller (see #376) project wide until dep issues can be
resolved.
* Adds a build dep on boost to avoid boost-too-old system dep issues
(with this hipBLASLt will use the boost in the super-project, which we
keep current).
* Adds a patch to hipBLASLt to only restore clocks on exit if rocm-smi
was found.

---------

Co-authored-by: Stella Laurenzo <stellaraccident@gmail.com>

Author: scottt

cmake/therock_amdgpu_targets.cmake

@@ -74,7 +74,6 @@ therock_add_amdgpu_target(gfx1150 "AMD Strix Point iGPU" FAMILY igpu-all gfx115X
 )
 therock_add_amdgpu_target(gfx1151 "AMD Strix Halo iGPU" FAMILY igpu-all gfx115X-all gfx115X-igpu
   EXCLUDE_TARGET_PROJECTS
-    hipBLASLt # https://github.com/ROCm/TheRock/issues/154
     rccl  # https://github.com/ROCm/TheRock/issues/150
 )
 

math-libs/BLAS/CMakeLists.txt

@@ -62,11 +62,14 @@ therock_cmake_subproject_declare(hipBLASLt
     -DHIPBLASLT_ENABLE_MARKER=OFF
     -DBUILD_CLIENTS_TESTS=${THEROCK_BUILD_TESTING}
     -DBUILD_CLIENTS_BENCHMARKS=${THEROCK_BUILD_TESTING}
+    # TODO(#376): rocRoller is broken
+    -DUSE_ROCROLLER=OFF
   COMPILER_TOOLCHAIN
     amd-hip
   BUILD_DEPS
     hipBLAS-common
     rocm-cmake
+    therock-boost
     therock-msgpack-cxx
   RUNTIME_DEPS
     hip-clr

math-libs/BLAS/hipBLASLt

@@ -1,7 +1,7 @@
-From e148639debe62686b3572b5aeccb50608464120a Mon Sep 17 00:00:00 2001
+From 28d221b97cf54db4b88bfbb923d6569a7065dd9f Mon Sep 17 00:00:00 2001
 From: Stella Laurenzo <stellaraccident@gmail.com>
 Date: Mon, 3 Feb 2025 20:28:58 -0800
-Subject: [PATCH 1/5] Update find_package for msgpack to work with 5.x and 6.x.
+Subject: [PATCH 1/7] Update find_package for msgpack to work with 5.x and 6.x.
 
 Adapted from: https://github.com/msgpack/msgpack-c/wiki/Q%26A#how-to-support-both-msgpack-c-c-version-5x-and-6x-
 
@@ -11,10 +11,10 @@ Note that 6.x also defines compile definitions that must be respected.
  1 file changed, 13 insertions(+), 15 deletions(-)
 
 diff --git a/tensilelite/Tensile/Source/lib/CMakeLists.txt b/tensilelite/Tensile/Source/lib/CMakeLists.txt
-index d055b635..a8a98690 100644
+index d6ad1df3..813f4848 100644
 --- a/tensilelite/Tensile/Source/lib/CMakeLists.txt
 +++ b/tensilelite/Tensile/Source/lib/CMakeLists.txt
-@@ -96,22 +96,20 @@ if(TENSILE_USE_LLVM OR TENSILE_USE_MSGPACK)
+@@ -105,22 +105,20 @@ if(TENSILE_USE_LLVM OR TENSILE_USE_MSGPACK)
  endif()
 
  if(TENSILE_USE_MSGPACK)
@@ -51,5 +51,5 @@ index d055b635..a8a98690 100644
 
  if(TENSILE_USE_LLVM)
 -- 
-2.34.1
+2.43.0
 

patches/amd-mainline/hipBLASLt/0001-Update-find_package-for-msgpack-to-work-with-5.x-and.patch

@@ -1,7 +1,7 @@
-From 75d321b60ad464b77004edeab95b5a6ed792f19a Mon Sep 17 00:00:00 2001
+From aba1654c5317edcfcef22166f3e5073016627eeb Mon Sep 17 00:00:00 2001
 From: Stella Laurenzo <stellaraccident@gmail.com>
 Date: Fri, 7 Mar 2025 12:13:41 -0800
-Subject: [PATCH 2/5] Do not hard-code hipBLASLt to find tools in /opt/rocm and
+Subject: [PATCH 2/7] Do not hard-code hipBLASLt to find tools in /opt/rocm and
  propagate configure time PATH.
 
 * Makes gen_assembly.sh and compile_code_object.sh scripts fail on error vs silently succeeding (`set -e`). These scripts have a suspect design overall, but making them error on failure is a minimal thing that is needed.
@@ -12,14 +12,14 @@ In general, hipBLASLt is very sloppy in terms of how it finds its tools. I resis
 
 This will have a subtle change that if the project was not build with an explicit ROCM_PATH and was somehow relying on falling back to /opt/rocm implicitly, that will now fail. As it should.
 ---
- .../src/kernels/CompileSourceKernel.cmake          |  4 +++-
- .../rocblaslt/src/kernels/compile_code_object.sh   | 14 ++++++++++++--
- tensilelite/Tensile/Ops/gen_assembly.sh            | 14 +++++++++-----
- tensilelite/Tensile/cmake/TensileConfig.cmake      | 10 +++++++++-
- 4 files changed, 33 insertions(+), 9 deletions(-)
+ .../src/kernels/CompileSourceKernel.cmake        |  6 ++++--
+ .../rocblaslt/src/kernels/compile_code_object.sh | 16 +++++++++++++---
+ tensilelite/Tensile/Ops/gen_assembly.sh          | 14 +++++++++-----
+ tensilelite/Tensile/cmake/TensileConfig.cmake    |  8 +++++++-
+ 4 files changed, 33 insertions(+), 11 deletions(-)
 
 diff --git a/library/src/amd_detail/rocblaslt/src/kernels/CompileSourceKernel.cmake b/library/src/amd_detail/rocblaslt/src/kernels/CompileSourceKernel.cmake
-index 13cf5bd2..27a3172d 100644
+index 13cf5bd2..ed20e560 100644
 --- a/library/src/amd_detail/rocblaslt/src/kernels/CompileSourceKernel.cmake
 +++ b/library/src/amd_detail/rocblaslt/src/kernels/CompileSourceKernel.cmake
 @@ -31,6 +31,8 @@ function(CompileSourceKernel source archs buildIdKind outputFolder)
@@ -31,10 +31,11 @@ index 13cf5bd2..27a3172d 100644
 +                         "${CMAKE_COMMAND}" -E env "'PATH=$ENV{PATH}'" --
 +                         bash  ${CMAKE_CURRENT_SOURCE_DIR}/src/amd_detail/rocblaslt/src/kernels/compile_code_object.sh ${source} ${archs} ${CMAKE_BUILD_TYPE} ${buildIdKind} ${outputFolder}/hipblasltTransform.hsaco
                         COMMENT "Compiling source kernels")
- endfunction()
+-endfunction()
 \ No newline at end of file
++endfunction()
 diff --git a/library/src/amd_detail/rocblaslt/src/kernels/compile_code_object.sh b/library/src/amd_detail/rocblaslt/src/kernels/compile_code_object.sh
-index 4ff71ea0..04ab923d 100644
+index 4ff71ea0..c40c7002 100644
 --- a/library/src/amd_detail/rocblaslt/src/kernels/compile_code_object.sh
 +++ b/library/src/amd_detail/rocblaslt/src/kernels/compile_code_object.sh
 @@ -20,6 +20,8 @@
@@ -52,6 +53,8 @@ index 4ff71ea0..04ab923d 100644
 
 -rocm_path="${ROCM_PATH:-/opt/rocm}"
 -clang_path="${rocm_path}/bin/amdclang++"
+-$clang_path -x hip "$sources" --offload-arch="${archs}" -c --offload-device-only -Xoffload-linker --build-id=$build_id_kind $additional_options -o "$dest"
+\ No newline at end of file
 +if [[ -z "$ROCM_PATH" ]]; then
 +    clang_path="$(which amdclang++)"
 +    if [[ -z "$clang_path" ]]; then
@@ -62,10 +65,9 @@ index 4ff71ea0..04ab923d 100644
 +    clang_path="${ROCM_PATH}/bin/amdclang++"
 +fi
 +
- $clang_path -x hip "$sources" --offload-arch="${archs}" -c --offload-device-only -Xoffload-linker --build-id=$build_id_kind $additional_options -o "$dest"
-\ No newline at end of file
++$clang_path -x hip "$sources" --offload-arch="${archs}" -c --offload-device-only -Xoffload-linker --build-id=$build_id_kind $additional_options -o "$dest"
 diff --git a/tensilelite/Tensile/Ops/gen_assembly.sh b/tensilelite/Tensile/Ops/gen_assembly.sh
-index fdc5960d..889e1410 100755
+index 4f0bf97e..9329f848 100755
 --- a/tensilelite/Tensile/Ops/gen_assembly.sh
 +++ b/tensilelite/Tensile/Ops/gen_assembly.sh
 @@ -22,19 +22,23 @@
@@ -98,33 +100,31 @@ index fdc5960d..889e1410 100755
 
  IFS=';' read -r -a archs <<< "$archStr"
 diff --git a/tensilelite/Tensile/cmake/TensileConfig.cmake b/tensilelite/Tensile/cmake/TensileConfig.cmake
-index cf6685a7..a5cb874a 100644
+index d542c584..1d57885f 100644
 --- a/tensilelite/Tensile/cmake/TensileConfig.cmake
 +++ b/tensilelite/Tensile/cmake/TensileConfig.cmake
 @@ -210,6 +210,12 @@ function(TensileCreateLibraryFiles
    endif()
 
-   set(CommandLine ${VIRTUALENV_BIN_DIR}/${VIRTUALENV_PYTHON_EXENAME} ${Script} ${Options} ${Tensile_LOGIC_PATH} ${Tensile_OUTPUT_PATH} HIP)
-+  # Tensile relies on the tools from the path, so capture the configure time 
-+  # path. It would be better if tool paths were explicit, but that would be a pretty 
+   set(CommandLine ${CMAKE_COMMAND} -E env PYTHONPATH=${PROJECT_BINARY_DIR}/lib -- ${VIRTUALENV_BIN_DIR}/${VIRTUALENV_PYTHON_EXENAME} ${Script} ${Options} ${Tensile_LOGIC_PATH} ${Tensile_OUTPUT_PATH} HIP)
++  # Tensile relies on the tools from the path, so capture the configure time
++  # path. It would be better if tool paths were explicit, but that would be a pretty
 +  # big change.
-+  set(CommandLine 
++  set(CommandLine
 +     "${CMAKE_COMMAND}" -E env "'PATH=$ENV{PATH}'" --
-+      ${CommandLine})  
++      ${CommandLine})
    message(STATUS "Tensile_CREATE_COMMAND: ${CommandLine}")
 
    if(Tensile_EMBED_LIBRARY)
-@@ -275,7 +281,9 @@ function(TensileCreateExtOpLibraries OutputFolder ArchStr)
+@@ -270,7 +276,7 @@ function(TensileCreateExtOpLibraries OutputFolder ArchStr TensileExt_LIBRARY_TAR
      COMMAND ${CMAKE_COMMAND} -E rm -rf ${build_tmp_dir}
      COMMAND ${CMAKE_COMMAND} -E make_directory ${build_tmp_dir}
      COMMAND ${CMAKE_COMMAND} -E make_directory ${OutputFolder}
--    COMMAND bash "${script}" "\"${Archs}\"" "${build_tmp_dir}" "${VIRTUALENV_HOME_DIR}" "${Tensile_BUILD_ID}"
-+    COMMAND
-+      "${CMAKE_COMMAND}" -E env "'PATH=$ENV{PATH}'" --    
-+      bash "${script}" "\"${Archs}\"" "${build_tmp_dir}" "${VIRTUALENV_HOME_DIR}" "${Tensile_BUILD_ID}"
+-    COMMAND ${CMAKE_COMMAND} -E env PYTHONPATH=${PROJECT_BINARY_DIR}/lib -- bash "${script}" "\"${Archs}\"" "${build_tmp_dir}" "${VIRTUALENV_HOME_DIR}" "${Tensile_BUILD_ID}"
++    COMMAND ${CMAKE_COMMAND} -E env "'PATH=$ENV{PATH}'" PYTHONPATH=${PROJECT_BINARY_DIR}/lib -- bash "${script}" "\"${Archs}\"" "${build_tmp_dir}" "${VIRTUALENV_HOME_DIR}" "${Tensile_BUILD_ID}"
      COMMAND ${CMAKE_COMMAND} -E copy ${ext_op_library_path} ${build_tmp_dir}/extop_*.co ${OutputFolder}
    )
 
 -- 
-2.34.1
+2.43.0
 

patches/amd-mainline/hipBLASLt/0002-Do-not-hard-code-hipBLASLt-to-find-tools-in-opt-rocm.patch

@@ -1,7 +1,7 @@
-From d150c267b4808eb0abab141116add579158d482d Mon Sep 17 00:00:00 2001
+From a50ba02dda81965d094aded79852b0d06fb5a8d6 Mon Sep 17 00:00:00 2001
 From: Marius Brehler <marius.brehler@amd.com>
 Date: Wed, 12 Mar 2025 22:39:56 +0000
-Subject: [PATCH 3/5] Find `rocm_smi` via config files
+Subject: [PATCH 3/7] Find `rocm_smi` via config files
 
 Use the config files provided by upstream instead of a custom finder.
 ---
@@ -11,7 +11,7 @@ Use the config files provided by upstream instead of a custom finder.
  delete mode 100644 tensilelite/Tensile/Source/cmake/FindROCmSMI.cmake
 
 diff --git a/clients/CMakeLists.txt b/clients/CMakeLists.txt
-index 02ae47d3..6b662094 100755
+index 31a72fe8..96d908e4 100755
 --- a/clients/CMakeLists.txt
 +++ b/clients/CMakeLists.txt
 @@ -97,8 +97,8 @@ if( BUILD_CLIENTS_BENCHMARKS OR BUILD_CLIENTS_TESTS)
@@ -88,5 +88,5 @@ index a6fa375a..00000000
 -
 -#set(rocm_smi_root "${hip_LIB_INSTALL_DIR}/../../rocm_smi")
 -- 
-2.34.1
+2.43.0
 

patches/amd-mainline/hipBLASLt/0003-Find-rocm_smi-via-config-files.patch

@@ -1,7 +1,7 @@
-From c92ac8d8319a6820b3d30f6bef7d70a4bcc64571 Mon Sep 17 00:00:00 2001
+From 88f840418f8421862fb83411f8c2eeb85cc8800e Mon Sep 17 00:00:00 2001
 From: Marius Brehler <marius.brehler@amd.com>
 Date: Wed, 12 Mar 2025 22:41:27 +0000
-Subject: [PATCH 4/5] Fix finding BLAS and LAPACK
+Subject: [PATCH 4/7] Fix finding BLAS and LAPACK
 
 So far it is assumed that NETLIB BLAS and LAPACK are avilable without
 checking. If they are nore, the build fails when trying to link.
@@ -14,7 +14,7 @@ OpenBLAS.
  3 files changed, 7 insertions(+), 5 deletions(-)
 
 diff --git a/clients/CMakeLists.txt b/clients/CMakeLists.txt
-index 6b662094..b24ca1c3 100755
+index 96d908e4..f69dee9f 100755
 --- a/clients/CMakeLists.txt
 +++ b/clients/CMakeLists.txt
 @@ -92,7 +92,9 @@ if( BUILD_CLIENTS_BENCHMARKS OR BUILD_CLIENTS_TESTS)
@@ -29,10 +29,10 @@ index 6b662094..b24ca1c3 100755
 
    # Find the package ROCmSMI
 diff --git a/clients/benchmarks/CMakeLists.txt b/clients/benchmarks/CMakeLists.txt
-index 99ec31f3..0d70d528 100644
+index e05b1f09..39aa4e8c 100644
 --- a/clients/benchmarks/CMakeLists.txt
 +++ b/clients/benchmarks/CMakeLists.txt
-@@ -76,9 +76,9 @@ endif()
+@@ -80,9 +80,9 @@ endif()
  # target_compile_options does not go to linker like CMAKE_CXX_FLAGS does, so manually add
  if (NOT WIN32)
    if (BUILD_CUDA)
@@ -45,10 +45,10 @@ index 99ec31f3..0d70d528 100644
    list( APPEND COMMON_LINK_LIBS "-lm -lstdc++fs")
 
 diff --git a/clients/gtest/CMakeLists.txt b/clients/gtest/CMakeLists.txt
-index cbedcfa4..eec47b6f 100644
+index 694670a3..64cd576a 100644
 --- a/clients/gtest/CMakeLists.txt
 +++ b/clients/gtest/CMakeLists.txt
-@@ -85,9 +85,9 @@ target_compile_options(hipblaslt-test PRIVATE $<$<COMPILE_LANGUAGE:CXX>:${COMMON
+@@ -90,9 +90,9 @@ target_compile_options(hipblaslt-test PRIVATE $<$<COMPILE_LANGUAGE:CXX>:${COMMON
  # target_compile_options does not go to linker like CMAKE_CXX_FLAGS does, so manually add
 
  if (BUILD_CUDA)
@@ -61,5 +61,5 @@ index cbedcfa4..eec47b6f 100644
 
  list( APPEND COMMON_LINK_LIBS "-lm -lstdc++fs")
 -- 
-2.34.1
+2.43.0
 

patches/amd-mainline/hipBLASLt/0004-Fix-finding-BLAS-and-LAPACK.patch

@@ -1,15 +1,15 @@
-From 4311b6c3ec948bd819a3413d83435307378ccaa8 Mon Sep 17 00:00:00 2001
+From a6becad02a9019ebef81ca4a5f451cab3e36b9c9 Mon Sep 17 00:00:00 2001
 From: Scott Todd <scott.todd0@gmail.com>
 Date: Thu, 6 Mar 2025 15:17:08 -0800
-Subject: [PATCH 5/5] Update 'packaging' in pip venv to stabilize setuptools.
+Subject: [PATCH 5/7] Update 'packaging' in pip venv to stabilize setuptools.
 
 When using setuptools>=71.0.0 and packaging<22.0, setuptools may fail with `TypeError: canonicalize_version() got an unexpected keyword argument 'strip_trailing_zero'`.
 ---
  cmake/virtualenv.cmake | 4 ++++
  1 file changed, 4 insertions(+)
 
 diff --git a/cmake/virtualenv.cmake b/cmake/virtualenv.cmake
-index 82bb6d94..a80fb337 100644
+index 7b4a78fe..96531ad2 100644
 --- a/cmake/virtualenv.cmake
 +++ b/cmake/virtualenv.cmake
 @@ -35,6 +35,10 @@ function(virtualenv_install)
@@ -24,5 +24,5 @@ index 82bb6d94..a80fb337 100644
        COMMAND ${VIRTUALENV_BIN_DIR}/${VIRTUALENV_PYTHON_EXENAME} -m pip install ${ARGN}
        COMMAND_ECHO STDOUT
 -- 
-2.34.1
+2.43.0
 

patches/amd-mainline/hipBLASLt/0005-Update-packaging-in-pip-venv-to-stabilize-setuptools.patch

@@ -1,20 +1,20 @@
-From fa0eb8bcd1c5cf46243a1e1d070172d7309dfb4a Mon Sep 17 00:00:00 2001
+From 294d4db2ad360dd6603cf7358d3410ad19f75e72 Mon Sep 17 00:00:00 2001
 From: Marius Brehler <marius.brehler@amd.com>
 Date: Tue, 18 Mar 2025 12:14:21 +0000
-Subject: [PATCH 6/6] Build `hipblas-sequence` on non-Ubuntu systems
+Subject: [PATCH 6/7] Build `hipblas-sequence` on non-Ubuntu systems
 
 ---
  clients/benchmarks/CMakeLists.txt | 64 +++++++++++++++----------------
  1 file changed, 30 insertions(+), 34 deletions(-)
 
 diff --git a/clients/benchmarks/CMakeLists.txt b/clients/benchmarks/CMakeLists.txt
-index 0d70d528..269cea79 100644
+index 39aa4e8c..1e0f8940 100644
 --- a/clients/benchmarks/CMakeLists.txt
 +++ b/clients/benchmarks/CMakeLists.txt
-@@ -107,42 +107,38 @@ add_executable( hipblaslt-bench-extop-matrixtransform client_extop_matrixtransfo
+@@ -118,42 +118,38 @@ add_executable( hipblaslt-bench-extop-matrixtransform client_extop_matrixtransfo
  add_executable( hipblaslt-bench-extop-softmax client_extop_softmax.cpp ../common/hipblaslt_random.cpp)
  add_executable( hipblaslt-bench-extop-amax client_extop_amax.cpp  ../common/hipblaslt_random.cpp)
- set(ext_bench_list_all hipblaslt-bench-groupedgemm-fixed-mk hipblaslt-bench-extop-layernorm hipblaslt-bench-extop-matrixtransform hipblaslt-bench-extop-softmax hipblaslt-bench-extop-amax)
+ set(ext_bench_list_all hipblaslt-api-overhead hipblaslt-bench-groupedgemm-fixed-mk hipblaslt-bench-extop-layernorm hipblaslt-bench-extop-matrixtransform hipblaslt-bench-extop-softmax hipblaslt-bench-extop-amax)
 -# Currently only build this for Ubuntu because other distro's llvm-dev version is not matching the rocm's version.
  if(${CMAKE_SYSTEM_NAME} STREQUAL "Linux")
 -  file(READ "/etc/os-release" OS_RELEASE)

patches/amd-mainline/hipBLASLt/0006-Build-hipblas-sequence-on-non-Ubuntu-systems.patch

@@ -0,0 +1,44 @@
+From 121342dcded7994a5fdf19d3256b50ae86f92454 Mon Sep 17 00:00:00 2001
+From: Stella Laurenzo <stellaraccident@gmail.com>
+Date: Tue, 8 Apr 2025 14:34:29 -0700
+Subject: [PATCH 7/7] Skip clock reset if rocm-smi is not available.
+
+* Pinning clocks was already conditioned on this, so the cleanup logic should also be.
+---
+ tensilelite/Tensile/Common/GlobalParameters.py | 12 +++++++++---
+ 1 file changed, 9 insertions(+), 3 deletions(-)
+
+diff --git a/tensilelite/Tensile/Common/GlobalParameters.py b/tensilelite/Tensile/Common/GlobalParameters.py
+index 8ec5f326..58dd69f4 100644
+--- a/tensilelite/Tensile/Common/GlobalParameters.py
++++ b/tensilelite/Tensile/Common/GlobalParameters.py
+@@ -549,7 +549,10 @@ def assignGlobalParameters(config, isaInfoMap: Dict[IsaVersion, IsaInfo]):
+         globalParameters["CmakeCCompiler"] = os.environ.get("CMAKE_C_COMPILER")
+ 
+     globalParameters["ROCmBinPath"] = os.path.join(globalParameters["ROCmPath"], "bin")
+-    globalParameters["ROCmSMIPath"] = locateExe(globalParameters["ROCmBinPath"], "rocm-smi")
++    try:
++        globalParameters["ROCmSMIPath"] = locateExe(globalParameters["ROCmBinPath"], "rocm-smi")
++    except OSError:
++        printWarning("Could not find rocm-smi: features relying on it will be ignored")
+     globalParameters["ROCmLdPath"] = locateExe(
+         os.path.join(globalParameters["ROCmPath"], "llvm/bin"), "ld.lld"
+     )
+@@ -624,9 +627,12 @@ def setupRestoreClocks():
+ 
+     def restoreClocks():
+         if globalParameters["PinClocks"]:
++            # Clocks will only be pinned if rocm-smi is available, therefore
++            # we only need to restore if found.
+             rsmi = globalParameters["ROCmSMIPath"]
+-            subprocess.call([rsmi, "-d", "0", "--resetclocks"])
+-            subprocess.call([rsmi, "-d", "0", "--setfan", "50"])
++            if rsmi is not None:
++                subprocess.call([rsmi, "-d", "0", "--resetclocks"])
++                subprocess.call([rsmi, "-d", "0", "--setfan", "50"])
+ 
+     atexit.register(restoreClocks)
+ 
+-- 
+2.43.0
+