read_launch.cc

#include <hip/hip_runtime.h>

#include<string>
#include<iostream>
#include<fstream>
#include<vector>



hipFunction_t read_shared( std::string path , const char* func_name )
{
  std::ostringstream sstream;
  std::ifstream fin(path, std::ios::binary);
  sstream << fin.rdbuf();
  std::string shared(sstream.str());

  std::cout << "shared object file read back in. size = " << shared.size() << "\n";

  hipModule_t module;
  hipError_t ret;

  ret = hipModuleLoadData(&module, shared.data() );

  if (ret != hipSuccess) {
    std::cerr << "could not load module with hip\n";
    exit(1);
  }

  std::cout << "shared object file loaded as hip module\n";

  hipFunction_t tmp;

  ret = hipModuleGetFunction(&tmp, module, func_name );

  if (ret != hipSuccess) {
    std::cerr << "could not find function in module with hip\n";
    exit(1);
  }

  std::cout << "Got function!\n";
  return tmp;
}




int main()
{
  hipFunction_t func = read_shared( "module.so" , "test" );

  const int N = 128;

  hipDeviceptr_t dev_ptr_dest;
  hipDeviceptr_t dev_ptr_a;

  float* ptr_dest = new float[N];
  float* ptr_a = new float[N];

  for (int i = 0 ; i < N ; ++i )
    {
      ptr_dest[i] = 0.0;
      ptr_a[i] = 2.0;
    }

  hipMalloc(&dev_ptr_dest, sizeof(float)*N);
  hipMalloc(&dev_ptr_a, sizeof(float)*N);

  hipMemcpyHtoD(dev_ptr_dest, ptr_dest, sizeof(float)*N );
  hipMemcpyHtoD(dev_ptr_a, ptr_a, sizeof(float)*N );

  struct {
    int   N;
    hipDeviceptr_t ptr_dest;
    hipDeviceptr_t ptr_a;
  } args{
     N,
     dev_ptr_dest,
     dev_ptr_a,
  };

  auto size = sizeof(args);
  void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
            HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
            HIP_LAUNCH_PARAM_END};


  hipError_t ret = hipModuleLaunchKernel( func ,
                      1, 1, 1,
                      N, 1, 1,
                      0, nullptr, nullptr,
                      config);

  if (ret != hipSuccess)
    {
      std::cerr << "launch failed\n";
      exit(1);
    }
  else
    {
      std::cerr << "launch successful.\n";
    }


  hipMemcpyDtoH( ptr_dest , dev_ptr_dest , sizeof(float)*N );
  hipMemcpyDtoH( ptr_a , dev_ptr_a , sizeof(float)*N );

  std::cout << "dest: ";
  for (int i = 0 ; i < N ; ++i )
    {
      std::cout << ptr_dest[i] << " ";
    }
  std::cout << "\n";

  std::cout << "a: ";
  for (int i = 0 ; i < N ; ++i )
    {
      std::cout << ptr_a[i] << " ";
    }
  std::cout << "\n";

  return 0;
}