repo/gentoo:master commit in: sci-ml/pytorch/, sci-ml/pytorch/files/
"Alfredo Tupone" <[email protected]> Wed, 05 Aug 2026 13:52:39 +0000 (UTC)
| Newsgroups | gmane.linux.gentoo.cvs |
|---|---|
| Message-ID | <1785937935.5d7bac9f2846fbce36798d587b3c8bc64028622f.tupone@gentoo> |
commit: 5d7bac9f2846fbce36798d587b3c8bc64028622f
Author: Sv. Lockal <lockalsash <AT> gmail <DOT> com>
AuthorDate: Wed Aug 5 09:51:10 2026 +0000
Commit: Alfredo Tupone <tupone <AT> gentoo <DOT> org>
CommitDate: Wed Aug 5 13:52:15 2026 +0000
URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=5d7bac9f
sci-ml/pytorch: more rocm fixes
- readd gfx1200 patch from caffe2
- sync CK_COMMIT with actual submodule commit of pytorch-2.13.0
- composable-kernel expand-isa is not needed, gfx101x was supported in upstream
- process mkl/openblas USE flags properly; otherwise it links automagically to whatever pytorch finds
- dev-cpp/cpp-httplib usage is header-only
- more libc++ memcpy fixes
Signed-off-by: Sv. Lockal <lockalsash <AT> gmail.com>
Part-of: https://github.com/gentoo/gentoo/pull/46649
Closes: https://github.com/gentoo/gentoo/pull/46649
Signed-off-by: Alfredo Tupone <tupone <AT> gentoo.org>
sci-ml/pytorch/Manifest | 2 +-
.../composable-kernel-7fe50dc-expand-isa.patch | 141 ---------------------
.../files/pytorch-2.13.0-rocm-assert-fix.patch | 47 +++++++
...h-2.13.0-r1.ebuild => pytorch-2.13.0-r2.ebuild} | 36 ++++--
4 files changed, 73 insertions(+), 153 deletions(-)
diff --git a/sci-ml/pytorch/Manifest b/sci-ml/pytorch/Manifest
index 1811ce5b902d..3433497817d6 100644
--- a/sci-ml/pytorch/Manifest
+++ b/sci-ml/pytorch/Manifest
@@ -1,4 +1,4 @@
-DIST composable_kernel-7fe50dc3.tar.gz 5380728 BLAKE2B c89c346d8e2d7a93a9cf26409e477fcdd25c43bc3f99d904c3bfe1bc282c6844ef2f2c80aceabe3bf4494db3457285384d5de5a22281aa426ba7479af82b0caf SHA512 a62f92e2dd7da944bd34bab6cf3bf624f630dc316d29c755e9fd523343c3f7648b0b7e0c9a0c8f5e9654477599ae8be9dac687d4054b0390f064ac2e40fc1cd3
+DIST composable_kernel-f1746955.tar.gz 7687217 BLAKE2B 973bf5b8eece2944e053c17eb3f8bf2d7f9a7744a651382c3d9677d6d4098f0107d64f3790899f74644681cdfb84e4989600c1e76bc1e96ee81579f870d55210 SHA512 dd03532d0487735d0ca6af5a37bf6a87917ab3ee81f70b14e29329d67a95f05bbccf18fb706ac9020918e7f6ac235e68b49164589fde073ff407a755daf26950
DIST flash-attention-2.7.4.gh.tar.gz 5841323 BLAKE2B 432999d763f2b3d732580ddfea5d3e01370351db0656546259a5e500a07516dd03c98828bfb55855dabe4adc651033b5d97ea4725ca46158b9970f0fbc662710 SHA512 05a4afb09e666f7404d6a3f8b5256e7bed6eba60a6f1bde2b7dbb96d318975f0b458c2521c7a38d88e97b6e4c27f29077cf787849daf82586e33f43a3d9a84b3
DIST pytorch-2.11.0.tar.gz 63504636 BLAKE2B d7227bd2d5475acf696bb48a3ca8cd8b7b6cd75fb99e64ebc71434e23817f94dea5164f48a176dcad1acd12186c33b02452582c1b41d6265b89cf54b29630670 SHA512 a5b27efdbce0506b49808e9dc5fbf1a56a22b1651b60236570823e1fe2c7395a0de9d3d7a01a568a201d593c52e7aa402e5d7645a54cb5932d5866e374ebf945
DIST pytorch-2.12.0.tar.gz 64740318 BLAKE2B e5e46a93422454da420563d0a8057c87d1ed4b27e0bd9b3b30e46b5e40cfbb489072144ac00267af2f085432d55e12adf750daf15ea165a1252449b00dca1ff9 SHA512 f5ab0f6933d88271f772b416f8c9b3b0d3e1ffaf8d00838b455206266b40d7c805e34003d84b01cddc7f1ad917dd72d6f79a05063b0962cbf223d9042cff3206
diff --git a/sci-ml/pytorch/files/composable-kernel-7fe50dc-expand-isa.patch b/sci-ml/pytorch/files/composable-kernel-7fe50dc-expand-isa.patch
deleted file mode 100644
index 53dcaf71fb44..000000000000
--- a/sci-ml/pytorch/files/composable-kernel-7fe50dc-expand-isa.patch
+++ /dev/null
@@ -1,141 +0,0 @@
-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
-@@ -78,7 +78,7 @@
- #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
-@@ -86,12 +86,12 @@
-
- // 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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const Block2CTileMap block_2_ctile_map)
- {
- #if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || \
-- defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__))
-+ defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__))
-
- const index_t num_blocks_per_batch =
- __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count);
---- 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
-@@ -50,7 +50,7 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const CGridDesc_M0_M10_M11_N0_N10_N11 e_grid_desc_m0_m10_m11_n0_n10_n11,
- const Block2CTileMap block_2_ctile_map)
- {
--#if(defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) || defined(__gfx11__) || \
-+#if(defined(__gfx906__) || defined(__gfx9__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || \
- defined(__gfx12__))
-
- constexpr index_t shared_block_size =
---- 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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const Block2CTileMap block_2_ctile_map,
- const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch)
- {
--#if(defined(__gfx906__) || defined(__gfx103__) || defined(__gfx90a__) || defined(__gfx908__) || \
-+#if(defined(__gfx906__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx90a__) || defined(__gfx908__) || \
- defined(__gfx94__) || defined(__gfx11__) || defined(__gfx12__))
- const index_t num_blocks_per_batch =
- __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count);
---- 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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const Block2CTileMap block_2_ctile_map,
- const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch)
- {
--#if(defined(__gfx906__) || defined(__gfx103__) || defined(__gfx90a__) || defined(__gfx908__) || \
-+#if(defined(__gfx906__) || defined(__gfx101__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx90a__) || defined(__gfx908__) || \
- defined(__gfx94__) || 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_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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const Block2CTileMap block_2_ctile_map,
- const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch)
- {
--#if(defined(__gfx906__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__))
-+#if(defined(__gfx906__) || defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__))
- // offset base pointer for each work-group
- const index_t num_blocks_per_batch =
- __builtin_amdgcn_readfirstlane(get_grid_size() / batch_count);
---- 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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const BElementwiseOperation b_element_op,
- const CDEElementwiseOperation cde_element_op)
- {
--#if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx103__) || \
-+#if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx101__) || 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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- #endif
- kernel_gemm_dpp(const typename GridwiseGemm::Argument karg)
- {
--#if(defined(__gfx103__) || defined(__gfx11__))
-+#if(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 @@ __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
- const ComputePtrOffsetOfStridedBatch compute_ptr_offset_of_batch)
- {
- #if(defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx94__) || \
-- defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__))
-+ defined(__gfx101__) || defined(__gfx103__) || defined(__gfx11__) || defined(__gfx12__))
- GridwiseTensorRearrangeKernel::Run(in_grid_desc,
- p_in_global,
- out_grid_desc,
---- a/include/ck_tile/core/config.hpp
-+++ b/include/ck_tile/core/config.hpp
-@@ -9,6 +9,9 @@
- #if 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__)
-@@ -200,7 +203,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
diff --git a/sci-ml/pytorch/files/pytorch-2.13.0-rocm-assert-fix.patch b/sci-ml/pytorch/files/pytorch-2.13.0-rocm-assert-fix.patch
new file mode 100644
index 000000000000..d9ba276fe777
--- /dev/null
+++ b/sci-ml/pytorch/files/pytorch-2.13.0-rocm-assert-fix.patch
@@ -0,0 +1,47 @@
+Fix compilation failures for targets like gfx1100.
+Loss.cu has a distinctive pattern of using CUDA_KERNEL_ASSERT after 64-bit getelementptr,
+which causes LLVM backend error. The problem is addressed in https://github.com/llvm/llvm-project/pull/174774,
+but requires a workaround for LLVM-22 (e.g., using __assert_fail, which does not emit llvm.trap()).
+
+See also: https://bugs.gentoo.org/978229
+--- a/aten/src/ATen/native/cuda/Loss.cu
++++ b/aten/src/ATen/native/cuda/Loss.cu
+@@ -157,12 +157,38 @@ int nll_loss_threads(int64_t nframe){
+ AT_PRIVATE_CASE_TYPE_USING_HINT(at::ScalarType::Byte, index_t, __VA_ARGS__) \
+ AT_PRIVATE_CASE_TYPE_USING_HINT(at::ScalarType::Long, index_t, __VA_ARGS__))
+
++// Workaround for LLVM bug in insertSimulatedTrap (fixed in
++// https://github.com/llvm/llvm-project/pull/174774): abort() generates
++// llvm.trap() which causes FinalizeISel to skip subsequent blocks containing
++// S_ADD_U64_PSEUDO, leaving it unexpanded. Using __assert_fail avoids
++// llvm.trap() while preserving the same assertion behavior.
++#ifdef USE_ROCM
++#define CHECK_INDEX_IN_CLASS(INDEX, N_CLASSES) \
++ if constexpr(std::is_unsigned_v<decltype(INDEX)>) { \
++ if (C10_UNLIKELY(!(INDEX < N_CLASSES))) { \
++ __assert_fail( \
++ #INDEX " < " #N_CLASSES, \
++ __FILE__, \
++ static_cast<unsigned int>(__LINE__), \
++ __func__); \
++ } \
++ } else { \
++ if (C10_UNLIKELY(!(INDEX >= 0 && INDEX < N_CLASSES))) { \
++ __assert_fail( \
++ #INDEX " >= 0 && " #INDEX " < " #N_CLASSES, \
++ __FILE__, \
++ static_cast<unsigned int>(__LINE__), \
++ __func__); \
++ } \
++ }
++#else
+ #define CHECK_INDEX_IN_CLASS(INDEX, N_CLASSES) \
+ if constexpr(std::is_unsigned_v<decltype(INDEX)>) { \
+ CUDA_KERNEL_ASSERT(INDEX < N_CLASSES); \
+ } else { \
+ CUDA_KERNEL_ASSERT(INDEX >= 0 && INDEX < N_CLASSES); \
+ }
++#endif
+
+ template <typename scalar_t, typename index_t>
+ __global__ void nll_loss_forward_no_reduce_cuda_kernel(
diff --git a/sci-ml/pytorch/pytorch-2.13.0-r1.ebuild b/sci-ml/pytorch/pytorch-2.13.0-r2.ebuild
similarity index 92%
rename from sci-ml/pytorch/pytorch-2.13.0-r1.ebuild
rename to sci-ml/pytorch/pytorch-2.13.0-r2.ebuild
index 3c39801b3546..55126d79388d 100644
--- a/sci-ml/pytorch/pytorch-2.13.0-r1.ebuild
+++ b/sci-ml/pytorch/pytorch-2.13.0-r2.ebuild
@@ -11,9 +11,9 @@ DISTUTILS_EXT=1
ROCM_VERSION=6.1
inherit distutils-r1 prefix cuda flag-o-matic rocm multiprocessing
-# pytorch-2.13.0 depends on future version of composable kernel
+# pytorch-2.13.0 depends on specific commit of composable kernel
# TODO: replace it with DEPEND in the future
-CK_COMMIT=7fe50dc3da2069d6645d9deb8c017a876472a977
+CK_COMMIT=f1746955fdaf80a3414de814bf32437686dac347
CK_P=composable_kernel-${CK_COMMIT:0:8}
# Starting from 2.7.0 pytorch moved flash attention out-of-tree,
@@ -76,6 +76,7 @@ RDEPEND="
dev-python/sympy[${PYTHON_USEDEP}]
dev-python/typing-extensions[${PYTHON_USEDEP}]
')
+ !mkl? ( !openblas? ( virtual/blas ) )
cuda? (
dev-libs/cudnn
>=sci-ml/cudnn-frontend-1.12.0:=
@@ -84,12 +85,12 @@ RDEPEND="
)
distributed? (
!rocm? ( sci-ml/tensorpipe[cuda?] )
- dev-cpp/cpp-httplib:=
)
fbgemm? ( >=sci-ml/FBGEMM-1.4 )
gloo? ( >=sci-ml/gloo-2025.06.04[cuda?,rocm?] )
kineto? ( ~sci-ml/kineto-0.4.0_p20260603 )
mimalloc? ( dev-libs/mimalloc )
+ mkl? ( sci-libs/mkl )
mpi? ( virtual/mpi )
nnpack? (
sci-ml/NNPACK
@@ -99,6 +100,7 @@ RDEPEND="
dev-python/numpy[${PYTHON_USEDEP}]
') )
onednn? ( sci-ml/oneDNN )
+ openblas? ( sci-libs/openblas )
opencl? ( virtual/opencl )
qnnpack? (
!sci-libs/QNNPACK
@@ -130,8 +132,6 @@ RDEPEND="
>=sci-ml/XNNPACK-2024.11
dev-libs/pthreadpool
)
- mkl? ( sci-libs/mkl )
- openblas? ( sci-libs/openblas )
"
DEPEND="
@@ -148,6 +148,7 @@ DEPEND="
dev-python/typing-extensions[${PYTHON_USEDEP}]
')
cuda? ( ~dev-libs/cutlass-4.4.2[tools(+)] )
+ distributed? ( dev-cpp/cpp-httplib:= )
onednn? ( sci-ml/ideep )
rocm? (
>=sci-libs/hipCUB-6.3:= <sci-libs/hipCUB-7.3:=
@@ -163,6 +164,7 @@ PATCHES=(
"${FILESDIR}"/${P}-aotriton-fixes.patch
"${FILESDIR}"/${P}-glog.patch
"${FILESDIR}"/${P}-removekineto-pr178960.patch
+ "${FILESDIR}"/${P}-rocm-assert-fix.patch
"${FILESDIR}"/${P}-unbundle_fbgemm.patch
"${FILESDIR}"/${P}-unbundle_flatbuffers.patch
"${FILESDIR}"/${P}-unbundle_fmt.patch
@@ -238,13 +240,10 @@ 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-7fe50dc-expand-isa.patch
- popd > /dev/null || die
-
# Workaround for libc++ issue https://github.com/llvm/llvm-project/issues/100802
- sed -e 's/std::memcpy/memcpy/g' -i torch/headeronly/util/Half.h || die
+ sed -e 's/std::memcpy/memcpy/g' \
+ -i torch/headeronly/util/Half.h \
+ -i aten/src/ATen/native/cuda/int4mm.cu || die
ebegin "HIPifying cuda sources"
FBCODE_BUILD_TOOL="buck" ${EPYTHON} tools/amd_build/build_amd.py || die
@@ -305,6 +304,21 @@ python_compile() {
local -x USE_XNNPACK=$(usex xnnpack)
local -x USE_XPU=OFF
+ if use mkl; then
+ local -x USE_BLAS=ON
+ local -x BLAS=MKL
+ elif use openblas; then
+ local -x USE_BLAS=ON
+ local -x BLAS=OpenBLAS
+ else
+ # Note: USE_BLAS=OFF does not properly work, instead pytorch still searches for libraries.
+ # Even though it is easy to patch pytorch to disable BLAS usage completely,
+ # it is practically not useful, as it would render the library barely usable.
+ local -x USE_BLAS=ON
+ local -x BLAS=Generic
+ local -x GENERIC_BLAS_LIBRARIES=cblas
+ fi
+
if use cuda; then
# bug 867706 926116
cuda_add_sandbox