commit: 9355a33a68cb7f92182362c46cbb4bf5a6499ee7 Author: Sv. Lockal <lockalsash <AT> gmail <DOT> com> AuthorDate: Sun Oct 12 07:24:21 2025 +0000 Commit: Sam James <sam <AT> gentoo <DOT> org> CommitDate: Tue Oct 21 18:47:31 2025 +0000 URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=9355a33a
sci-libs/composable-kernel: add 7.0.2 Bug: https://bugs.gentoo.org/964799 Signed-off-by: Sv. Lockal <lockalsash <AT> gmail.com> Part-of: https://github.com/gentoo/gentoo/pull/44165 Signed-off-by: Sam James <sam <AT> gentoo.org> sci-libs/composable-kernel/Manifest | 1 + .../composable-kernel-7.0.2.ebuild | 162 ++++++++ ...sable-kernel-7.0.1-conditional-ckprofiler.patch | 12 + .../composable-kernel-7.0.1-libcxx-includes.patch | 420 +++++++++++++++++++++ 4 files changed, 595 insertions(+) diff --git a/sci-libs/composable-kernel/Manifest b/sci-libs/composable-kernel/Manifest index 832cbc4fb864..2387faa9bed1 100644 --- a/sci-libs/composable-kernel/Manifest +++ b/sci-libs/composable-kernel/Manifest @@ -1,3 +1,4 @@ DIST composable-kernel-6.3.0.tar.gz 3376552 BLAKE2B ad3dde27d0134fa8bf6b5935603f7cac17c959ab7b123e8d69cfe82bee188ade0fe9d3666d46d9819de479e875e46c0fdad2574f68a3bcf0e9acbb871f163b24 SHA512 4fa7be3f54d06d894dca7ef9fa8e3db24971fb857fee96e7ccff89844fa466789756d5f140d2791a984311d59a044b5351d0358e676c218ea0b0aadf60954d8e DIST composable-kernel-6.4.3.tar.gz 4373178 BLAKE2B 1128f5de53e31365ff44d2c0b540ab662cd5c1659018d15f35e4c7036c8cf613ba5b95f910b7b3af94673a08ee36917fc38f9f19140140f2779252abda976fa4 SHA512 7073889f9af312a366eae9924771336e40b378255150a48308ce3975e254378da07bb49d615fbda7cc89dab35a14269df776ca88f483b803663cce325e463bb9 +DIST composable-kernel-7.0.2.tar.gz 5171930 BLAKE2B 2f54c21888ab8a0e76c0b4e5734b8e0c556b313e36ededed0610ea00ea8f0e4e72a757392cc2f47c84fa7fa7da0a72c41bd18b74b3a88b40e53ab7dd6ff2efbd SHA512 84c85ecad99d94a6c0bb159c1853d81e39fcb4d02c63a109cada84297f62b90aa9eb3d9102cb8e50b0c4f84982d8e5c8c8fdc5377acd014493621c9aec5439e0 DIST gtest-1.14.0_p20220421.tar.gz 849107 BLAKE2B e9c32d9c1d98959583b696430a6be411f5d69e3db96669b16cffc0f1dcad42512391c0a4733f6e829c82953ff33fa243ee782f3cf6e86436f0d290f45d1a3db7 SHA512 967aac7d85da0e216fe13b17d10f894a31d763d9b88201bac7d3dfc37600552b5472d30a166a6ef27f2778677e73fd3e43d082695a48f47824262eb9cd4fcd2f diff --git a/sci-libs/composable-kernel/composable-kernel-7.0.2.ebuild b/sci-libs/composable-kernel/composable-kernel-7.0.2.ebuild new file mode 100644 index 000000000000..e499d773ed60 --- /dev/null +++ b/sci-libs/composable-kernel/composable-kernel-7.0.2.ebuild @@ -0,0 +1,162 @@ +# Copyright 1999-2025 Gentoo Authors +# Distributed under the terms of the GNU General Public License v2 + +EAPI=8 + +ROCM_VERSION=${PV} +PYTHON_COMPAT=( python3_{10..14} python3_13t ) + +inherit check-reqs cmake flag-o-matic multiprocessing python-r1 rocm + +GTEST_COMMIT="b85864c64758dec007208e56af933fc3f52044ee" +GTEST_FILE="gtest-1.14.0_p20220421.tar.gz" + +DESCRIPTION="High Performance Composable Kernel for AMD GPUs" +HOMEPAGE="https://github.com/ROCm/composable_kernel"; +SRC_URI="https://github.com/ROCm/composable_kernel/archive/rocm-${PV}.tar.gz -> ${P}.tar.gz + test? ( https://github.com/google/googletest/archive/${GTEST_COMMIT}.tar.gz -> ${GTEST_FILE} )" +S="${WORKDIR}/composable_kernel-rocm-${PV}" + +LICENSE="MIT" +SLOT="0/$(ver_cut 1-2)" +KEYWORDS="~amd64" + +IUSE="debug profiler test" +REQUIRED_USE="${ROCM_REQUIRED_USE} ${PYTHON_REQUIRED_USE}" +RESTRICT="!test? ( test )" + +RDEPEND=" + dev-util/hip:${SLOT} + ${PYTHON_DEPS} +" + +DEPEND="${RDEPEND}" + +BDEPEND=" + dev-build/rocm-cmake +" + +PATCHES=( + "${FILESDIR}"/${PN}-6.1.1-no-git-no-hash.patch + "${FILESDIR}"/${PN}-6.3.0-conditional-kernels.patch + "${FILESDIR}"/${PN}-7.0.1-conditional-ckprofiler.patch + "${FILESDIR}"/${PN}-6.4.1-expand-isa.patch + "${FILESDIR}"/${PN}-7.0.1-libcxx-includes.patch +) + +ck_check-reqs() { + [[ ${MERGE_TYPE} == binary ]] && return + + targets=($AMDGPU_TARGETS) + if [[ ${#targets[@]} -gt 1 ]]; then + ewarn "composable-kernel will be compiled for multiple GPU architectures," + ewarn "which will take a significant amount of time." + ewarn "Please consider setting AMDGPU_TARGETS USE_EXPAND variable to a single architecture." + fi + + # It takes ~2Gb of RAM per build thread + local user_jobs=$(makeopts_jobs) + local available_memory_mb=$(free -m | awk '/Mem:/ {print $7}') + local max_jobs=$(( available_memory_mb / 2048 )) + max_jobs=$(( max_jobs < 1 ? 1 : max_jobs )) + local limited_jobs=$(( user_jobs < max_jobs ? user_jobs : max_jobs )) + if [[ "${max_jobs}" -lt "${user_jobs}" ]]; then + ewarn "${available_memory_mb} MB of free RAM is not enough for ${user_jobs} parallel build jobs (~2Gb per job)." + ewarn "Please consider setting MAKEOPTS=\"-j${limited_jobs}\" for this package." + fi + + local CHECKREQS_MEMORY=$((user_jobs*2048))M + check-reqs_${EBUILD_PHASE_FUNC} +} + +pkg_pretend() { + ck_check-reqs +} + +pkg_setup() { + ck_check-reqs +} + +src_prepare() { + sed -e '/-Werror/d' -i cmake/EnableCompilerWarnings.cmake || die + + # don't build examples + sed -e "/add_subdirectory(example)/d" -i CMakeLists.txt || die + + # Flag -amdgpu-early-inline-all explodes memory consumption + # https://github.com/llvm/llvm-project/issues/86332 + sed -e "/-amdgpu-early-inline-all/d" -e "/-amdgpu-function-calls/d" -i CMakeLists.txt || die + + cmake_src_prepare +} + +src_configure() { + rocm_use_hipcc + + if ! use debug; then + append-cflags "-DNDEBUG" + append-cxxflags "-DNDEBUG" + CMAKE_BUILD_TYPE="Release" + else + CMAKE_BUILD_TYPE="Debug" + fi + + local mycmakeargs=( + -DCMAKE_SKIP_RPATH=ON + -DBUILD_DEV=OFF + -DGPU_TARGETS="$(get_amdgpu_flags)" + -DCMAKE_INSTALL_PREFIX="${EPREFIX}/usr" + -DBUILD_TESTING=$(usex test ON OFF) + -DCK_USE_PROFILER=$(usex profiler ON OFF) + -Wno-dev + ) + + # Since 6.4.1 "fallback" DL kernels should be enabled manually... + if use amdgpu_targets_gfx1010 || use amdgpu_targets_gfx1011 || use amdgpu_targets_gfx1012 \ + || use amdgpu_targets_gfx1030 || use amdgpu_targets_gfx1031 ; then + mycmakeargs+=(-DDL_KERNELS=ON) + fi + + if use test; then + mycmakeargs+=( + -DFETCHCONTENT_SOURCE_DIR_GTEST="${WORKDIR}/googletest-${GTEST_COMMIT}" + ) + fi + + # rocminfo call during configuration; should not happen + # Bug: https://github.com/ROCm/composable_kernel/issues/2994 + check_amdgpu + addpredict /dev/random + + cmake_src_configure +} + +src_install() { + cmake_src_install + + # shellcheck disable=SC2329 + installation() { + python_domodule python/ck4inductor + + # install package-data manually, as there is no PEP517 compliance + shopt -s globstar + package_data=( + include/ck/**/*.hpp + library/src/tensor_operation_instance/gpu/gemm_universal/**/*.hpp + ) + shopt -u globstar + + inst_path="${D}$(python_get_sitedir)/ck4inductor" + for file in "${package_data[@]}"; do + location="${inst_path}/$(dirname "$file")" + mkdir -p "${location}" + cp "${file}" "${location}" + done + } + python_foreach_impl installation +} + +src_test() { + check_amdgpu + LD_LIBRARY_PATH="${BUILD_DIR}"/lib cmake_src_test -j1 +} diff --git a/sci-libs/composable-kernel/files/composable-kernel-7.0.1-conditional-ckprofiler.patch b/sci-libs/composable-kernel/files/composable-kernel-7.0.1-conditional-ckprofiler.patch new file mode 100644 index 000000000000..c9a762aae5ee --- /dev/null +++ b/sci-libs/composable-kernel/files/composable-kernel-7.0.1-conditional-ckprofiler.patch @@ -0,0 +1,12 @@ +Add a flag to build without ckprofiler, which takes few GB of space and not used in many cases. +--- a/CMakeLists.txt ++++ b/CMakeLists.txt +@@ -645,7 +645,7 @@ if(NOT GPU_ARCHS AND USER_GPU_TARGETS AND NOT MIOPEN_REQ_LIBS_ONLY) + endif() + endif() + +-if (NOT MIOPEN_REQ_LIBS_ONLY) ++if (NOT MIOPEN_REQ_LIBS_ONLY OR CK_USE_PROFILER) + rocm_package_setup_component(profiler + LIBRARY_NAME composablekernel + PACKAGE_NAME ckprofiler diff --git a/sci-libs/composable-kernel/files/composable-kernel-7.0.1-libcxx-includes.patch b/sci-libs/composable-kernel/files/composable-kernel-7.0.1-libcxx-includes.patch new file mode 100644 index 000000000000..f17d44aa614f --- /dev/null +++ b/sci-libs/composable-kernel/files/composable-kernel-7.0.1-libcxx-includes.patch @@ -0,0 +1,420 @@ +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_wmma_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_wmma_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <numeric> + #include <initializer_list> +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_wmma_cshuffle_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_wmma_cshuffle_v3.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl_fpAintB_b_scale.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl_fpAintB_b_scale.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_contraction_multiple_abd_xdl_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_contraction_multiple_abd_xdl_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <vector> + +--- a/include/ck/tensor_operation/gpu/device/impl/device_fpAintB_gemm_wmma.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_fpAintB_gemm_wmma.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_dpp.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_dpp.hpp +@@ -3,6 +3,7 @@ + + #pragma once + ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_abd_xdl_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_abd_xdl_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp +@@ -5,6 +5,7 @@ + + #ifndef __HIPCC_RTC__ + #include <iostream> ++#include <map> + #include <sstream> + #include "ck/host_utility/device_prop.hpp" + #include "ck/host_utility/kernel_launch.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_ab_scale.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_ab_scale.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_blockscale_bpreshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3_blockscale_bpreshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_lds_direct_load.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_lds_direct_load.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_streamk_v3.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v2.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v2.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_preshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_preshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_scale.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_scale.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_mx.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_mx.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3r1.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3r1.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <typeinfo> + +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_splitk_c_shuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_splitk_c_shuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_splitk_c_shuffle_lds_direct_load.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_splitk_c_shuffle_lds_direct_load.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <numeric> + #include <sstream> + +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp +@@ -6,6 +6,7 @@ + #include <functional> + #include <iostream> + #include <iterator> ++#include <map> + #include <numeric> + #include <sstream> + +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_xdl_cshuffle_tile_loop.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_xdl_cshuffle_tile_loop.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <tuple> + +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_query_attention_forward_wmma.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_query_attention_forward_wmma.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <numeric> + #include <initializer_list> +--- a/include/ck/tensor_operation/gpu/device/impl/device_moe_gemm.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_moe_gemm.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_moe_gemm_blockscale.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_moe_gemm_blockscale.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <hip/hip_runtime.h> + +--- a/include/ck/tensor_operation/gpu/device/impl/device_moe_mx_gemm.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_moe_mx_gemm.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_moe_mx_gemm_bns.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_moe_mx_gemm_bns.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_moe_mx_gemm_bpreshuffle.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_moe_mx_gemm_bpreshuffle.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + + #include "ck/utility/common_header.hpp" +--- a/include/ck/tensor_operation/gpu/device/impl/device_multi_query_attention_forward_wmma.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_multi_query_attention_forward_wmma.hpp +@@ -4,6 +4,7 @@ + #pragma once + + #include <iostream> ++#include <map> + #include <sstream> + #include <numeric> + #include <initializer_list>
