Does the hipimpl::hip_init() affect get_kernargs_size_align?

[yulingao]

I found that in hip version3.3, if hiplaunchkernelggl doesn't call hipimpl::hip_init(), it will result in not finding the arguments to the defined kernel function when get_kernargs_size_align, which ultimately results in the error "Missing metadata for global function: ***".

After reading the source code, there doesn't seem to be a direct correlation between these two functions. What would cause this to happen?

If I want to be able to run get_kernargs_size_align normally without calling hipimpl::hip_init() instead of the whole hiplaunchkernelggl, what can I do?

[cjatin]

3.x is now almost 4 year old release.

Can you upgrade it to 5.6+ and check.

Also can you provide a bit more data here like: your gpu name/architecture, a sample code that reproduces this issue, compiler command you used to compile.

[yulingao]

3.x is now almost 4 year old release.

Can you upgrade it to 5.6+ and check.

Also can you provide a bit more data here like: your gpu name/architecture, a sample code that reproduces this issue, compiler command you used to compile.

Sorry, only version 3.3 and 2.9 are available on the hpc I'm using.

If I comment out hip_impl::hip_init() in the hipLaunchKernelGGL of functional_grid_launch.hpp, then ps.get_kernargs_size_align(reinterpret_caststd::uintptr_t(kernel)) in make_kernarg() will report an error. If I don't comment it out, everything is fine. But now, I have to comment it out.

template<typename... Args, typename F = void (*)(Args...)> inline void hipLaunchKernelGGL(F kernel, const dim3 &numBlocks, const dim3 &dimBlocks, std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) { // hip_impl::hip_init(); auto kernarg = hip_impl::make_kernarg(kernel, std::tuple<Args...>{std::move(args)...}); std::size_t kernarg_size = kernarg.size(); ... }

After my analysis, the problem is in the read_kernarg_metadata() in program_state.inl. In my opinion, hip_init should be unnecessary to read the arguments to the kernel function. But In hip 2.9, things work well when commenting out the hip_init(). Can you help me to analyse why the above happens.

[cjatin]

in 3.x hip_impl::hip_init() was internal hip init (not to be confused with hipInit()). It is required before you query kernel arguments because it initializes visible devices which is required by program_state. You can not comment that line out, it is required.

[yulingao]

in 3.x hip_impl::hip_init() was internal hip init (not to be confused with hipInit()). It is required before you query kernel arguments because it initializes visible devices which is required by program_state. You can not comment that line out, it is required.

Well, that's a real pity. I wanna be able to run get_kernargs_size_align without calling hipimpl::hip_init(). Looks like there's no way to implement my features.

[b-sumner]

Hi @yulingao can you comment on your features? kernarg is not a concept of HIP, its an implementation detail. So it's hard to understand why you care about kernarg if you are writing HIP programs.

[yulingao]

@b-sumner Apologize for the delay in my response. What I would like to do is intercept the hiplaunchkernelGGL function in order to achieve remote execution of the HIP program.

[b-sumner]

Hi @yulingao thanks for the information. I can think of a few ways to do that, but only with some pretty significant restrictions. For example, many HIP (and Cuda) programs allow device functions to directly access host memory or the memory of other GPUs. Are your applications of interest not using such features?

[ppanchad-amd]

@yulingao Do you still assistance with this ticket? If not, please close. Thanks!