Fix namespace used for hip integration, update hip and cuda examples

.gitignore

@@ -5,3 +5,5 @@ doc\html
 cscope.*
 examples/cuda/multiple_streams*
 examples/cuda/single_stream*
+examples/hip/multiple_streams*
+examples/hip/single_stream*

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.]]
 ]
 
 

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).
 
@@ -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
@@ -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.]]
 ]
 
 

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
 

examples/cuda/multiple_streams.cu

@@ -8,15 +8,12 @@
 #include <chrono>
 #include <cstdlib>
 #include <iostream>
-#include <memory>
 #include <random>
 #include <tuple>
 
 #include <cuda_runtime_api.h>
 
 #include <boost/assert.hpp>
-#include <boost/bind.hpp>
-#include <boost/intrusive_ptr.hpp>
 
 #include <boost/fiber/all.hpp>
 #include <boost/fiber/cuda/waitfor.hpp>

examples/cuda/single_stream.cu

@@ -7,15 +7,12 @@
 #include <chrono>
 #include <cstdlib>
 #include <iostream>
-#include <memory>
 #include <random>
 #include <tuple>
 
 #include <cuda_runtime_api.h>
 
 #include <boost/assert.hpp>
-#include <boost/bind.hpp>
-#include <boost/intrusive_ptr.hpp>
 
 #include <boost/fiber/all.hpp>
 #include <boost/fiber/cuda/waitfor.hpp>

examples/hip/Makefile

@@ -1,23 +1,23 @@
-HIP_PATH := /opt/rocm/hip
+HIP_PATH := /opt/rocm
 
 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
 
 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

examples/hip/multiple_streams.cpp → examples/hip/multiple_streams.hip

@@ -7,21 +7,18 @@
 #include <chrono>
 #include <cstdlib>
 #include <iostream>
-#include <memory>
 #include <random>
 #include <tuple>
 
 #include <hip/hip_runtime.h>
 
 #include <boost/assert.hpp>
-#include <boost/bind.hpp>
-#include <boost/intrusive_ptr.hpp>
 
 #include <boost/fiber/all.hpp>
-#include <boost/fiber/hip/waitfor.hpp>
+#include "boost/fiber/hip/waitfor.hpp"
 
 __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];
@@ -62,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);
                     }

examples/hip/single_stream.cpp → examples/hip/single_stream.hip

@@ -7,21 +7,18 @@
 #include <chrono>
 #include <cstdlib>
 #include <iostream>
-#include <memory>
 #include <random>
 #include <tuple>
 
 #include <hip/hip_runtime.h>
 
 #include <boost/assert.hpp>
-#include <boost/bind.hpp>
-#include <boost/intrusive_ptr.hpp>
 
 #include <boost/fiber/all.hpp>
-#include <boost/fiber/hip/waitfor.hpp>
+#include "boost/fiber/hip/waitfor.hpp"
 
 __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];
@@ -55,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);

include/boost/fiber/cuda/waitfor.hpp

@@ -9,14 +9,10 @@
 
 #include <initializer_list>
 #include <mutex>
-#include <iostream>
 #include <set>
 #include <tuple>
 #include <vector>
 
-#include <boost/assert.hpp>
-#include <boost/config.hpp>
-
 #include <cuda_runtime_api.h>
 
 #include <boost/fiber/detail/config.hpp>

include/boost/fiber/hip/waitfor.hpp

@@ -4,19 +4,15 @@
 //    (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 <initializer_list>
 #include <mutex>
-#include <iostream>
 #include <set>
 #include <tuple>
 #include <vector>
 
-#include <boost/assert.hpp>
-#include <boost/config.hpp>
-
 #include <hip/hip_runtime.h>
 
 #include <boost/fiber/detail/config.hpp>
@@ -30,7 +26,7 @@
 
 namespace boost {
 namespace fibers {
-namespace cuda {
+namespace hip {
 namespace detail {
 
 template< typename Rendezvous >
@@ -136,4 +132,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