townforge/fiber
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

HIP binding

Browse Source 
- waitfor stream(s) to complete while fiber is suspended
This commit is contained in:
Oliver Kowalke 2017-11-02 19:55:48 +01:00
parent 07b1d88be6
commit 2401e6b70c
8 changed files with 373 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

29
examples/hip/Makefile Normal file
View File

@ -0,0 +1,29 @@
HIP_PATH := /opt/rocm/hip
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
all: build
build: single_stream multiple_streams
single_stream.o:single_stream.cpp
$(HIPCC) $(INCLUDES) $(CPPFLAGS) -o $@ -c $<
single_stream: single_stream.o
$(HIPCC) $(LDFLAGS) -o $@ $+ $(LIBRARIES)
multiple_streams.o:multiple_streams.cpp
$(HIPCC) $(INCLUDES) $(CPPFLAGS) -o $@ -c $<
multiple_streams: multiple_streams.o
$(HIPCC) $(LDFLAGS) -o $@ $+ $(LIBRARIES)
clean:
rm -f single_stream single_stream.o multiple_streams multiple_streams.o
clobber: clean

BIN
examples/hip/multiple_streams Executable file
View File

Binary file not shown.

110
examples/hip/multiple_streams.cpp Normal file
View File

@ -0,0 +1,110 @@
#include "hip/hip_runtime.h"
#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>
__global__
void kernel( hipLaunchParm lp, int size, int * a, int * b, int * c) {
int idx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
if ( idx < size) {
int idx1 = (idx + 1) % 256;
int idx2 = (idx + 2) % 256;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / 2;
}
}
int main() {
try {
bool done = false;
boost::fibers::fiber f1( [&done]{
std::cout << "f1: entered" << std::endl;
try {
hipStream_t stream0, stream1;
hipStreamCreate( & stream0);
hipStreamCreate( & stream1);
int size = 1024 * 1024;
int full_size = 20 * size;
int * host_a, * host_b, * host_c;
hipHostMalloc( & host_a, full_size * sizeof( int), hipHostMallocDefault);
hipHostMalloc( & host_b, full_size * sizeof( int), hipHostMallocDefault);
hipHostMalloc( & host_c, full_size * sizeof( int), hipHostMallocDefault);
int * dev_a0, * dev_b0, * dev_c0;
int * dev_a1, * dev_b1, * dev_c1;
hipMalloc( & dev_a0, size * sizeof( int) );
hipMalloc( & dev_b0, size * sizeof( int) );
hipMalloc( & dev_c0, size * sizeof( int) );
hipMalloc( & dev_a1, size * sizeof( int) );
hipMalloc( & dev_b1, size * sizeof( int) );
hipMalloc( & dev_c1, size * sizeof( int) );
std::minstd_rand generator;
std::uniform_int_distribution<> distribution(1, 6);
for ( int i = 0; i < full_size; ++i) {
host_a[i] = distribution( generator);
host_b[i] = distribution( generator);
}
for ( int i = 0; i < full_size; i += 2 * size) {
hipMemcpyAsync( dev_a0, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream0);
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(kernel, dim3(size / 256), dim3(256), 0, stream0 , size, dev_a0, dev_b0, dev_c0);
hipLaunchKernel(kernel, dim3(size / 256), dim3(256), 0, stream1 , size, dev_a1, dev_b1, dev_c1);
hipMemcpyAsync( host_c + i, dev_c0, size * sizeof( int), hipMemcpyDeviceToHost, stream0);
hipMemcpyAsync( host_c + i + size, dev_c1, size * sizeof( int), hipMemcpyDeviceToHost, stream1);
}
auto results = boost::fibers::hip::waitfor_all( stream0, stream1);
for ( auto & result : results) {
BOOST_ASSERT( stream0 == std::get< 0 >( result) || stream1 == std::get< 0 >( result) );
BOOST_ASSERT( hipSuccess == std::get< 1 >( result) );
}
std::cout << "f1: GPU computation finished" << std::endl;
hipHostFree( host_a);
hipHostFree( host_b);
hipHostFree( host_c);
hipFree( dev_a0);
hipFree( dev_b0);
hipFree( dev_c0);
hipFree( dev_a1);
hipFree( dev_b1);
hipFree( dev_c1);
hipStreamDestroy( stream0);
hipStreamDestroy( stream1);
done = true;
} catch ( std::exception const& ex) {
std::cerr << "exception: " << ex.what() << std::endl;
}
std::cout << "f1: leaving" << std::endl;
});
boost::fibers::fiber f2([&done]{
std::cout << "f2: entered" << std::endl;
while ( ! done) {
std::cout << "f2: sleeping" << std::endl;
boost::this_fiber::sleep_for( std::chrono::milliseconds( 1 ) );
}
std::cout << "f2: leaving" << std::endl;
});
f1.join();
f2.join();
std::cout << "done." << std::endl;
return EXIT_SUCCESS;
} catch ( std::exception const& e) {
std::cerr << "exception: " << e.what() << std::endl;
} catch (...) {
std::cerr << "unhandled exception" << std::endl;
}
return EXIT_FAILURE;
}

BIN
examples/hip/multiple_streams.o Normal file
View File

Binary file not shown.

BIN
examples/hip/single_stream Executable file
View File

Binary file not shown.

95
examples/hip/single_stream.cpp Normal file
View File

@ -0,0 +1,95 @@
#include "hip/hip_runtime.h"
#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>
__global__
void kernel( hipLaunchParm lp, int size, int * a, int * b, int * c) {
int idx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
if ( idx < size) {
int idx1 = (idx + 1) % 256;
int idx2 = (idx + 2) % 256;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / 2;
}
}
int main() {
try {
bool done = false;
boost::fibers::fiber f1([&done]{
std::cout << "f1: entered" << std::endl;
try {
hipStream_t stream;
hipStreamCreate( & stream);
int size = 1024 * 1024;
int full_size = 20 * size;
int * host_a, * host_b, * host_c;
hipHostMalloc( & host_a, full_size * sizeof( int), hipHostMallocDefault);
hipHostMalloc( & host_b, full_size * sizeof( int), hipHostMallocDefault);
hipHostMalloc( & host_c, full_size * sizeof( int), hipHostMallocDefault);
int * dev_a, * dev_b, * dev_c;
hipMalloc( & dev_a, size * sizeof( int) );
hipMalloc( & dev_b, size * sizeof( int) );
hipMalloc( & dev_c, size * sizeof( int) );
std::minstd_rand generator;
std::uniform_int_distribution<> distribution(1, 6);
for ( int i = 0; i < full_size; ++i) {
host_a[i] = distribution( generator);
host_b[i] = distribution( generator);
}
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);
hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream);
}
auto result = boost::fibers::hip::waitfor_all( stream);
BOOST_ASSERT( stream == std::get< 0 >( result) );
BOOST_ASSERT( hipSuccess == std::get< 1 >( result) );
std::cout << "f1: GPU computation finished" << std::endl;
hipHostFree( host_a);
hipHostFree( host_b);
hipHostFree( host_c);
hipFree( dev_a);
hipFree( dev_b);
hipFree( dev_c);
hipStreamDestroy( stream);
done = true;
} catch ( std::exception const& ex) {
std::cerr << "exception: " << ex.what() << std::endl;
}
std::cout << "f1: leaving" << std::endl;
});
boost::fibers::fiber f2([&done]{
std::cout << "f2: entered" << std::endl;
while ( ! done) {
std::cout << "f2: sleeping" << std::endl;
boost::this_fiber::sleep_for( std::chrono::milliseconds( 1 ) );
}
std::cout << "f2: leaving" << std::endl;
});
f1.join();
f2.join();
std::cout << "done." << std::endl;
return EXIT_SUCCESS;
} catch ( std::exception const& e) {
std::cerr << "exception: " << e.what() << std::endl;
} catch (...) {
std::cerr << "unhandled exception" << std::endl;
}
return EXIT_FAILURE;
}

BIN
examples/hip/single_stream.o Normal file
View File

Binary file not shown.

139
include/boost/fiber/hip/waitfor.hpp Normal file
View File

@ -0,0 +1,139 @@
// Copyright Oliver Kowalke 2017.
// Distributed under the Boost Software License, Version 1.0.
// (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
#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>
#include <boost/fiber/detail/is_all_same.hpp>
#include <boost/fiber/condition_variable.hpp>
#include <boost/fiber/mutex.hpp>
#ifdef BOOST_HAS_ABI_HEADERS
# include BOOST_ABI_PREFIX
#endif
namespace boost {
namespace fibers {
namespace cuda {
namespace detail {
template< typename Rendezvous >
static void trampoline( hipStream_t st, hipError_t status, void * vp) {
Rendezvous * data = static_cast< Rendezvous * >( vp);
data->notify( st, status);
}
class single_stream_rendezvous {
public:
single_stream_rendezvous( hipStream_t st) {
unsigned int flags = 0;
hipError_t status = ::hipStreamAddCallback( st, trampoline< single_stream_rendezvous >, this, flags);
if ( hipSuccess != status) {
st_ = st;
status_ = status;
done_ = true;
}
}
void notify( hipStream_t st, hipError_t status) noexcept {
std::unique_lock< mutex > lk{ mtx_ };
st_ = st;
status_ = status;
done_ = true;
lk.unlock();
cv_.notify_one();
}
std::tuple< hipStream_t, hipError_t > wait() {
std::unique_lock< mutex > lk{ mtx_ };
cv_.wait( lk, [this]{ return done_; });
return std::make_tuple( st_, status_);
}
private:
mutex mtx_{};
condition_variable cv_{};
hipStream_t st_{};
hipError_t status_{ hipErrorUnknown };
bool done_{ false };
};
class many_streams_rendezvous {
public:
many_streams_rendezvous( std::initializer_list< hipStream_t > l) :
stx_{ l } {
results_.reserve( stx_.size() );
for ( hipStream_t st : stx_) {
unsigned int flags = 0;
hipError_t status = ::hipStreamAddCallback( st, trampoline< many_streams_rendezvous >, this, flags);
if ( hipSuccess != status) {
std::unique_lock< mutex > lk{ mtx_ };
stx_.erase( st);
results_.push_back( std::make_tuple( st, status) );
}
}
}
void notify( hipStream_t st, hipError_t status) noexcept {
std::unique_lock< mutex > lk{ mtx_ };
stx_.erase( st);
results_.push_back( std::make_tuple( st, status) );
if ( stx_.empty() ) {
lk.unlock();
cv_.notify_one();
}
}
std::vector< std::tuple< hipStream_t, hipError_t > > wait() {
std::unique_lock< mutex > lk{ mtx_ };
cv_.wait( lk, [this]{ return stx_.empty(); });
return results_;
}
private:
mutex mtx_{};
condition_variable cv_{};
std::set< hipStream_t > stx_;
std::vector< std::tuple< hipStream_t, hipError_t > > results_;
};
}
void waitfor_all();
inline
std::tuple< hipStream_t, hipError_t > waitfor_all( hipStream_t st) {
detail::single_stream_rendezvous rendezvous( st);
return rendezvous.wait();
}
template< typename ... STP >
std::vector< std::tuple< hipStream_t, hipError_t > > waitfor_all( hipStream_t st0, STP ... stx) {
static_assert( boost::fibers::detail::is_all_same< hipStream_t, STP ...>::value, "all arguments must be of type `CUstream*`.");
detail::many_streams_rendezvous rendezvous{ st0, stx ... };
return rendezvous.wait();
}
}}}
#ifdef BOOST_HAS_ABI_HEADERS
# include BOOST_ABI_SUFFIX
#endif
#endif // BOOST_FIBERS_CUDA_WAITFOR_H