ceph/src/boost/libs/fiber/examples/hip/multiple_streams.cpp

   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__
  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
  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 stream0, stream1;
  38                     hipStreamCreate( & stream0);
  39                     hipStreamCreate( & stream1);
  40                     int size = 1024 * 1024;
  41                     int full_size = 20 * size;
  42                     int * host_a, * host_b, * host_c;
  43                     hipHostMalloc( & host_a, full_size * sizeof( int), hipHostMallocDefault);
  44                     hipHostMalloc( & host_b, full_size * sizeof( int), hipHostMallocDefault);
  45                     hipHostMalloc( & host_c, full_size * sizeof( int), hipHostMallocDefault);
  46                     int * dev_a0, * dev_b0, * dev_c0;
  47                     int * dev_a1, * dev_b1, * dev_c1;
  48                     hipMalloc( & dev_a0, size * sizeof( int) );
  49                     hipMalloc( & dev_b0, size * sizeof( int) );
  50                     hipMalloc( & dev_c0, size * sizeof( int) );
  51                     hipMalloc( & dev_a1, size * sizeof( int) );
  52                     hipMalloc( & dev_b1, size * sizeof( int) );
  53                     hipMalloc( & dev_c1, size * sizeof( int) );
  54                     std::minstd_rand generator;
  55                     std::uniform_int_distribution<> distribution(1, 6);
  56                     for ( int i = 0; i < full_size; ++i) {
  57                         host_a[i] = distribution( generator);
  58                         host_b[i] = distribution( generator);
  59                     }
  60                     for ( int i = 0; i < full_size; i += 2 * size) {
  61                         hipMemcpyAsync( dev_a0, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream0);
  62                         hipMemcpyAsync( dev_a1, host_a + i + size, size * sizeof( int), hipMemcpyHostToDevice, stream1);
  63                         hipMemcpyAsync( dev_b0, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream0);
  64                         hipMemcpyAsync( dev_b1, host_b + i + size, size * sizeof( int), hipMemcpyHostToDevice, stream1);
  65                         hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream0, dev_a0, dev_b0, dev_c0, size);
  66                         hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream1, dev_a1, dev_b1, dev_c1, size);
  67                         hipMemcpyAsync( host_c + i, dev_c0, size * sizeof( int), hipMemcpyDeviceToHost, stream0);
  68                         hipMemcpyAsync( host_c + i + size, dev_c1, size * sizeof( int), hipMemcpyDeviceToHost, stream1);
  69                     }
  70                     auto results = boost::fibers::hip::waitfor_all( stream0, stream1);
  71                     for ( auto & result : results) {
  72                         BOOST_ASSERT( stream0 == std::get< 0 >( result) || stream1 == std::get< 0 >( result) );
  73                         BOOST_ASSERT( hipSuccess == std::get< 1 >( result) );
  74                     }
  75                     std::cout << "f1: GPU computation finished" << std::endl;
  76                     hipHostFree( host_a);
  77                     hipHostFree( host_b);
  78                     hipHostFree( host_c);
  79                     hipFree( dev_a0);
  80                     hipFree( dev_b0);
  81                     hipFree( dev_c0);
  82                     hipFree( dev_a1);
  83                     hipFree( dev_b1);
  84                     hipFree( dev_c1);
  85                     hipStreamDestroy( stream0);
  86                     hipStreamDestroy( stream1);
  87                     done = true;
  88                 } catch ( std::exception const& ex) {
  89                     std::cerr << "exception: " << ex.what() << std::endl;
  90                 }
  91                 std::cout << "f1: leaving" << std::endl;
  92         });
  93         boost::fibers::fiber f2([&done]{
  94             std::cout << "f2: entered" << std::endl;
  95             while ( ! done) {
  96                 std::cout << "f2: sleeping" << std::endl;
  97                 boost::this_fiber::sleep_for( std::chrono::milliseconds( 1 ) );
  98             }
  99             std::cout << "f2: leaving" << std::endl;
 100         });
 101         f1.join();
 102         f2.join();
 103         std::cout << "done." << std::endl;
 104         return EXIT_SUCCESS;
 105     } catch ( std::exception const& e) {
 106         std::cerr << "exception: " << e.what() << std::endl;
 107     } catch (...) {
 108         std::cerr << "unhandled exception" << std::endl;
 109     }
 110         return EXIT_FAILURE;
 111 }