1
2 // Copyright Oliver Kowalke 2017.
3 // Distributed under the Boost Software License, Version 1.0.
4 // (See accompanying file LICENSE_1_0.txt or copy at
5 // http://www.boost.org/LICENSE_1_0.txt)
6
7 #include <chrono>
8 #include <cstdlib>
9 #include <iostream>
10 #include <memory>
11 #include <random>
12 #include <tuple>
13
14 #include <hip/hip_runtime.h>
15
16 #include <boost/assert.hpp>
17 #include <boost/bind.hpp>
18 #include <boost/intrusive_ptr.hpp>
19
20 #include <boost/fiber/all.hpp>
21 #include <boost/fiber/hip/waitfor.hpp>
22
23 __global__
vector_add(hipLaunchParm lp,int * a,int * b,int * c,int size)24 void vector_add(hipLaunchParm lp, int * a, int * b, int * c, int size) {
25 int idx = threadIdx.x + blockIdx.x * blockDim.x;
26 if ( idx < size) {
27 c[idx] = a[idx] + b[idx];
28 }
29 }
30
main()31 int main() {
32 try {
33 bool done = false;
34 boost::fibers::fiber f1([&done]{
35 std::cout << "f1: entered" << std::endl;
36 try {
37 hipStream_t stream;
38 hipStreamCreate( & stream);
39 int size = 1024 * 1024;
40 int full_size = 20 * size;
41 int * host_a, * host_b, * host_c;
42 hipHostMalloc( & host_a, full_size * sizeof( int), hipHostMallocDefault);
43 hipHostMalloc( & host_b, full_size * sizeof( int), hipHostMallocDefault);
44 hipHostMalloc( & host_c, full_size * sizeof( int), hipHostMallocDefault);
45 int * dev_a, * dev_b, * dev_c;
46 hipMalloc( & dev_a, size * sizeof( int) );
47 hipMalloc( & dev_b, size * sizeof( int) );
48 hipMalloc( & dev_c, size * sizeof( int) );
49 std::minstd_rand generator;
50 std::uniform_int_distribution<> distribution(1, 6);
51 for ( int i = 0; i < full_size; ++i) {
52 host_a[i] = distribution( generator);
53 host_b[i] = distribution( generator);
54 }
55 for ( int i = 0; i < full_size; i += size) {
56 hipMemcpyAsync( dev_a, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream);
57 hipMemcpyAsync( dev_b, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream);
58 hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream, dev_a, dev_b, dev_c, size);
59 hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream);
60 }
61 auto result = boost::fibers::hip::waitfor_all( stream);
62 BOOST_ASSERT( stream == std::get< 0 >( result) );
63 BOOST_ASSERT( hipSuccess == std::get< 1 >( result) );
64 std::cout << "f1: GPU computation finished" << std::endl;
65 hipHostFree( host_a);
66 hipHostFree( host_b);
67 hipHostFree( host_c);
68 hipFree( dev_a);
69 hipFree( dev_b);
70 hipFree( dev_c);
71 hipStreamDestroy( stream);
72 done = true;
73 } catch ( std::exception const& ex) {
74 std::cerr << "exception: " << ex.what() << std::endl;
75 }
76 std::cout << "f1: leaving" << std::endl;
77 });
78 boost::fibers::fiber f2([&done]{
79 std::cout << "f2: entered" << std::endl;
80 while ( ! done) {
81 std::cout << "f2: sleeping" << std::endl;
82 boost::this_fiber::sleep_for( std::chrono::milliseconds( 1 ) );
83 }
84 std::cout << "f2: leaving" << std::endl;
85 });
86 f1.join();
87 f2.join();
88 std::cout << "done." << std::endl;
89 return EXIT_SUCCESS;
90 } catch ( std::exception const& e) {
91 std::cerr << "exception: " << e.what() << std::endl;
92 } catch (...) {
93 std::cerr << "unhandled exception" << std::endl;
94 }
95 return EXIT_FAILURE;
96 }
97