commit: 24b6befbdc895b920d307417a5067bfb90e5d5fa Author: Sv. Lockal <lockalsash <AT> gmail <DOT> com> AuthorDate: Fri Jul 11 07:37:55 2025 +0000 Commit: Alfredo Tupone <tupone <AT> gentoo <DOT> org> CommitDate: Fri Jul 11 10:04:59 2025 +0000 URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=24b6befb
sci-ml/caffe2: fix gfx101x compilation and memefficient linkage Closes: https://bugs.gentoo.org/959808 Bug: https://bugs.gentoo.org/956674 Signed-off-by: Sv. Lockal <lockalsash <AT> gmail.com> Part-of: https://github.com/gentoo/gentoo/pull/42956 Closes: https://github.com/gentoo/gentoo/pull/42956 Signed-off-by: Alfredo Tupone <tupone <AT> gentoo.org> ...ffe2-2.7.1-r2.ebuild => caffe2-2.7.1-r3.ebuild} | 6 + .../caffe2/files/caffe2-2.7.1-aotriton-fixes.patch | 27 ++++ .../files/composable-kernel-6.4.1-expand-isa.patch | 141 +++++++++++++++++++++ 3 files changed, 174 insertions(+) diff --git a/sci-ml/caffe2/caffe2-2.7.1-r2.ebuild b/sci-ml/caffe2/caffe2-2.7.1-r3.ebuild similarity index 97% rename from sci-ml/caffe2/caffe2-2.7.1-r2.ebuild rename to sci-ml/caffe2/caffe2-2.7.1-r3.ebuild index 4ccb6c07061c..c314b266cdc3 100644 --- a/sci-ml/caffe2/caffe2-2.7.1-r2.ebuild +++ b/sci-ml/caffe2/caffe2-2.7.1-r3.ebuild @@ -147,6 +147,7 @@ PATCHES=( "${FILESDIR}"/${PN}-2.7.0-glog-0.7.1.patch "${FILESDIR}"/${PN}-2.7.0-llvm.patch "${FILESDIR}"/${PN}-2.7.1-ck-config.patch + "${FILESDIR}"/${PN}-2.7.1-aotriton-fixes.patch ) src_prepare() { @@ -221,6 +222,11 @@ src_prepare() { sed -e "s:third_party/composable_kernel:../composable_kernel-${CK_COMMIT}:g" \ -i aten/src/ATen/CMakeLists.txt || die + # Bug 959808: fix for gfx101x targets + pushd "${WORKDIR}/composable_kernel-${CK_COMMIT}" > /dev/null || die + eapply "${FILESDIR}"/composable-kernel-6.4.1-expand-isa.patch + popd > /dev/null || die + if tc-is-clang; then # Systemwide gcc (for absl and at::TensorBase) + hipcc (llvm>=18) need abi-compat=17. # But systemwide clang>=18 + hipcc (>=llvm-18) need opposite! diff --git a/sci-ml/caffe2/files/caffe2-2.7.1-aotriton-fixes.patch b/sci-ml/caffe2/files/caffe2-2.7.1-aotriton-fixes.patch new file mode 100644 index 000000000000..1d2c7bf8f89d --- /dev/null +++ b/sci-ml/caffe2/files/caffe2-2.7.1-aotriton-fixes.patch @@ -0,0 +1,27 @@ +Fix installation with aotriton + +Upstream bug: https://github.com/pytorch/pytorch/issues/158109 +--- a/cmake/External/aotriton.cmake ++++ b/cmake/External/aotriton.cmake +@@ -43,10 +43,6 @@ if(NOT __AOTRITON_INCLUDED) + + # Note it is INSTALL"ED" + if(DEFINED ENV{AOTRITON_INSTALLED_PREFIX}) +- install(DIRECTORY +- $ENV{AOTRITON_INSTALLED_PREFIX}/lib64 +- $ENV{AOTRITON_INSTALLED_PREFIX}/include +- DESTINATION ${__AOTRITON_INSTALL_DIR}) + set(__AOTRITON_INSTALL_DIR "$ENV{AOTRITON_INSTALLED_PREFIX}") + message(STATUS "Using Preinstalled AOTriton at ${__AOTRITON_INSTALL_DIR}") + elseif(DEFINED ENV{AOTRITON_INSTALL_FROM_SOURCE}) +--- a/caffe2/CMakeLists.txt ++++ b/caffe2/CMakeLists.txt +@@ -921,7 +921,7 @@ if(USE_ROCM) + set(CUDA_LINK_LIBRARIES_KEYWORD PRIVATE) + list(APPEND Caffe2_HIP_SRCS ${GENERATED_CXX_TORCH_CUDA}) + hip_add_library(torch_hip ${Caffe2_HIP_SRCS}) +- if(USE_FLASH_ATTENTION) ++ if(USE_FLASH_ATTENTION OR USE_MEM_EFF_ATTENTION) + target_link_libraries(torch_hip PRIVATE __caffe2_aotriton) + endif() + set(CUDA_LINK_LIBRARIES_KEYWORD) diff --git a/sci-ml/caffe2/files/composable-kernel-6.4.1-expand-isa.patch b/sci-ml/caffe2/files/composable-kernel-6.4.1-expand-isa.patch new file mode 100644 index 000000000000..8a3fb4e1ec6d --- /dev/null +++ b/sci-ml/caffe2/files/composable-kernel-6.4.1-expand-isa.patch @@ -0,0 +1,141 @@ +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 +Bug: https://bugs.gentoo.org/show_bug.cgi?id=959808 +--- 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
