multiple_streams.cpp 5.0 KB

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