commit: 3046b420e43c7c4f4455ecde322fd339a30e7711 Author: Sv. Lockal <lockalsash <AT> gmail <DOT> com> AuthorDate: Wed Jun 11 09:01:08 2025 +0000 Commit: Sam James <sam <AT> gentoo <DOT> org> CommitDate: Sat Jun 14 22:37:11 2025 +0000 URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=3046b420
sci-libs/composable-kernel: add 6.4.1 Signed-off-by: Sv. Lockal <lockalsash <AT> gmail.com> Part-of: https://github.com/gentoo/gentoo/pull/42554 Signed-off-by: Sam James <sam <AT> gentoo.org> sci-libs/composable-kernel/Manifest | 1 + .../composable-kernel-6.4.1.ebuild | 150 +++++++++++++++++++++ .../files/composable-kernel-6.4.1-expand-isa.patch | 140 +++++++++++++++++++ 3 files changed, 291 insertions(+) diff --git a/sci-libs/composable-kernel/Manifest b/sci-libs/composable-kernel/Manifest index 7bdbbc67114a..4cab3c392dca 100644 --- a/sci-libs/composable-kernel/Manifest +++ b/sci-libs/composable-kernel/Manifest @@ -1,4 +1,5 @@ DIST composable-kernel-5.7.1.tar.gz 2158402 BLAKE2B 660c205c2a99cd17f29a9077e9e7b8b88e41f42ab40d92d0a235e7a60683a04234cce6b2eb6110265c2f461dd132c40151b2521b7d64df4cebb7f27094d8b2b0 SHA512 3931c5dbcab59513541103ebf7446311a1e318421a9e12aeafa7a98cf0717a75e1519140e6b2fbcaf4cb90829a11f03a724ff3ccc6636139c61667ac82558c55 DIST composable-kernel-6.1.1.tar.gz 2520967 BLAKE2B 549d76c5c36b0a273a4c51159abd162403ad035520f97c43dd4b66560ad30021199d2f4cf7e2c12ffe33beee2fb49604db6d7c40b3b33fdcbfde111b9d4833aa SHA512 3f7667cea91760447c3f6ac0bfc261c0bca95abafbfce209d1c889e04f01ec29849e8d7f0e1cb1d77a4309ac06b3c2557721a773e4a67be383de91668cb279b2 DIST composable-kernel-6.3.0.tar.gz 3376552 BLAKE2B ad3dde27d0134fa8bf6b5935603f7cac17c959ab7b123e8d69cfe82bee188ade0fe9d3666d46d9819de479e875e46c0fdad2574f68a3bcf0e9acbb871f163b24 SHA512 4fa7be3f54d06d894dca7ef9fa8e3db24971fb857fee96e7ccff89844fa466789756d5f140d2791a984311d59a044b5351d0358e676c218ea0b0aadf60954d8e +DIST composable-kernel-6.4.1.tar.gz 4373219 BLAKE2B d00b795a85c36f4199fcbcef7716699865fbbbde72e987dc378e19abb3866d6ccb2fd89095beab3de39009bfa036dda652d9ddf291d0734921c37c410b6c0bae SHA512 a36854e2f884d458847d053f8d1e4254f1f3227690e447d5d75f574df4d604b2a4b7a1b9d0d0c8048b2596ab6633b5a435a2e38352cc6c35117558a60d73a6f4 DIST gtest-1.14.0_p20220421.tar.gz 849107 BLAKE2B e9c32d9c1d98959583b696430a6be411f5d69e3db96669b16cffc0f1dcad42512391c0a4733f6e829c82953ff33fa243ee782f3cf6e86436f0d290f45d1a3db7 SHA512 967aac7d85da0e216fe13b17d10f894a31d763d9b88201bac7d3dfc37600552b5472d30a166a6ef27f2778677e73fd3e43d082695a48f47824262eb9cd4fcd2f diff --git a/sci-libs/composable-kernel/composable-kernel-6.4.1.ebuild b/sci-libs/composable-kernel/composable-kernel-6.4.1.ebuild new file mode 100644 index 000000000000..a05ecce8c100 --- /dev/null +++ b/sci-libs/composable-kernel/composable-kernel-6.4.1.ebuild @@ -0,0 +1,150 @@ +# Copyright 1999-2025 Gentoo Authors +# Distributed under the terms of the GNU General Public License v2 + +# shellcheck disable=SC2317 +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-enable-examples.patch + "${FILESDIR}"/${PN}-6.1.1-no-git-no-hash.patch + "${FILESDIR}"/${PN}-6.3.0-no-inline-all.patch + "${FILESDIR}"/${PN}-6.3.0-conditional-kernels.patch + "${FILESDIR}"/${PN}-6.3.0-conditional-ckprofiler.patch + "${FILESDIR}"/${PN}-6.4.1-expand-isa.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 free_memory_mb=$(free -m | awk '/Mem:/ {print $4}') + local max_jobs=$(( free_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 "${free_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 + 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 + + cmake_src_configure +} + +src_install() { + cmake_src_install + + 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-6.4.1-expand-isa.patch b/sci-libs/composable-kernel/files/composable-kernel-6.4.1-expand-isa.patch new file mode 100644 index 000000000000..48ffb25369f0 --- /dev/null +++ b/sci-libs/composable-kernel/files/composable-kernel-6.4.1-expand-isa.patch @@ -0,0 +1,140 @@ +Fix for "undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'" for AMDGPU_TARGETS="gfx1012". +Combines of 3 patches from https://github.com/ROCm/composable_kernel/issues/775#issuecomment-2726315348 + +Bug: https://bugs.gentoo.org/947583 +--- a/include/ck/ck.hpp ++++ b/include/ck/ck.hpp +@@ -82,7 +82,7 @@ CK_DECLARE_ENV_VAR_BOOL(CK_LOGGING) + #define CK_BUFFER_RESOURCE_3RD_DWORD -1 + #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__) + #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000 +-#elif defined(__gfx103__) ++#elif defined(__gfx101__) || defined(__gfx103__) + #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000 + #elif defined(__gfx11__) || defined(__gfx12__) + #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000 +@@ -90,12 +90,12 @@ CK_DECLARE_ENV_VAR_BOOL(CK_LOGGING) + + // FMA instruction + #ifndef __HIP_DEVICE_COMPILE__ // for host code, define nothing +-#elif defined(__gfx803__) || defined(__gfx900__) // for GPU code +-#define CK_USE_AMD_V_MAC_F32 +-#elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) // for GPU code ++#elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) || defined(__gfx1011__) || defined(__gfx1012__) // for GPU code + #define CK_USE_AMD_V_FMAC_F32 + #define CK_USE_AMD_V_DOT2_F32_F16 + #define CK_USE_AMD_V_DOT4_I32_I8 ++#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx101__) // for GPU code ++#define CK_USE_AMD_V_MAC_F32 + #elif defined(__gfx11__) || defined(__gfx12__) + #define CK_USE_AMD_V_FMAC_F32 + #define CK_USE_AMD_V_DOT2_F32_F16 +--- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_dl.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_dl.hpp +@@ -71,7 +71,7 @@ __global__ void + const Block2CTileMap block_2_ctile_map) + { + #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx908__) || \ +- defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx103__) || defined(__gfx11__) || \ ++ defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || \ + defined(__gfx12__)) + + const index_t num_blocks_per_batch = +--- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_dl.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_dl.hpp +@@ -51,7 +51,7 @@ __global__ void + const Block2CTileMap block_2_ctile_map) + { + #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx9__) || \ +- defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) ++ defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__)) + + constexpr index_t shared_block_size = + GridwiseGemm::GetSharedMemoryNumberOfByte() / sizeof(ABDataType); +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp +@@ -48,7 +48,7 @@ __global__ void + const Block2CTileMap block_2_ctile_map, + const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) + { +-#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx103__) || \ ++#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx101__) || defined(__gfx103__) || \ + defined(__gfx90a__) || defined(__gfx908__) || defined(__gfx94__) || defined(__gfx11__) || \ + defined(__gfx12__)) + const index_t num_blocks_per_batch = +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp +@@ -90,7 +90,7 @@ __global__ void + const Block2CTileMap block_2_ctile_map, + const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) + { +-#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx103__) || \ ++#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx101__) || defined(__gfx103__) || \ + defined(__gfx90a__) || defined(__gfx908__) || defined(__gfx94__) || defined(__gfx11__) || \ + defined(__gfx12__)) + // offset base pointer for each work-group +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp +@@ -106,7 +106,7 @@ __global__ void + const Block2CTileMap block_2_ctile_map, + const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch) + { +-#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx103__) || \ ++#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx101__) || defined(__gfx103__) || \ + defined(__gfx11__) || defined(__gfx12__)) + // offset base pointer for each work-group + const index_t num_blocks_per_batch = +--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp ++++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp +@@ -40,7 +40,7 @@ __global__ void + const CDEElementwiseOperation cde_element_op) + { + #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx908__) || \ +- defined(__gfx90a__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx94__) || \ ++ defined(__gfx90a__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx94__) || \ + defined(__gfx12__)) + __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()]; + +--- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp ++++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp +@@ -28,7 +28,7 @@ __global__ void + #endif + kernel_gemm_dpp(const typename GridwiseGemm::Argument karg) + { +-#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx103__) || defined(__gfx11__)) ++#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__)) + __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()]; + + const auto a_grid_desc_ak0_m_ak1 = amd_wave_read_first_lane( +--- a/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp ++++ b/include/ck/tensor_operation/gpu/grid/gridwise_tensor_rearrange.hpp +@@ -36,7 +36,7 @@ __global__ void + const ComputePtrOffsetOfStridedBatch compute_ptr_offset_of_batch) + { + #if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx906__) || defined(__gfx908__) || \ +- defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx103__) || defined(__gfx11__) || \ ++ defined(__gfx90a__) || defined(__gfx94__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || \ + defined(__gfx12__)) + GridwiseTensorRearrangeKernel::Run(in_grid_desc, + p_in_global, +--- a/include/ck_tile/core/config.hpp ++++ b/include/ck_tile/core/config.hpp +@@ -10,6 +10,9 @@ + #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__) + #define __gfx94__ + #endif ++#if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) ++#define __gfx101__ ++#endif + #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \ + defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \ + defined(__gfx10_3_generic__) +@@ -199,7 +202,7 @@ + #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || \ + defined(__gfx9__) // for GPU code + #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x00020000 +-#elif defined(__gfx103__) // for GPU code ++#elif defined(__gfx101__) || defined(__gfx103__) // for GPU code + #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31014000 + #elif defined(__gfx11__) || defined(__gfx12__) // for GPU code + #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31004000
