From ab4b421d9885a2c797b7510653337e556d61ce8f Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 14 Aug 2026 00:46:37 +0200 Subject: [PATCH 1/9] fix namespace used for hip integration --- include/boost/fiber/hip/waitfor.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/include/boost/fiber/hip/waitfor.hpp b/include/boost/fiber/hip/waitfor.hpp index e05099ef..a568e5ed 100644 --- a/include/boost/fiber/hip/waitfor.hpp +++ b/include/boost/fiber/hip/waitfor.hpp @@ -4,8 +4,8 @@ // (See accompanying file LICENSE_1_0.txt or copy at // http://www.boost.org/LICENSE_1_0.txt) -#ifndef BOOST_FIBERS_CUDA_WAITFOR_H -#define BOOST_FIBERS_CUDA_WAITFOR_H +#ifndef BOOST_FIBERS_HIP_WAITFOR_H +#define BOOST_FIBERS_HIP_WAITFOR_H #include #include @@ -30,7 +30,7 @@ namespace boost { namespace fibers { -namespace cuda { +namespace hip { namespace detail { template< typename Rendezvous > @@ -136,4 +136,4 @@ std::vector< std::tuple< hipStream_t, hipError_t > > waitfor_all( hipStream_t st # include BOOST_ABI_SUFFIX #endif -#endif // BOOST_FIBERS_CUDA_WAITFOR_H +#endif // BOOST_FIBERS_HIP_WAITFOR_H From acb6bee5571d60a34168055ebf592faed0078643 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 14 Aug 2026 00:47:50 +0200 Subject: [PATCH 2/9] remove unused includes from cuda, hip headers --- include/boost/fiber/cuda/waitfor.hpp | 4 ---- include/boost/fiber/hip/waitfor.hpp | 4 ---- 2 files changed, 8 deletions(-) diff --git a/include/boost/fiber/cuda/waitfor.hpp b/include/boost/fiber/cuda/waitfor.hpp index e258c9d3..340e0e8d 100644 --- a/include/boost/fiber/cuda/waitfor.hpp +++ b/include/boost/fiber/cuda/waitfor.hpp @@ -9,14 +9,10 @@ #include #include -#include #include #include #include -#include -#include - #include #include diff --git a/include/boost/fiber/hip/waitfor.hpp b/include/boost/fiber/hip/waitfor.hpp index a568e5ed..fb590901 100644 --- a/include/boost/fiber/hip/waitfor.hpp +++ b/include/boost/fiber/hip/waitfor.hpp @@ -9,14 +9,10 @@ #include #include -#include #include #include #include -#include -#include - #include #include From 0c8740f4fb00c610f23d3242971ec40ed5ecb190 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 14 Aug 2026 01:05:50 +0200 Subject: [PATCH 3/9] remove unused includes, libs, silence deprecation warning from bind in hip,cuda examples --- examples/cuda/Makefile | 2 +- examples/cuda/multiple_streams.cu | 3 --- examples/cuda/single_stream.cu | 3 --- examples/hip/Makefile | 2 +- examples/hip/multiple_streams.cpp | 3 --- examples/hip/single_stream.cpp | 3 --- 6 files changed, 2 insertions(+), 14 deletions(-) diff --git a/examples/cuda/Makefile b/examples/cuda/Makefile index 60a3098e..0931bec4 100644 --- a/examples/cuda/Makefile +++ b/examples/cuda/Makefile @@ -5,7 +5,7 @@ NVCC := $(CUDA_PATH)/bin/nvcc CPPFLAGS := -O2 -std=c++11 LDFLAGS := -g -L/usr/local/lib INCLUDES := -I/usr/local/include -I$(CUDA_PATH)/include -LIBRARIES := -lboost_fiber -lboost_context -lboost_system -lboost_filesystem +LIBRARIES := -lboost_fiber -lboost_context all: build diff --git a/examples/cuda/multiple_streams.cu b/examples/cuda/multiple_streams.cu index ec66b567..0645cf26 100644 --- a/examples/cuda/multiple_streams.cu +++ b/examples/cuda/multiple_streams.cu @@ -8,15 +8,12 @@ #include #include #include -#include #include #include #include #include -#include -#include #include #include diff --git a/examples/cuda/single_stream.cu b/examples/cuda/single_stream.cu index 54781811..64bf698f 100644 --- a/examples/cuda/single_stream.cu +++ b/examples/cuda/single_stream.cu @@ -7,15 +7,12 @@ #include #include #include -#include #include #include #include #include -#include -#include #include #include diff --git a/examples/hip/Makefile b/examples/hip/Makefile index d40e8c1a..9c1f67a3 100644 --- a/examples/hip/Makefile +++ b/examples/hip/Makefile @@ -5,7 +5,7 @@ HIPCC := $(HIP_PATH)/bin/hipcc CPPFLAGS := -O2 -std=c++11 LDFLAGS := -L/usr/local/lib INCLUDES := -I/usr/local/include -I$(HIP_PATH)/include -LIBRARIES := -lboost_fiber -lboost_context -lboost_system -lboost_filesystem +LIBRARIES := -lboost_fiber -lboost_context all: build diff --git a/examples/hip/multiple_streams.cpp b/examples/hip/multiple_streams.cpp index 75a7449c..b1773547 100644 --- a/examples/hip/multiple_streams.cpp +++ b/examples/hip/multiple_streams.cpp @@ -7,15 +7,12 @@ #include #include #include -#include #include #include #include #include -#include -#include #include #include diff --git a/examples/hip/single_stream.cpp b/examples/hip/single_stream.cpp index 1959528e..583ab7c9 100644 --- a/examples/hip/single_stream.cpp +++ b/examples/hip/single_stream.cpp @@ -7,15 +7,12 @@ #include #include #include -#include #include #include #include #include -#include -#include #include #include From 2559774ede97745550f9ec7fe963376b03d87494 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 14 Aug 2026 01:38:30 +0200 Subject: [PATCH 4/9] use modern hip in examples --- examples/hip/Makefile | 2 +- examples/hip/multiple_streams.cpp | 6 +++--- examples/hip/single_stream.cpp | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/examples/hip/Makefile b/examples/hip/Makefile index 9c1f67a3..7262e8ce 100644 --- a/examples/hip/Makefile +++ b/examples/hip/Makefile @@ -1,4 +1,4 @@ -HIP_PATH := /opt/rocm/hip +HIP_PATH := /opt/rocm HIPCC := $(HIP_PATH)/bin/hipcc diff --git a/examples/hip/multiple_streams.cpp b/examples/hip/multiple_streams.cpp index b1773547..5308bce7 100644 --- a/examples/hip/multiple_streams.cpp +++ b/examples/hip/multiple_streams.cpp @@ -18,7 +18,7 @@ #include __global__ -void vector_add(hipLaunchParm lp, int * a, int * b, int * c, int size) { +void vector_add(int * a, int * b, int * c, int size) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if ( idx < size) { c[idx] = a[idx] + b[idx]; @@ -59,8 +59,8 @@ int main() { hipMemcpyAsync( dev_a1, host_a + i + size, size * sizeof( int), hipMemcpyHostToDevice, stream1); hipMemcpyAsync( dev_b0, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream0); hipMemcpyAsync( dev_b1, host_b + i + size, size * sizeof( int), hipMemcpyHostToDevice, stream1); - hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream0, dev_a0, dev_b0, dev_c0, size); - hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream1, dev_a1, dev_b1, dev_c1, size); + hipLaunchKernelGGL( vector_add, dim3(size / 256), dim3(256), 0, stream0, dev_a0, dev_b0, dev_c0, size); + hipLaunchKernelGGL( vector_add, dim3(size / 256), dim3(256), 0, stream1, dev_a1, dev_b1, dev_c1, size); hipMemcpyAsync( host_c + i, dev_c0, size * sizeof( int), hipMemcpyDeviceToHost, stream0); hipMemcpyAsync( host_c + i + size, dev_c1, size * sizeof( int), hipMemcpyDeviceToHost, stream1); } diff --git a/examples/hip/single_stream.cpp b/examples/hip/single_stream.cpp index 583ab7c9..955540fd 100644 --- a/examples/hip/single_stream.cpp +++ b/examples/hip/single_stream.cpp @@ -18,7 +18,7 @@ #include __global__ -void vector_add(hipLaunchParm lp, int * a, int * b, int * c, int size) { +void vector_add(int * a, int * b, int * c, int size) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if ( idx < size) { c[idx] = a[idx] + b[idx]; @@ -52,7 +52,7 @@ int main() { for ( int i = 0; i < full_size; i += size) { hipMemcpyAsync( dev_a, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream); hipMemcpyAsync( dev_b, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream); - hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream, dev_a, dev_b, dev_c, size); + hipLaunchKernelGGL( vector_add, dim3(size / 256), dim3(256), 0, stream, dev_a, dev_b, dev_c, size); hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream); } auto result = boost::fibers::hip::waitfor_all( stream); From 6df1e9b68340d1abe27144ab4add47953bafa7a9 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 14 Aug 2026 14:05:11 +0200 Subject: [PATCH 5/9] use hip extension for files with hip kernels --- examples/hip/Makefile | 4 ++-- examples/hip/{multiple_streams.cpp => multiple_streams.hip} | 2 +- examples/hip/{single_stream.cpp => single_stream.hip} | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) rename examples/hip/{multiple_streams.cpp => multiple_streams.hip} (99%) rename examples/hip/{single_stream.cpp => single_stream.hip} (98%) diff --git a/examples/hip/Makefile b/examples/hip/Makefile index 7262e8ce..f59fd362 100644 --- a/examples/hip/Makefile +++ b/examples/hip/Makefile @@ -11,13 +11,13 @@ all: build build: single_stream multiple_streams -single_stream.o:single_stream.cpp +single_stream.o:single_stream.hip $(HIPCC) $(INCLUDES) $(CPPFLAGS) -o $@ -c $< single_stream: single_stream.o $(HIPCC) $(LDFLAGS) -o $@ $+ $(LIBRARIES) -multiple_streams.o:multiple_streams.cpp +multiple_streams.o:multiple_streams.hip $(HIPCC) $(INCLUDES) $(CPPFLAGS) -o $@ -c $< multiple_streams: multiple_streams.o diff --git a/examples/hip/multiple_streams.cpp b/examples/hip/multiple_streams.hip similarity index 99% rename from examples/hip/multiple_streams.cpp rename to examples/hip/multiple_streams.hip index 5308bce7..69ef98b3 100644 --- a/examples/hip/multiple_streams.cpp +++ b/examples/hip/multiple_streams.hip @@ -15,7 +15,7 @@ #include #include -#include +#include "boost/fiber/hip/waitfor.hpp" __global__ void vector_add(int * a, int * b, int * c, int size) { diff --git a/examples/hip/single_stream.cpp b/examples/hip/single_stream.hip similarity index 98% rename from examples/hip/single_stream.cpp rename to examples/hip/single_stream.hip index 955540fd..2db279dd 100644 --- a/examples/hip/single_stream.cpp +++ b/examples/hip/single_stream.hip @@ -15,7 +15,7 @@ #include #include -#include +#include "boost/fiber/hip/waitfor.hpp" __global__ void vector_add(int * a, int * b, int * c, int size) { From 58b790b94e7a3294f81358a20433c4d8431f19ea Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Fri, 14 Aug 2026 14:08:22 +0200 Subject: [PATCH 6/9] ignore hip example artifacts --- .gitignore | 2 ++ 1 file changed, 2 insertions(+) diff --git a/.gitignore b/.gitignore index e8631914..012aebed 100644 --- a/.gitignore +++ b/.gitignore @@ -5,3 +5,5 @@ doc\html cscope.* examples/cuda/multiple_streams* examples/cuda/single_stream* +examples/hip/multiple_streams* +examples/hip/single_stream* From d23ac5a48148bdec0bece60fe115376b8e91f239 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Sun, 16 Aug 2026 01:46:29 +0200 Subject: [PATCH 7/9] update HIP snippet in docs --- doc/hip.qbk | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/hip.qbk b/doc/hip.qbk index b5302c24..84755d26 100644 --- a/doc/hip.qbk +++ b/doc/hip.qbk @@ -57,7 +57,7 @@ streams has finished. for ( int i = 0; i < full_size; i += size) { hipMemcpyAsync( dev_a, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream); hipMemcpyAsync( dev_b, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream); - hipLaunchKernel(kernel, dim3(size / 256), dim3(256), 0, stream, size, dev_a, dev_b, dev_c); + hipLaunchKernelGGL(kernel, dim3(size / 256), dim3(256), 0, stream, size, dev_a, dev_b, dev_c); hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream); } auto result = boost::fibers::hip::waitfor_all( stream); // suspend fiber till HIP stream has finished From 5d8f3a353b996d6e47b9d38f8e61448d2d90f52c Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Sun, 16 Aug 2026 01:47:15 +0200 Subject: [PATCH 8/9] fix format --- doc/hip.qbk | 2 +- examples/hip/multiple_streams.hip | 2 +- examples/hip/single_stream.hip | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/doc/hip.qbk b/doc/hip.qbk index 84755d26..9a02f3f7 100644 --- a/doc/hip.qbk +++ b/doc/hip.qbk @@ -57,7 +57,7 @@ streams has finished. for ( int i = 0; i < full_size; i += size) { hipMemcpyAsync( dev_a, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream); hipMemcpyAsync( dev_b, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream); - hipLaunchKernelGGL(kernel, dim3(size / 256), dim3(256), 0, stream, size, dev_a, dev_b, dev_c); + hipLaunchKernelGGL( kernel, dim3(size / 256), dim3(256), 0, stream, size, dev_a, dev_b, dev_c); hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream); } auto result = boost::fibers::hip::waitfor_all( stream); // suspend fiber till HIP stream has finished diff --git a/examples/hip/multiple_streams.hip b/examples/hip/multiple_streams.hip index 69ef98b3..a0254289 100644 --- a/examples/hip/multiple_streams.hip +++ b/examples/hip/multiple_streams.hip @@ -18,7 +18,7 @@ #include "boost/fiber/hip/waitfor.hpp" __global__ -void vector_add(int * a, int * b, int * c, int size) { +void vector_add( int * a, int * b, int * c, int size) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if ( idx < size) { c[idx] = a[idx] + b[idx]; diff --git a/examples/hip/single_stream.hip b/examples/hip/single_stream.hip index 2db279dd..5a706f40 100644 --- a/examples/hip/single_stream.hip +++ b/examples/hip/single_stream.hip @@ -18,7 +18,7 @@ #include "boost/fiber/hip/waitfor.hpp" __global__ -void vector_add(int * a, int * b, int * c, int size) { +void vector_add( int * a, int * b, int * c, int size) { int idx = threadIdx.x + blockIdx.x * blockDim.x; if ( idx < size) { c[idx] = a[idx] + b[idx]; From 10ac60b849f055fef1802151ea9128c91489cc86 Mon Sep 17 00:00:00 2001 From: Mateusz Jakub Fila Date: Mon, 24 Aug 2026 10:56:32 +0200 Subject: [PATCH 9/9] update links and CUDA/HIP documentation --- doc/cuda.qbk | 6 +++--- doc/hip.qbk | 8 ++++---- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/doc/cuda.qbk b/doc/cuda.qbk index 261ff1e6..1ca061b1 100644 --- a/doc/cuda.qbk +++ b/doc/cuda.qbk @@ -8,7 +8,7 @@ [#cuda] [section:cuda CUDA] -[@http://developer.nvidia.com/cuda-zone/ CUDA (Compute Unified Device Architecture)] is a platform for parallel computing +[@https://developer.nvidia.com/cuda/ CUDA (Compute Unified Device Architecture)] is a platform for parallel computing on NVIDIA GPUs. The application programming interface of CUDA gives access to GPU's instruction set and computation resources (Execution of compute kernels). @@ -102,8 +102,8 @@ streams has finished. }}} [variablelist -[[Effects:] [Suspends active fiber till CUDA stream has finished its operations.]] -[[Returns:] [tuple of stream reference and the CUDA stream status]] +[[Effects:] [Suspends active fiber till CUDA stream or streams have finished their operations.]] +[[Returns:] [For a single stream, returns a tuple containing the stream and its CUDA status. For multiple streams, returns a vector of tuples containing each stream and its CUDA status.]] ] diff --git a/doc/hip.qbk b/doc/hip.qbk index 9a02f3f7..20925b23 100644 --- a/doc/hip.qbk +++ b/doc/hip.qbk @@ -8,8 +8,8 @@ [#hip] [section:hip ROCm/HIP] -[@http://github.com/ROCm-Developer-Tools/HIP/tree/roc-1.6.0/ HIP] is part of the -[@http://rocm.github.io/ ROC (Radeon Open Compute)] platform for parallel computing +[@https://github.com/ROCm-Developer-Tools/HIP/tree/roc-1.6.0/ HIP] is part of the +[@https://www.amd.com/en/products/software/rocm.html ROCm] platform for parallel computing on AMD and NVIDIA GPUs. The application programming interface of HIP gives access to GPU's instruction set and computation resources (Execution of compute kernels). @@ -103,8 +103,8 @@ streams has finished. }}} [variablelist -[[Effects:] [Suspends active fiber till HIP stream has finished its operations.]] -[[Returns:] [tuple of stream reference and the HIP stream status]] +[[Effects:] [Suspends active fiber till HIP stream or streams have finished their operations.]] +[[Returns:] [For a single stream, returns a tuple containing the stream and its status. For multiple streams, returns a vector of tuples containing each stream and its status.]] ]