Possiblity to import HSACO files during runtime like cuModuleLoadData(&hmod, hsaco)?

[ghostplant]

With the extractkernel tool, I can get the HSACO code from the executable program and saving it as a binary file (e.g. example.hsaco).

Are there way to load this binary file into AMD devices at HIP runtime and launch it, which is like what cuModuleLoadData() does:

CUmodule hmod;
cuModuleLoadData(&hmod, p_code_data);
..
cuLaunchKernel(..);

[dragontamer]

I've never done this before. But this blog post seems to be doing what you want?

https://gpuopen.com/rocm-with-harmony-combining-opencl-hcc-hsa-in-a-single-program/

EDIT: In particular, this .png file seems to have all the major hsa functions you need to call to do what you're asking for:

[ghostplant]

@dragontamer Great information! I will take a quick look and see whether it works for my purpose.

[ghostplant]

@dragontamer I found that hipModuleLoadData() has already been available and it worked well, while aql-based management is too much complex. Thanks!

[ghostplant]

Recently, I found that hipModuleLoadData() only works for HSACO codes with no global variables used, which would otherwise fail and report:

// Just after calling hipModuleLoadData(&hmod, hsaco_blob_data)
terminate called after throwing an instance of 'std::runtime_error'
  what():  Global symbol: K is undefined.

hipModuleLoadData() is expected to handle global variables import automatically. How to handle this case?