hipMemcpyFromSymbol fails with -fgpu-default-stream=per-thread

[torrance]

I'm trying to use hipMemcpyFromSymbol() with the hipcc option -fgpu-default-stream=per-thread enabled, all as part of setting hipFFT callbacks

The API is defined as:

hipMemcpyFromSymbol (
    void *dst, const void *symbolName, size_t sizeBytes,
    size_t offset=0, hipMemcpyKind kind=hipMemcpyDeviceToHost
)

And in my own code I call it as:

hipMemcpyFromSymbol(
    &callback_load_z_hptr,
    HIP_SYMBOL(callback_load_cf64_dptr),
    sizeof(hipfftCallbackLoadZ)
)

This works just fine with -fgpu-default-stream=legacy, and allows me to set callbacks for hipFFT. My tests pass.

However, when compiled with hipcc and -fgpu-default-stream=per-thread, the optional parameters are no longer defined, as per:

https://github.com/ROCm-Developer-Tools/hipamd/blob/4209792929ddf54ba9530813b7879cfdee42df14/include/hip/amd_detail/amd_hip_runtime_pt_api.h#L96-L97

Fine, ok, I'll call it with optional parameters specified:

hipMemcpyFromSymbol(
    &callback_load_z_hptr,
    HIP_SYMBOL(callback_load_cf64_dptr),
    sizeof(hipfftCallbackLoadZ),
    0, hipMemcpyDeviceToHost
)

But now I get the error:

note: candidate function not viable: no known conversion from 'hipfftCallbackLoadZ' (aka 'HIP_vector_type<double, 2> (*)(void *, unsigned long, void *, void *)') to 'const void *' for 2nd argument; take the address of the argument with &
hipError_t hipMemcpyFromSymbol_spt(void* dst, const void* symbol,size_t sizeBytes

It no longer likes the type returned by HIP_SYMBOL anymore, despite this being OK with the legacy stream API. Let me take the address of the HIP_SYMBOL:

hipMemcpyFromSymbol(
    &callback_load_z_hptr,
    & HIP_SYMBOL(callback_load_cf64_dptr),
    sizeof(hipfftCallbackLoadZ),
    0, hipMemcpyDeviceToHost
)

Now this compiles and works with -fgpu-default-stream=per-thread with my tests passing. But the same call now results in failing tests with -fgpu-default-stream=legacy.

My expectation is that HIP will expose a stable API irrespective of the default stream option.

[ppanchad-amd]

@torrance Apologies for the lack of response. Can you please test with latest ROCm 6.1.0 (HIP 6.1)? If resolved, please close ticket. Thanks!

[darren-amd]

Hi @torrance,

Good news, an update was made in ROCm 5.7 that ensures consistency between the spt API declarations and the non-spt API declarations, as seen here: SWDEV-407692.

Please update to the latest version of ROCm (6.2.2) and you should see the changes requested, thanks!