Advertisement
fwinter

read_launch.cc

Jan 23rd, 2021
135
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 2.46 KB | None | 0 0
  1. #include <hip/hip_runtime.h>
  2.  
  3. #include<string>
  4. #include<iostream>
  5. #include<fstream>
  6. #include<vector>
  7.  
  8.  
  9.  
  10. hipFunction_t read_shared( std::string path , const char* func_name )
  11. {
  12.   std::ostringstream sstream;
  13.   std::ifstream fin(path, std::ios::binary);
  14.   sstream << fin.rdbuf();
  15.   std::string shared(sstream.str());
  16.  
  17.   std::cout << "shared object file read back in. size = " << shared.size() << "\n";
  18.  
  19.   hipModule_t module;
  20.   hipError_t ret;
  21.  
  22.   ret = hipModuleLoadData(&module, shared.data() );
  23.  
  24.   if (ret != hipSuccess) {
  25.     std::cerr << "could not load module with hip\n";
  26.     exit(1);
  27.   }
  28.    
  29.   std::cout << "shared object file loaded as hip module\n";
  30.  
  31.   hipFunction_t tmp;
  32.  
  33.   ret = hipModuleGetFunction(&tmp, module, func_name );
  34.      
  35.   if (ret != hipSuccess) {
  36.     std::cerr << "could not find function in module with hip\n";
  37.     exit(1);
  38.   }
  39.  
  40.   std::cout << "Got function!\n";
  41.   return tmp;
  42. }
  43.  
  44.  
  45.  
  46.  
  47. int main()
  48. {
  49.   hipFunction_t func = read_shared( "module.so" , "test" );
  50.  
  51.   const int N = 128;
  52.  
  53.   hipDeviceptr_t dev_ptr_dest;
  54.   hipDeviceptr_t dev_ptr_a;
  55.  
  56.   float* ptr_dest = new float[N];
  57.   float* ptr_a = new float[N];
  58.  
  59.   for (int i = 0 ; i < N ; ++i )
  60.     {
  61.       ptr_dest[i] = 0.0;
  62.       ptr_a[i] = 2.0;
  63.     }
  64.  
  65.   hipMalloc(&dev_ptr_dest, sizeof(float)*N);
  66.   hipMalloc(&dev_ptr_a, sizeof(float)*N);
  67.  
  68.   hipMemcpyHtoD(dev_ptr_dest, ptr_dest, sizeof(float)*N );
  69.   hipMemcpyHtoD(dev_ptr_a, ptr_a, sizeof(float)*N );
  70.  
  71.   struct {
  72.     int   N;
  73.     hipDeviceptr_t ptr_dest;
  74.     hipDeviceptr_t ptr_a;
  75.   } args{    
  76.      N,
  77.      dev_ptr_dest,
  78.      dev_ptr_a,
  79.   };
  80.  
  81.   auto size = sizeof(args);
  82.   void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
  83.             HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
  84.             HIP_LAUNCH_PARAM_END};
  85.  
  86.  
  87.   hipError_t ret = hipModuleLaunchKernel( func ,
  88.                       1, 1, 1,
  89.                       N, 1, 1,
  90.                       0, nullptr, nullptr,
  91.                       config);
  92.  
  93.   if (ret != hipSuccess)
  94.     {
  95.       std::cerr << "launch failed\n";
  96.       exit(1);
  97.     }
  98.   else
  99.     {
  100.       std::cerr << "launch successful.\n";
  101.     }
  102.  
  103.  
  104.   hipMemcpyDtoH( ptr_dest , dev_ptr_dest , sizeof(float)*N );
  105.   hipMemcpyDtoH( ptr_a , dev_ptr_a , sizeof(float)*N );
  106.  
  107.   std::cout << "dest: ";
  108.   for (int i = 0 ; i < N ; ++i )
  109.     {
  110.       std::cout << ptr_dest[i] << " ";
  111.     }
  112.   std::cout << "\n";
  113.  
  114.   std::cout << "a: ";
  115.   for (int i = 0 ; i < N ; ++i )
  116.     {
  117.       std::cout << ptr_a[i] << " ";
  118.     }
  119.   std::cout << "\n";
  120.  
  121.   return 0;
  122. }
  123.  
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement