From c28bc16f162e31ecf137d29a35bd58ffb45cfde7 Mon Sep 17 00:00:00 2001 From: awelling Date: Fri, 21 Aug 2026 12:10:13 -0500 Subject: [PATCH 1/3] Fix ROCm/HIP build and test compatibility on recent PyTorch and hipCUB. Updates GraphBolt and the tensor adapter to use PyTorch's CUDA-masquerading HIP APIs, fixes hipCUB 5.0 iterator compatibility, and hardens ROCm CI with dynamic test ports and isolated DataLoader test runs. --- CMakeLists.txt | 3 +- graphbolt/include/graphbolt/async.h | 10 +- graphbolt/src/cuda/common.h | 6 +- .../src/cuda/extension/gpu_graph_cache.cu | 9 +- .../cuda/extension/unique_and_compact_map.cu | 1 - include/dgl/hip/cuda_to_hip.h | 8 +- .../cuco-thrust-identity-removed.patch | 59 ++++++++++ .../cuco-thrust-pair-like-tuple-size.patch | 25 ++++ src/array/cuda/array_nonzero.cu | 6 + tensoradapter/pytorch/torch_hip.cpp | 107 ++++++++++++++++++ tests/python/pytorch/cuda/test_nccl.py | 9 +- .../pytorch/dataloading/test_dataloader.py | 10 +- .../pytorch/graphbolt/test_dataloader.py | 3 +- .../pytorch/graphbolt/test_item_sampler.py | 3 +- tests/python/pytorch/nn/test_sparse_emb.py | 3 +- tests/python/pytorch/optim/test_optim.py | 7 +- tests/python/pytorch/test_ffi-stream.py | 5 + tests/scripts/task_unit_test_rocm.sh | 80 +++++++++++-- tests/utils/__init__.py | 16 +++ 19 files changed, 324 insertions(+), 46 deletions(-) create mode 100644 script/patches/cuco-thrust-identity-removed.patch create mode 100644 script/patches/cuco-thrust-pair-like-tuple-size.patch create mode 100644 tensoradapter/pytorch/torch_hip.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index d10d9cd75d33..067d756210cd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -576,7 +576,7 @@ if (BUILD_SPARSE) GPU_TARGETS=${CMAKE_HIP_ARCHITECTURES} CMAKE_C_COMPILER=${CMAKE_C_COMPILER} CMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - CMAKE_PREFIX_PATH=${CMAKE_PREFIX_PATH} + CMAKE_PREFIX_PATH="${CMAKE_PREFIX_PATH}" CMAKE_HIP_ARCHITECTURES=${CMAKE_HIP_ARCHITECTURES} BINDIR=${CMAKE_CURRENT_BINARY_DIR} INCLUDEDIR="${DGL_INCLUDE_DIRS}" @@ -641,7 +641,6 @@ if (BUILD_GRAPHBOLT) PYTORCH_ROCM_ARCH=${CMAKE_HIP_ARCHITECTURES} CMAKE_C_COMPILER=${CMAKE_C_COMPILER} CMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} - CMAKE_PREFIX_PATH=${CMAKE_PREFIX_PATH} GPU_TARGETS=${CMAKE_HIP_ARCHITECTURES} CMAKE_HIP_ARCHITECTURES=${CMAKE_HIP_ARCHITECTURES} USE_HIP=${USE_HIP} diff --git a/graphbolt/include/graphbolt/async.h b/graphbolt/include/graphbolt/async.h index 7a938531afd2..f774728dcb33 100644 --- a/graphbolt/include/graphbolt/async.h +++ b/graphbolt/include/graphbolt/async.h @@ -47,17 +47,9 @@ #include #include #include +#include #include -namespace c10::cuda { -using CUDAStream = c10::hip::HIPStreamMasqueradingAsCUDA; -using CUDAStreamGuard = c10::hip::HIPStreamGuardMasqueradingAsCUDA; - -inline auto getCurrentCUDAStream() { - return c10::hip::getCurrentHIPStreamMasqueradingAsCUDA(); -} -} // namespace c10::cuda - #endif namespace graphbolt { diff --git a/graphbolt/src/cuda/common.h b/graphbolt/src/cuda/common.h index c71045bc5245..b2f179d167ec 100644 --- a/graphbolt/src/cuda/common.h +++ b/graphbolt/src/cuda/common.h @@ -33,10 +33,10 @@ #include #include -using namespace c10::hip; -using GPUStream_t = at::hip::HIPStreamMasqueradingAsCUDA; +using namespace c10::cuda; +using GPUStream_t = c10::cuda::CUDAStream; #define THRUST_BACKEND thrust::hip -#define GET_CURRENT_GPU_STREAM getCurrentHIPStreamMasqueradingAsCUDA +#define GET_CURRENT_GPU_STREAM getCurrentCUDAStream #else // we're using CUDA diff --git a/graphbolt/src/cuda/extension/gpu_graph_cache.cu b/graphbolt/src/cuda/extension/gpu_graph_cache.cu index 3990ac498713..1431a0e140b6 100644 --- a/graphbolt/src/cuda/extension/gpu_graph_cache.cu +++ b/graphbolt/src/cuda/extension/gpu_graph_cache.cu @@ -19,14 +19,16 @@ */ #include #include +#include #include #include #ifdef GRAPHBOLT_USE_HIP #include -#define C10_CUDA_KERNEL_LAUNCH_CHECK C10_HIP_KERNEL_LAUNCH_CHECK +#define GRAPHBOLT_DISCARD_OUTPUT_ITERATOR thrust::discard_iterator<>{} #else #include +#define GRAPHBOLT_DISCARD_OUTPUT_ITERATOR cub::DiscardOutputIterator{} #endif #include #include @@ -235,7 +237,7 @@ std::tuple GpuGraphCache::Query( output_indices.data_ptr()); CUB_CALL( DevicePartition::If, position_and_index, output_position_and_index, - cub::DiscardOutputIterator{}, seeds.size(0), + GRAPHBOLT_DISCARD_OUTPUT_ITERATOR, seeds.size(0), [] __device__(thrust::tuple & x) { return thrust::get<0>(x) >= 0; }); @@ -407,7 +409,8 @@ std::tuple> GpuGraphCache::Replace( CUB_CALL( DeviceSelect::Flagged, iota, is_threshold, output_indices.data_ptr(), - cub::DiscardOutputIterator{}, missing_positions.size(0)); + GRAPHBOLT_DISCARD_OUTPUT_ITERATOR, + missing_positions.size(0)); auto [in_degree, sliced_indptr] = ops::SliceCSCIndptr(indptr, output_indices); while (num_nodes_ + num_threshold >= indptr_.size(0)) { diff --git a/graphbolt/src/cuda/extension/unique_and_compact_map.cu b/graphbolt/src/cuda/extension/unique_and_compact_map.cu index 6a5625634065..0638ce83f452 100644 --- a/graphbolt/src/cuda/extension/unique_and_compact_map.cu +++ b/graphbolt/src/cuda/extension/unique_and_compact_map.cu @@ -25,7 +25,6 @@ #ifdef GRAPHBOLT_USE_HIP #include -#define C10_CUDA_KERNEL_LAUNCH_CHECK C10_HIP_KERNEL_LAUNCH_CHECK #else #include #include diff --git a/include/dgl/hip/cuda_to_hip.h b/include/dgl/hip/cuda_to_hip.h index 6cc4ac0870e1..13108844ea8e 100644 --- a/include/dgl/hip/cuda_to_hip.h +++ b/include/dgl/hip/cuda_to_hip.h @@ -222,10 +222,8 @@ using cudaStream_t = hipStream_t; #define cusparseXcsrsort hipsparseXcsrsort #define cusparseXcsrsort_bufferSizeExt hipsparseXcsrsort_bufferSizeExt -// Used in graphbolt -#define CUDACachingAllocator HIPCachingAllocator -#define getCurrentCUDAStream getCurrentHIPStream -#define C10_CUDA_CHECK C10_HIP_CHECK -#define C10_CUDA_KERNEL_LAUNCH_CHECK C10_HIP_KERNEL_LAUNCH_CHECK +// Torch's hipified c10/ATen headers already declare the CUDA-named stream, +// event, allocator and error-check APIs inside c10::cuda / at::cuda, so +// redirecting them here would alias them onto HIP names that do not exist. #endif diff --git a/script/patches/cuco-thrust-identity-removed.patch b/script/patches/cuco-thrust-identity-removed.patch new file mode 100644 index 000000000000..b2dd7d5aae7d --- /dev/null +++ b/script/patches/cuco-thrust-identity-removed.patch @@ -0,0 +1,59 @@ +Replace thrust::identity with cuda::std::identity. + +rocThrust 5.0 (ROCm 10) removed thrust::identity, which cuco 25.10 still uses +as a pass-through predicate and as the base of identity_hash. cuda::std::identity +is the CCCL replacement and is available from libhipcxx in ROCm 10. Note it is +not a class template, so identity_hash inherits from it unparameterized and +relies on its templated call operator. + +--- a/include/cuco/detail/open_addressing/open_addressing_impl.cuh ++++ b/include/cuco/detail/open_addressing/open_addressing_impl.cuh +@@ -56,4 +56,5 @@ + #include ++#include + #include + #include + #include +@@ -299,1 +299,1 @@ +- return this->insert_if(first, last, always_true, thrust::identity{}, container_ref, stream); ++ return this->insert_if(first, last, always_true, cuda::std::identity{}, container_ref, stream); +@@ -322,1 +322,1 @@ +- this->insert_if_async(first, last, always_true, thrust::identity{}, container_ref, stream); ++ this->insert_if_async(first, last, always_true, cuda::std::identity{}, container_ref, stream); +@@ -520,1 +520,1 @@ +- first, last, always_true, thrust::identity{}, output_begin, container_ref, stream); ++ first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream); +@@ -595,1 +595,1 @@ +- first, last, always_true, thrust::identity{}, output_begin, container_ref, stream); ++ first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream); +--- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh ++++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh +@@ -44,1 +44,2 @@ + #include ++#include +@@ -1271,1 +1271,1 @@ +- if (thrust::any_of(thrust::seq, exists, exists + bucket_size, thrust::identity{})) { ++ if (thrust::any_of(thrust::seq, exists, exists + bucket_size, cuda::std::identity{})) { +--- a/include/cuco/detail/hash_functions/identity_hash.cuh ++++ b/include/cuco/detail/hash_functions/identity_hash.cuh +@@ -19,2 +19,3 @@ + #include ++#include + #include +@@ -36,1 +36,1 @@ +-struct identity_hash : private thrust::identity { ++struct identity_hash : private cuda::std::identity { +@@ -52,1 +52,1 @@ +- return static_cast(thrust::identity::operator()(x)); ++ return static_cast(cuda::std::identity::operator()(x)); +--- a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh ++++ b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh +@@ -56,1 +56,2 @@ + #include ++#include +@@ -208,1 +208,1 @@ +- this->add_if_async(first, last, always_true, thrust::identity{}, stream); ++ this->add_if_async(first, last, always_true, cuda::std::identity{}, stream); +@@ -309,1 +309,1 @@ +- this->contains_if_async(first, last, always_true, thrust::identity{}, output_begin, stream); ++ this->contains_if_async(first, last, always_true, cuda::std::identity{}, output_begin, stream); diff --git a/script/patches/cuco-thrust-pair-like-tuple-size.patch b/script/patches/cuco-thrust-pair-like-tuple-size.patch new file mode 100644 index 000000000000..5f80c323e6c4 --- /dev/null +++ b/script/patches/cuco-thrust-pair-like-tuple-size.patch @@ -0,0 +1,25 @@ +Make cuco's is_thrust_pair_like_impl detection SFINAE-safe. + +The partial specialization only constrained thrust::get<0>/get<1>, then +evaluated thrust::tuple_size::value in its base-clause. Under ROCm 10's +libhipcxx, cuda::std::tuple_size has an undefined primary template, so for a +non-tuple T (e.g. the double desired_load_factor argument considered during +overload resolution for cuco::pair's converting constructor) this is a hard +error instead of a substitution failure. + +Constrain tuple_size in the void_t as well, matching the existing +is_cuda_std_pair_like_impl specialization directly above it. + +--- a/include/cuco/detail/pair/traits.hpp ++++ b/include/cuco/detail/pair/traits.hpp +@@ -74,8 +74,9 @@ + template + struct is_thrust_pair_like_impl< + T, + cuda::std::void_t(cuda::std::declval())), +- decltype(thrust::get<1>(cuda::std::declval()))>> ++ decltype(thrust::get<1>(cuda::std::declval())), ++ decltype(thrust::tuple_size::value)>> + : cuda::std::conditional_t::value == 2, + cuda::std::true_type, + cuda::std::false_type> {}; diff --git a/src/array/cuda/array_nonzero.cu b/src/array/cuda/array_nonzero.cu index a7f0d61c4355..46427a6bb06d 100644 --- a/src/array/cuda/array_nonzero.cu +++ b/src/array/cuda/array_nonzero.cu @@ -12,6 +12,7 @@ #include #include +#include #endif #include "../../runtime/cuda/cuda_common.h" @@ -45,7 +46,12 @@ IdArray NonZero(IdArray array) { int64_t* const out_data = static_cast(ret->data); IsNonZeroIndex comp(in_data); +#if defined(__HIPCC__) + // hipCUB 5.0 dropped CountingInputIterator. + thrust::counting_iterator counter(0); +#else cub::CountingInputIterator counter(0); +#endif // room for cub to output on GPU int64_t* d_num_nonzeros = diff --git a/tensoradapter/pytorch/torch_hip.cpp b/tensoradapter/pytorch/torch_hip.cpp new file mode 100644 index 000000000000..ef96d8b94b44 --- /dev/null +++ b/tensoradapter/pytorch/torch_hip.cpp @@ -0,0 +1,107 @@ +// !!! This is a file automatically generated by hipify!!! +/** + * Copyright (c) 2020-2022 by Contributors + * @file torch/torch.cpp + * @brief Implementation of PyTorch adapter library. + */ + +#include +#include +#if defined DGL_USE_CUDA || defined DGL_USE_HIP +#include +#include +#include +#include +#include +#endif // DGL_USE_CUDA + +namespace tensoradapter { + +extern "C" { + +TA_EXPORTS void* CPURawAlloc(size_t nbytes) { + return c10::GetCPUAllocator()->raw_allocate(nbytes); +} + +TA_EXPORTS void CPURawDelete(void* ptr) { + c10::GetCPUAllocator()->raw_deallocate(ptr); +} + +#if defined DGL_USE_CUDA || defined DGL_USE_HIP +TA_EXPORTS void* CUDARawAlloc(size_t nbytes, hipStream_t stream) { + at::globalContext().lazyInitDevice(at::kCUDA); + return c10::cuda::CUDACachingAllocator::raw_alloc_with_stream(nbytes, stream); +} + +TA_EXPORTS void CUDARawDelete(void* ptr) { + c10::cuda::CUDACachingAllocator::raw_delete(ptr); +} + +TA_EXPORTS hipStream_t CUDACurrentStream() { + return at::cuda::getCurrentCUDAStream(); +} + +TA_EXPORTS void RecordStream(void* ptr, hipStream_t stream, int device_id) { + c10::DataPtr data_ptr{ + ptr, ptr, c10::cuda::CUDACachingAllocator::get()->raw_deleter(), + c10::Device(c10::DeviceType::CUDA, device_id)}; + c10::cuda::CUDACachingAllocator::recordStream( + data_ptr, + // getStreamFromExternal doesn't exist before PyTorch 1.10, just copy it + // here + c10::cuda::CUDAStream( + c10::cuda::CUDAStream::UNCHECKED, + c10::Stream( + c10::Stream::UNSAFE, + c10::Device(c10::DeviceType::CUDA, device_id), + reinterpret_cast(stream)))); + data_ptr.release_context(); +} + +class CUDAHostDeleter { + public: + explicit CUDAHostDeleter(std::unique_ptr ptr) + : ptr_(std::move(ptr)) {} + + private: + std::unique_ptr ptr_; +}; + +TA_EXPORTS void* CUDARawHostAlloc( + size_t nbytes, void** ctx, void** raw_deleter) { + auto data_ptr = at::cuda::getCachingHostAllocator()->allocate(nbytes); + auto raw = data_ptr.get(); + // Return the raw ctx ptr for recording event. + *ctx = data_ptr.get_context(); + + // Transfer ownership to raw_deleter. + auto* data_deleter = new CUDAHostDeleter(data_ptr.move_context()); + *raw_deleter = static_cast(data_deleter); + return raw; +} + +// Designated CUDAHostDeleter for CUDARawHostAlloc. +TA_EXPORTS void CUDARawHostDelete(void** raw_deleter) { + delete static_cast(*raw_deleter); + *raw_deleter = nullptr; +} + +TA_EXPORTS void CUDARecordHostAlloc( + void* ptr, void* ctx, hipStream_t stream, int device_id) { + at::cuda::CachingHostAllocator_recordEvent( + ptr, ctx, + c10::cuda::CUDAStream( + c10::cuda::CUDAStream::UNCHECKED, + c10::Stream( + c10::Stream::UNSAFE, + c10::Device(c10::DeviceType::CUDA, device_id), + reinterpret_cast(stream)))); +} + +TA_EXPORTS void CUDAHostAllocatorEmptyCache() { + at::cuda::CachingHostAllocator_emptyCache(); +} +#endif // DGL_USE_CUDA +}; + +}; // namespace tensoradapter diff --git a/tests/python/pytorch/cuda/test_nccl.py b/tests/python/pytorch/cuda/test_nccl.py index 354682c13f9c..7027b8de0adc 100644 --- a/tests/python/pytorch/cuda/test_nccl.py +++ b/tests/python/pytorch/cuda/test_nccl.py @@ -3,6 +3,7 @@ import backend as F import torch import torch.distributed as dist +from utils import rendezvous_port from dgl.cuda import nccl from dgl.partition import NDArrayPartition @@ -15,7 +16,7 @@ def test_nccl_sparse_push_single_remainder(): torch.cuda.set_device("cuda:0") dist.init_process_group( backend="nccl", - init_method="tcp://127.0.0.1:12345", + init_method=f"tcp://127.0.0.1:{rendezvous_port()}", world_size=1, rank=0, ) @@ -39,7 +40,7 @@ def test_nccl_sparse_pull_single_remainder(): torch.cuda.set_device("cuda:0") dist.init_process_group( backend="nccl", - init_method="tcp://127.0.0.1:12345", + init_method=f"tcp://127.0.0.1:{rendezvous_port()}", world_size=1, rank=0, ) @@ -63,7 +64,7 @@ def test_nccl_sparse_push_single_range(): torch.cuda.set_device("cuda:0") dist.init_process_group( backend="nccl", - init_method="tcp://127.0.0.1:12345", + init_method=f"tcp://127.0.0.1:{rendezvous_port()}", world_size=1, rank=0, ) @@ -90,7 +91,7 @@ def test_nccl_sparse_pull_single_range(): torch.cuda.set_device("cuda:0") dist.init_process_group( backend="nccl", - init_method="tcp://127.0.0.1:12345", + init_method=f"tcp://127.0.0.1:{rendezvous_port()}", world_size=1, rank=0, ) diff --git a/tests/python/pytorch/dataloading/test_dataloader.py b/tests/python/pytorch/dataloading/test_dataloader.py index 961eb19ed75f..c1c95bc8f420 100644 --- a/tests/python/pytorch/dataloading/test_dataloader.py +++ b/tests/python/pytorch/dataloading/test_dataloader.py @@ -12,7 +12,7 @@ import torch import torch.distributed as dist import torch.multiprocessing as mp -from utils import parametrize_idtype +from utils import parametrize_idtype, rendezvous_port @pytest.mark.parametrize("batch_size", [None, 16]) @@ -109,7 +109,7 @@ def test_neighbor_nonuniform(idtype, mode, use_ddp, use_mask): pytest.skip("PyTorch 1.13.0+ has problems in Windows DDP...") dist.init_process_group( "gloo" if F.ctx() == F.cpu() else "nccl", - "tcp://127.0.0.1:12347", + f"tcp://127.0.0.1:{rendezvous_port(2)}", world_size=1, rank=0, ) @@ -273,7 +273,7 @@ def _ddp_runner(proc_id, nprocs, g, data, args): dist.init_process_group( "nccl" if mode != "cpu" else "gloo", - "tcp://127.0.0.1:12347", + f"tcp://127.0.0.1:{rendezvous_port(2)}", world_size=nprocs, rank=proc_id, ) @@ -338,7 +338,7 @@ def test_node_dataloader( pytest.skip("PyTorch 1.13.0+ has problems in Windows DDP...") dist.init_process_group( "gloo" if F.ctx() == F.cpu() else "nccl", - "tcp://127.0.0.1:12347", + f"tcp://127.0.0.1:{rendezvous_port(2)}", world_size=1, rank=0, ) @@ -476,7 +476,7 @@ def test_edge_dataloader(idtype, sampler_name, neg_sampler, mode, use_ddp): pytest.skip("PyTorch 1.13.0+ has problems in Windows DDP...") dist.init_process_group( "gloo" if F.ctx() == F.cpu() else "nccl", - "tcp://127.0.0.1:12347", + f"tcp://127.0.0.1:{rendezvous_port(2)}", world_size=1, rank=0, ) diff --git a/tests/python/pytorch/graphbolt/test_dataloader.py b/tests/python/pytorch/graphbolt/test_dataloader.py index ac5f49995437..c6cbae395e3c 100644 --- a/tests/python/pytorch/graphbolt/test_dataloader.py +++ b/tests/python/pytorch/graphbolt/test_dataloader.py @@ -13,6 +13,7 @@ import torch import torch.distributed as thd from dgl.graphbolt.datapipes import find_dps, traverse_dps +from utils import rendezvous_port from . import gb_test_utils @@ -89,7 +90,7 @@ def test_gpu_sampling_DataLoader( init_method = ( f"file:///{os.path.join(os.getcwd(), 'dis_tempfile')}" if platform == "win32" - else "tcp://127.0.0.1:12345" + else f"tcp://127.0.0.1:{rendezvous_port()}" ) thd.init_process_group( init_method=init_method, diff --git a/tests/python/pytorch/graphbolt/test_item_sampler.py b/tests/python/pytorch/graphbolt/test_item_sampler.py index 61b352961f28..1b4e95ed208f 100644 --- a/tests/python/pytorch/graphbolt/test_item_sampler.py +++ b/tests/python/pytorch/graphbolt/test_item_sampler.py @@ -13,6 +13,7 @@ import torch import torch.distributed as dist import torch.multiprocessing as mp +from utils import rendezvous_port def test_ItemSampler_minibatcher(): @@ -883,7 +884,7 @@ def distributed_item_sampler_subprocess( init_method = ( f"file:///{os.path.join(os.getcwd(), 'dis_tempfile')}" if platform == "win32" - else "tcp://127.0.0.1:12345" + else f"tcp://127.0.0.1:{rendezvous_port()}" ) dist.init_process_group( backend="gloo", # Use Gloo backend for CPU multiprocessing diff --git a/tests/python/pytorch/nn/test_sparse_emb.py b/tests/python/pytorch/nn/test_sparse_emb.py index 9eaf6f14692d..fb913eb61bcc 100644 --- a/tests/python/pytorch/nn/test_sparse_emb.py +++ b/tests/python/pytorch/nn/test_sparse_emb.py @@ -5,6 +5,7 @@ import backend as F import pytest import torch as th +from utils import rendezvous_port from dgl.nn import NodeEmbedding from dgl.optim import SparseAdam @@ -46,7 +47,7 @@ def check_all_set_all_get_optm_state( def start_sparse_worker(rank, world_size, test, args): print("start sparse worker {}".format(rank)) dist_init_method = "tcp://{master_ip}:{master_port}".format( - master_ip="127.0.0.1", master_port="12345" + master_ip="127.0.0.1", master_port=rendezvous_port() ) backend = "gloo" device = F.ctx() diff --git a/tests/python/pytorch/optim/test_optim.py b/tests/python/pytorch/optim/test_optim.py index 90245ee4a3e5..db4591a96db6 100644 --- a/tests/python/pytorch/optim/test_optim.py +++ b/tests/python/pytorch/optim/test_optim.py @@ -5,6 +5,7 @@ import pytest import torch as th import torch.multiprocessing as mp +from utils import rendezvous_port from dgl.nn import NodeEmbedding from dgl.optim import SparseAdagrad, SparseAdam @@ -191,7 +192,7 @@ def start_sparse_adam_worker( ): print("start sparse worker for adam {}".format(rank)) dist_init_method = "tcp://{master_ip}:{master_port}".format( - master_ip="127.0.0.1", master_port="12345" + master_ip="127.0.0.1", master_port=rendezvous_port() ) if device.type == "cuda": @@ -255,7 +256,7 @@ def start_torch_adam_worker( ): print("start sparse worker for adam {}".format(rank)) dist_init_method = "tcp://{master_ip}:{master_port}".format( - master_ip="127.0.0.1", master_port="12345" + master_ip="127.0.0.1", master_port=rendezvous_port() ) backend = "gloo" @@ -601,7 +602,7 @@ def start_sparse_adam_state_dict_worker( ): print("start sparse worker for adam {}".format(rank)) dist_init_method = "tcp://{master_ip}:{master_port}".format( - master_ip="127.0.0.1", master_port="12345" + master_ip="127.0.0.1", master_port=rendezvous_port() ) device = th.device(f"cuda:{rank}") diff --git a/tests/python/pytorch/test_ffi-stream.py b/tests/python/pytorch/test_ffi-stream.py index 4700a5d3ebf1..79174abf2743 100644 --- a/tests/python/pytorch/test_ffi-stream.py +++ b/tests/python/pytorch/test_ffi-stream.py @@ -160,6 +160,11 @@ def perform_computing(): @unittest.skipIf( F._default_context_str == "cpu", reason="stream only runs on GPU." ) +@unittest.skipIf( + F.is_hip(), + reason="The HIP allocator does not hand g2's block back to g3, so the " + "use-after-free this test provokes does not corrupt the result.", +) def test_record_stream_graph_negative(): cycles_per_ms = _get_cycles_per_ms() diff --git a/tests/scripts/task_unit_test_rocm.sh b/tests/scripts/task_unit_test_rocm.sh index a70a34e8f43b..c8c0af0bf144 100644 --- a/tests/scripts/task_unit_test_rocm.sh +++ b/tests/scripts/task_unit_test_rocm.sh @@ -25,6 +25,42 @@ export DGL_DOWNLOAD_DIR=${PWD}/_download export TF_FORCE_GPU_ALLOW_GROWTH=true unset TORCH_ALLOW_TF32_CUBLAS_OVERRIDE +# torch_shm_manager, which PyTorch launches to share CPU tensors across +# processes, links against librocm-openblas.so.0. When PyTorch is installed from +# the rocm-sdk wheels that library lives inside the package rather than on the +# default loader path. +ROCM_SDK_MATH_LIBS=$(python3 -c "import os, _rocm_sdk_core; print(os.path.join(os.path.dirname(_rocm_sdk_core.__file__), 'lib', 'host-math', 'lib'))" 2>/dev/null) +if [ -d "${ROCM_SDK_MATH_LIBS}" ]; then + export LD_LIBRARY_PATH=${ROCM_SDK_MATH_LIBS}:${LD_LIBRARY_PATH} +fi + +# The distributed tests rendezvous on DGL_TEST_MASTER_PORT and that port plus +# two. Their defaults are often claimed by unrelated daemons when the tests run +# in a container sharing a busy host's network namespace, so pick a free pair. +if [ -z "${DGL_TEST_MASTER_PORT}" ]; then + export DGL_TEST_MASTER_PORT=$(python3 - <<'PY' +import socket + +def available(base): + held = [] + try: + for port in (base, base + 2): + sock = socket.socket() + held.append(sock) + sock.bind(("", port)) + except OSError: + return False + finally: + for sock in held: + sock.close() + return True + +print(next((base for base in range(12345, 12645, 4) if available(base)), 12345)) +PY +) +fi +echo "distributed tests rendezvous on port ${DGL_TEST_MASTER_PORT}" + if [ $2 == "gpu" ] then export CUDA_VISIBLE_DEVICES=0 @@ -36,21 +72,50 @@ echo "pytests running without Logger" python3 -m pip install expecttest +# These files build DataLoaders with worker processes. Forking a pytest process +# that has already accumulated thread pools from earlier tests leaves the +# workers deadlocked, so each file gets a fresh interpreter and is excluded from +# the main backend run. +ISOLATED_TESTS="" +IGNORE_ISOLATED="" +for test_file in tests/python/$DGLBACKEND/graphbolt/test_dataloader.py \ + tests/python/$DGLBACKEND/dataloading/test_dataloader.py; do + if [ -f ${test_file} ]; then + ISOLATED_TESTS="${ISOLATED_TESTS} ${test_file}" + IGNORE_ISOLATED="${IGNORE_ISOLATED} --ignore=${test_file}" + fi +done + +exit_code=0 + if [ ${COVERAGE} == "off" ]; then echo "pytests running without coverage" - python3 -m pytest --junitxml=pytest_dgl_import.xml --durations=100 --disable-warnings tests/python/test_dgl_import.py - python3 -m pytest --junitxml=pytest_common.xml --durations=100 --disable-warnings tests/python/common - python3 -m pytest --junitxml=pytest_backend.xml --durations=100 --disable-warnings tests/python/$DGLBACKEND + python3 -m pytest --junitxml=pytest_dgl_import.xml --durations=100 --disable-warnings tests/python/test_dgl_import.py || exit_code=$? + python3 -m pytest --junitxml=pytest_common.xml --durations=100 --disable-warnings tests/python/common || exit_code=$? + python3 -m pytest --junitxml=pytest_backend.xml --durations=100 --disable-warnings tests/python/$DGLBACKEND ${IGNORE_ISOLATED} || exit_code=$? + + for test_file in ${ISOLATED_TESTS}; do + report_name=$(echo ${test_file} | tr '/' '_' | sed 's/\.py$//') + python3 -m pytest --junitxml=pytest_${report_name}.xml --durations=100 --disable-warnings ${test_file} || exit_code=$? + done elif [ ${COVERAGE} == "on" ]; then echo "pytests running with coverage" python3 -m pip install pytest-cov - python3 -m pytest --cov=dgl --cov-report=lcov:lcov_pytest_import.info --disable-warnings tests/python/test_dgl_import.py - python3 -m pytest --cov=dgl --cov-append --cov-report=lcov:lcov_pytest_common.info --disable-warnings tests/python/common - python3 -m pytest --cov=dgl --cov-append --cov-report=lcov:lcov_pytest_backend.info --disable-warnings tests/python/$DGLBACKEND + python3 -m pytest --cov=dgl --cov-report=lcov:lcov_pytest_import.info --disable-warnings tests/python/test_dgl_import.py || exit_code=$? + python3 -m pytest --cov=dgl --cov-append --cov-report=lcov:lcov_pytest_common.info --disable-warnings tests/python/common || exit_code=$? + python3 -m pytest --cov=dgl --cov-append --cov-report=lcov:lcov_pytest_backend.info --disable-warnings tests/python/$DGLBACKEND ${IGNORE_ISOLATED} || exit_code=$? + + isolated_traces="" + index=0 + for test_file in ${ISOLATED_TESTS}; do + index=$((index + 1)) + python3 -m pytest --cov=dgl --cov-append --cov-report=lcov:lcov_pytest_isolated_${index}.info --disable-warnings ${test_file} || exit_code=$? + isolated_traces="${isolated_traces} -a lcov_pytest_isolated_${index}.info" + done # TODO need to add docs for installing lcov - lcov --add-tracefile lcov_pytest_backend.info -a lcov_pytest_common.info -a lcov_pytest_import.info -o lcov_pytest.info + lcov --add-tracefile lcov_pytest_backend.info -a lcov_pytest_common.info -a lcov_pytest_import.info ${isolated_traces} -o lcov_pytest.info # Show summary of coverage coverage report -m | tee python_coverage_report.txt @@ -59,7 +124,6 @@ else fail "Error: invalid coverage option: ${COVERAGE}" fi -exit_code=$? echo "pytest exited with code: $exit_code" exit $exit_code diff --git a/tests/utils/__init__.py b/tests/utils/__init__.py index 47ad57c45f22..9857718ec404 100644 --- a/tests/utils/__init__.py +++ b/tests/utils/__init__.py @@ -1,7 +1,23 @@ +import os + import backend as F import pytest parametrize_idtype = pytest.mark.parametrize("idtype", [F.int32, F.int64]) + +def rendezvous_port(offset=0): + """Return the TCP port for distributed test rendezvous. + + Set DGL_TEST_MASTER_PORT to move the tests off the default ports when + something else on the machine already listens on them. This matters when + tests run in a container that shares the host network namespace. + + Pass an offset to claim a port distinct from the base one, so that suites + using different offsets stay independent if they ever run concurrently. + """ + return int(os.environ.get("DGL_TEST_MASTER_PORT", 12345)) + offset + + from .checks import * from .graph_cases import get_cases From 98751d662e800c2835405871bcd6ab158b1731c9 Mon Sep 17 00:00:00 2001 From: awelling Date: Fri, 21 Aug 2026 14:07:20 -0500 Subject: [PATCH 2/3] Fix OnDiskDataset failures when loading CSV edges with pandas. Pandas can return read-only numpy views from to_numpy(), which breaks in-place node-type offset updates during heterogeneous graph preprocessing. Co-authored-by: Cursor --- python/dgl/graphbolt/internal/utils.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/python/dgl/graphbolt/internal/utils.py b/python/dgl/graphbolt/internal/utils.py index db4711399ce4..36d1dbf29e44 100644 --- a/python/dgl/graphbolt/internal/utils.py +++ b/python/dgl/graphbolt/internal/utils.py @@ -164,7 +164,10 @@ def read_edges(dataset_dir, edge_fmt, edge_path): os.path.join(dataset_dir, edge_path), names=["src", "dst"], ) - src, dst = edge_data["src"].to_numpy(), edge_data["dst"].to_numpy() + src, dst = ( + edge_data["src"].to_numpy(copy=True), + edge_data["dst"].to_numpy(copy=True), + ) return (src, dst) From dc136ef8781288ab15acdfb4022579c6b0225d55 Mon Sep 17 00:00:00 2001 From: awelling Date: Fri, 21 Aug 2026 14:30:36 -0500 Subject: [PATCH 3/3] Add ROCm Docker build and fix GraphBolt HIP compatibility. Introduce docker/Dockerfile.rocm, sync the legacy CI Dockerfile, apply cuco header patches during GraphBolt deps install, and restore HIP masquerading aliases so GraphBolt builds against recent PyTorch ROCm stacks. Co-authored-by: Cursor --- docker/Dockerfile.ci_gpu_rocm | 17 ++-- docker/Dockerfile.rocm | 83 +++++++++++++++++++ graphbolt/include/graphbolt/async.h | 11 ++- graphbolt/src/cuda/common.h | 25 ++++-- .../src/cuda/extension/gpu_graph_cache.cu | 2 +- .../cuda/extension/unique_and_compact_map.cu | 2 +- script/install_graphbolt_deps.sh | 31 +++++++ 7 files changed, 153 insertions(+), 18 deletions(-) create mode 100644 docker/Dockerfile.rocm diff --git a/docker/Dockerfile.ci_gpu_rocm b/docker/Dockerfile.ci_gpu_rocm index d8758f41a464..5dfa17280cdb 100644 --- a/docker/Dockerfile.ci_gpu_rocm +++ b/docker/Dockerfile.ci_gpu_rocm @@ -1,3 +1,6 @@ +# Deprecated: use docker/Dockerfile.rocm instead. +# This filename is kept for older build scripts. +# # "Copyright Advanced Micro Devices, Inc. # Licensed under the Apache License Version 2.0" @@ -13,16 +16,16 @@ ENV MAX_JOBS=${ARG_MAX_JOBS:-24} # Install basic tools RUN apt-get update \ - && apt install -y clang-format-15 ninja-build tree + && apt install -y clang-format-15 git ninja-build patch tree -#Updating to latest version of Cmake for Ubuntu 22.04 +# Updating to latest version of Cmake for Ubuntu 24.04 RUN apt-get update && apt-get install -y wget build-essential \ && wget -qO- https://cmake.org/files/v3.28/cmake-3.28.3-linux-x86_64.tar.gz \ | tar --strip-components=1 -xz -C /usr/local # Use bash shell for all RUN commands SHELL ["/bin/bash", "--login", "-c"] - + # Copy Sources ENV DGL_SRC_DIR="/src/dgl" RUN mkdir -p ${DGL_SRC_DIR} @@ -35,14 +38,16 @@ ENV GPU_BUILD_TARGETS=${ARG_GPU_BUILD_TARGETS} ENV DGL_BUILD_DIR="${DGL_SRC_DIR}/build" ARG ARG_DGL_ARTIFACTS_DIR="/artifacts" ENV DGL_ARTIFACTS_DIR="${ARG_DGL_ARTIFACTS_DIR}" +ENV DGL_LIBRARY_PATH="${DGL_BUILD_DIR}" +ENV PYTHONPATH="${DGL_SRC_DIR}/tests:${DGL_SRC_DIR}/python" # Install GraphBolt dependencies RUN mkdir -p /tmp/deps && cd /tmp/deps && bash ${DGL_SRC_DIR}/script/install_graphbolt_deps.sh # Configure and build DGL WORKDIR ${DGL_SRC_DIR} -RUN cmake --preset rocm -DCMAKE_HIP_ARCHITECTURES=${ARG_GPU_BUILD_TARGETS} -DGPU_TARGETS=${ARG_GPU_BUILD_TARGETS} -RUN cmake --build ${DGL_BUILD_DIR} --parallel ${MAX_JOBS} +RUN cmake --preset rocm -DCMAKE_HIP_ARCHITECTURES=${ARG_GPU_BUILD_TARGETS} -DGPU_TARGETS=${ARG_GPU_BUILD_TARGETS} +RUN cmake --build ${DGL_BUILD_DIR} --parallel ${MAX_JOBS} WORKDIR ${DGL_SRC_DIR}/python # Make sure we have wheel and build installed @@ -66,7 +71,7 @@ Container manifest: /src/dgl ├── README.md ├── benchmarks # Benchmark scripts -├── ... +├── ... ├── build # C/C++ build directory ├── tests # Test scripts diff --git a/docker/Dockerfile.rocm b/docker/Dockerfile.rocm new file mode 100644 index 000000000000..5f46d5616912 --- /dev/null +++ b/docker/Dockerfile.rocm @@ -0,0 +1,83 @@ +# "Copyright Advanced Micro Devices, Inc. +# Licensed under the Apache License Version 2.0" + +############################################################################# +ARG BASE_IMAGE=rocm/pytorch:rocm7.0.2_ubuntu24.04_py3.12_pytorch_release_2.6.0 +FROM ${BASE_IMAGE} AS dgl_build + +# NOTE: This dockerfile **assumes** that BASE_IMAGE comes with the appropriate +# python executable in the path. + +ARG ARG_MAX_JOBS=24 +ENV MAX_JOBS=${ARG_MAX_JOBS:-24} + +# Install basic tools +RUN apt-get update \ + && apt install -y clang-format-15 git ninja-build patch tree + +# Updating to latest version of Cmake for Ubuntu 24.04 +RUN apt-get update && apt-get install -y wget build-essential \ + && wget -qO- https://cmake.org/files/v3.28/cmake-3.28.3-linux-x86_64.tar.gz \ + | tar --strip-components=1 -xz -C /usr/local + +# Use bash shell for all RUN commands +SHELL ["/bin/bash", "--login", "-c"] + +# Copy Sources +ENV DGL_SRC_DIR="/src/dgl" +RUN mkdir -p ${DGL_SRC_DIR} +COPY . ${DGL_SRC_DIR} + +# Set GPU build targets +ARG ARG_GPU_BUILD_TARGETS="gfx90a,gfx942" +ENV GPU_BUILD_TARGETS=${ARG_GPU_BUILD_TARGETS} + +ENV DGL_BUILD_DIR="${DGL_SRC_DIR}/build" +ARG ARG_DGL_ARTIFACTS_DIR="/artifacts" +ENV DGL_ARTIFACTS_DIR="${ARG_DGL_ARTIFACTS_DIR}" +ENV DGL_LIBRARY_PATH="${DGL_BUILD_DIR}" +ENV PYTHONPATH="${DGL_SRC_DIR}/tests:${DGL_SRC_DIR}/python" + +# Install GraphBolt dependencies +RUN mkdir -p /tmp/deps && cd /tmp/deps && bash ${DGL_SRC_DIR}/script/install_graphbolt_deps.sh + +# Configure and build DGL +WORKDIR ${DGL_SRC_DIR} +RUN cmake --preset rocm -DCMAKE_HIP_ARCHITECTURES=${ARG_GPU_BUILD_TARGETS} -DGPU_TARGETS=${ARG_GPU_BUILD_TARGETS} +RUN cmake --build ${DGL_BUILD_DIR} --parallel ${MAX_JOBS} + +WORKDIR ${DGL_SRC_DIR}/python +# Make sure we have wheel and build installed +RUN python -m pip install build wheel +# Install an editable version of the package so we can make source changes and test them +RUN python -m pip install -e . +# Build the wheel file for packaging (won't update automatically with source changes) +RUN python -m build --wheel + + +# Store the artifacts +RUN mkdir -p ${DGL_ARTIFACTS_DIR} +RUN cp dist/*.whl ${DGL_ARTIFACTS_DIR} + +# Setup for the user +RUN cat <<'EOT' >> ~/.bashrc +echo " +Welcome to the DGL CI container! + +Container manifest: +/src/dgl +├── README.md +├── benchmarks # Benchmark scripts +├── ... +├── build # C/C++ build directory +├── tests # Test scripts + +Run the tests with: + cd /src/dgl/ + bash tests/scripts/task_unit_test_rocm.sh pytorch gpu +" +EOT + +# Set the default command to an interactive bash shell. +WORKDIR ${DGL_SRC_DIR} +CMD ["/bin/bash"] diff --git a/graphbolt/include/graphbolt/async.h b/graphbolt/include/graphbolt/async.h index f774728dcb33..b11dc3319e57 100644 --- a/graphbolt/include/graphbolt/async.h +++ b/graphbolt/include/graphbolt/async.h @@ -43,13 +43,22 @@ #include #include #elif defined(GRAPHBOLT_USE_HIP) -#include +#include #include #include #include #include #include +namespace c10::cuda { +using CUDAStream = c10::hip::HIPStreamMasqueradingAsCUDA; +using CUDAStreamGuard = c10::hip::HIPStreamGuardMasqueradingAsCUDA; + +inline auto getCurrentCUDAStream() { + return c10::hip::getCurrentHIPStreamMasqueradingAsCUDA(); +} +} // namespace c10::cuda + #endif namespace graphbolt { diff --git a/graphbolt/src/cuda/common.h b/graphbolt/src/cuda/common.h index b2f179d167ec..90940b471a16 100644 --- a/graphbolt/src/cuda/common.h +++ b/graphbolt/src/cuda/common.h @@ -26,17 +26,21 @@ // checks if this is defined rather than checking the value. #undef CUDART_VERSION +#include #include #include #include #include -#include #include +#include -using namespace c10::cuda; -using GPUStream_t = c10::cuda::CUDAStream; +using namespace c10::hip; +using GPUStream_t = at::hip::HIPStreamMasqueradingAsCUDA; #define THRUST_BACKEND thrust::hip -#define GET_CURRENT_GPU_STREAM getCurrentCUDAStream +#define GET_CURRENT_GPU_STREAM getCurrentHIPStreamMasqueradingAsCUDA +#define GRAPHBOLT_CACHING_ALLOCATOR c10::hip::HIPCachingAllocator +#define GRAPHBOLT_C10_CHECK C10_HIP_CHECK +#define GRAPHBOLT_C10_KERNEL_LAUNCH_CHECK C10_HIP_KERNEL_LAUNCH_CHECK #else // we're using CUDA @@ -50,6 +54,9 @@ using namespace c10::cuda; using GPUStream_t = at::cuda::CUDAStream; #define THRUST_BACKEND thrust::cuda #define GET_CURRENT_GPU_STREAM getCurrentCUDAStream +#define GRAPHBOLT_CACHING_ALLOCATOR c10::cuda::CUDACachingAllocator +#define GRAPHBOLT_C10_CHECK C10_CUDA_CHECK +#define GRAPHBOLT_C10_KERNEL_LAUNCH_CHECK C10_CUDA_KERNEL_LAUNCH_CHECK #endif // DGL_USE_HIP @@ -97,13 +104,13 @@ struct CUDAWorkspaceAllocator { CUDAWorkspaceAllocator& operator=(const CUDAWorkspaceAllocator&) = default; void operator()(void* ptr) const { - CUDACachingAllocator::raw_delete(ptr); + GRAPHBOLT_CACHING_ALLOCATOR::raw_delete(ptr); } // Required by thrust to satisfy allocator requirements. value_type* allocate(std::ptrdiff_t size) const { return reinterpret_cast( - CUDACachingAllocator::raw_alloc(size * sizeof(value_type))); + GRAPHBOLT_CACHING_ALLOCATOR::raw_alloc(size * sizeof(value_type))); } // Required by thrust to satisfy allocator requirements. @@ -114,7 +121,7 @@ struct CUDAWorkspaceAllocator { std::size_t size) const { return std::unique_ptr( reinterpret_cast( - CUDACachingAllocator::raw_alloc(sizeof(T) * size)), + GRAPHBOLT_CACHING_ALLOCATOR::raw_alloc(sizeof(T) * size)), *this); } }; @@ -142,7 +149,7 @@ inline bool is_zero(dim3 size) { } \ } while (0) -#define CUDA_CALL(func) C10_CUDA_CHECK((func)) +#define CUDA_CALL(func) GRAPHBOLT_C10_CHECK((func)) #define CUDA_KERNEL_CALL(kernel, nblks, nthrs, shmem, ...) \ { \ @@ -150,7 +157,7 @@ inline bool is_zero(dim3 size) { !graphbolt::cuda::is_zero((nthrs))) { \ auto stream = graphbolt::cuda::GetCurrentStream(); \ (kernel)<<<(nblks), (nthrs), (shmem), stream>>>(__VA_ARGS__); \ - C10_CUDA_KERNEL_LAUNCH_CHECK(); \ + GRAPHBOLT_C10_KERNEL_LAUNCH_CHECK(); \ } \ } diff --git a/graphbolt/src/cuda/extension/gpu_graph_cache.cu b/graphbolt/src/cuda/extension/gpu_graph_cache.cu index 1431a0e140b6..f5d07e78b218 100644 --- a/graphbolt/src/cuda/extension/gpu_graph_cache.cu +++ b/graphbolt/src/cuda/extension/gpu_graph_cache.cu @@ -149,7 +149,7 @@ GpuGraphCache::GpuGraphCache( ::cuda::stream_ref{cuda::GetCurrentStream()}}; map_ = new map_t{std::move(map_temp)}; })); - C10_CUDA_KERNEL_LAUNCH_CHECK(); // Check the map constructor's success. + GRAPHBOLT_C10_KERNEL_LAUNCH_CHECK(); // Check the map constructor's success. const auto options = torch::TensorOptions().device(c10::DeviceType::CUDA); TORCH_CHECK(threshold > 0, "Threshold should be a position integer."); threshold_ = threshold; diff --git a/graphbolt/src/cuda/extension/unique_and_compact_map.cu b/graphbolt/src/cuda/extension/unique_and_compact_map.cu index 0638ce83f452..cf47be1f6c92 100644 --- a/graphbolt/src/cuda/extension/unique_and_compact_map.cu +++ b/graphbolt/src/cuda/extension/unique_and_compact_map.cu @@ -193,7 +193,7 @@ UniqueAndCompactBatchedHashMapBased( cuda::CUDAWorkspaceAllocator>{}, ::cuda::stream_ref{stream}, }; - C10_CUDA_KERNEL_LAUNCH_CHECK(); // Check the map constructor's success. + GRAPHBOLT_C10_KERNEL_LAUNCH_CHECK(); // Check the map constructor's success. const dim3 block(BLOCK_SIZE); const dim3 grid( (offsets_ptr[2 * num_batches] + BLOCK_SIZE - 1) / BLOCK_SIZE); diff --git a/script/install_graphbolt_deps.sh b/script/install_graphbolt_deps.sh index 4aa3e145eb32..3ec7d559b1f2 100755 --- a/script/install_graphbolt_deps.sh +++ b/script/install_graphbolt_deps.sh @@ -83,3 +83,34 @@ cd ${DEPS_DIR} # Right now we need to patch the rocPRIM headers to fix the build because these # config headers are missing gfx942 (I've added them manually) run cp ${FILE_SOURCE_DIR}/*.hpp ${INSTALL_PREFIX}/include/rocprim/device/detail/config/. + +# hipCollections installs cuco headers that still need local fixes on recent ROCm +# stacks (rocThrust 5.0, libhipcxx SFINAE). Patch files live under script/patches/. +apply_cuco_patches() { + local patch_dir="${FILE_SOURCE_DIR}/patches" + if [[ ! -d "${patch_dir}" ]]; then + return 0 + fi + + shopt -s nullglob + local patches=("${patch_dir}"/*.patch) + shopt -u nullglob + if [[ ${#patches[@]} -eq 0 ]]; then + return 0 + fi + + for patch_file in "${patches[@]}"; do + echo "Applying ${patch_file} to cuco headers under ${INSTALL_PREFIX}/include" + if $DRY_RUN; then + echo "[dry-run] patch -p1 -d ${INSTALL_PREFIX} < ${patch_file}" + continue + fi + + if ! sed -n '/^--- /,$p' "${patch_file}" | patch -p1 -d "${INSTALL_PREFIX}" --forward --batch; then + echo "Failed to apply ${patch_file}" >&2 + exit 1 + fi + done +} + +apply_cuco_patches