hipPointerGetAttributes fails on malloc() and stack pointers

[ibaned]

There is some incompatibility between cudaPointerGetAttributes and hipPointerGetAttributes. In particular, if the pointer given to hipPointerGetAttributes is the result of calling malloc() or if it is simply the address of a stack object, cudaPointerGetAttributes will run successfully and indicate that these are host pointers, however hipPointerGetAttributes will fail with hipErrorInvalidValue. This is problematic for codes trying to figure out whether memory should be copied using hipMemcpy or memcpy. The deeper irony is that hipMemcpy(a, b, hipMemcpyDefault) seems to work with such pointers still, suggesting that HIP does have a way to identify that those types of pointers are in host space.

Here is some example code that maybe helps illustrate the issue:

  void* from_hipMalloc = nullptr;
  auto hip_err = hipMalloc(&from_hipMalloc, sizeof(double));
  if (hip_err != hipSuccess) {
    fprintf(stderr, "hipMalloc failed\n");
    std::terminate();
  }
  hipPointerAttribute_t attributes;
  hip_err = hipPointerGetAttributes(&attributes, from_hipMalloc);
  if (hip_err != hipSuccess) {
    fprintf(stderr, "failed for from_hipMalloc\n");
  } else {
    fprintf(stderr, "succeeded for from_hipMalloc\n");
  }
  void* from_hipHostMalloc = nullptr;
  hip_err = hipHostMalloc(&from_hipHostMalloc, sizeof(double));
  if (hip_err != hipSuccess) {
    fprintf(stderr, "hipHostMalloc failed\n");
    std::terminate();
  }
  hip_err = hipPointerGetAttributes(&attributes, from_hipHostMalloc);
  if (hip_err != hipSuccess) {
    fprintf(stderr, "failed for from_hipHostMalloc\n");
  } else {
    fprintf(stderr, "succeeded for from_hipHostMalloc\n");
  }
  void* from_malloc = malloc(sizeof(double));
  if (from_malloc == nullptr) {
    fprintf(stderr, "malloc failed\n");
    std::terminate();
  }
  hip_err = hipPointerGetAttributes(&attributes, from_malloc);
  if (hip_err != hipSuccess) {
    fprintf(stderr, "failed for from_malloc\n");
  } else {
    fprintf(stderr, "succeeded for from_malloc\n");
  }
  double stack_variable;
  void* from_stack = &stack_variable;
  hip_err = hipPointerGetAttributes(&attributes, from_stack);
  if (hip_err != hipSuccess) {
    fprintf(stderr, "failed for from_stack\n");
  } else {
    fprintf(stderr, "succeeded for from_stack\n");
  }

[GZGavinZhao]

This is issue has not been fixed as of ROCm 6.0. If you run the command with AMD_LOG_LEVEL=1, you will see:

succeeded for from_hipMalloc
succeeded for from_hipHostMalloc
:1:hip_memory.cpp           :3421: 9218710751 us: [pid:109671 tid:0x7fd72ce9ba80] Cannot get amd_mem_obj for ptr: 0x6dc83c10 

succeeded for from_malloc
:1:hip_memory.cpp           :3421: 9218710760 us: [pid:109671 tid:0x7fd72ce9ba80] Cannot get amd_mem_obj for ptr: 0x11ef71b8 

succeeded for from_stack

Which comes from hipamd/src/hip_memory.cpp line 3424:

  } else {
    attributes->type = hipMemoryTypeUnregistered;
    attributes->devicePointer = nullptr;
    attributes->hostPointer = nullptr;
    attributes->isManaged = false;
    attributes->allocationFlags = 0;
    attributes->device = hipInvalidDeviceId;
    LogPrintfError("Cannot get amd_mem_obj for ptr: 0x%x \n", ptr);
  }

Currently this is causing Blender to crash:

ic| status: 4097
:1:rocdevice.cpp            :3232: 8987584665 us: [pid:108950 tid:0x7fffad5fe680] hsa_amd_pointer_info() failed
[New Thread 0x7ffd73dff680 (LWP 108997)]
[New Thread 0x7ffd735fe680 (LWP 108998)]
[New Thread 0x7ffd725ff680 (LWP 108999)]
[New Thread 0x7ffd715ff680 (LWP 109000)]
[New Thread 0x7ffd705ff680 (LWP 109001)]
ic| status: 4097
[New Thread 0x7ffd6f7ff680 (LWP 109002)]
:1:rocdevice.cpp            :3232: 8987587428 us: [pid:108950 tid:0x7fffad5fe680] hsa_amd_pointer_info() failed
ic| srcMemoryType: 0, dstMemoryType: 2
blender: /XXX/clr/hipamd/src/hip_memory.cpp:2197: hipError_t ihipGetMemcpyParam3DCommand(amd::Command*&, const HIP_MEMCPY3D*, hip::Stream*): Assertion `false && "ShouldNotReachHere()"' failed.
[New Thread 0x7ffd6e9ff680 (LWP 109003)]
[New Thread 0x7ffd6e1fe680 (LWP 109004)]

Thread 38 "blender" received signal SIGABRT, Aborted.
[Switching to Thread 0x7fffad5fe680 (LWP 108993)]
__pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at pthread_kill.c:44
warning: 44	pthread_kill.c: No such file or directory
(gdb) bt
#0  __pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at pthread_kill.c:44
#1  __pthread_kill_internal (signo=6, threadid=<optimized out>) at pthread_kill.c:78
#2  __GI___pthread_kill (threadid=<optimized out>, signo=signo@entry=6) at pthread_kill.c:89
#3  0x00007ffff7a45196 in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26
#4  0x00007ffff7a29835 in __GI_abort () at abort.c:79
#5  0x00007ffff7a29759 in __assert_fail_base (fmt=0x7ffff7bbeb38 "%s%s%s:%u: %s%sAssertion `%s' failed.\n%n", assertion=assertion@entry=0x7fff9a0775c8 "false && \"ShouldNotReachHere()\"", 
    file=file@entry=0x7fff9a076d00 "/XXX/clr/hipamd/src/hip_memory.cpp", line=line@entry=2197, 
    function=function@entry=0x7fff9a077550 "hipError_t ihipGetMemcpyParam3DCommand(amd::Command*&, const HIP_MEMCPY3D*, hip::Stream*)") at assert.c:92
#6  0x00007ffff7a3be86 in __assert_fail (assertion=0x7fff9a0775c8 "false && \"ShouldNotReachHere()\"", file=0x7fff9a076d00 "/XXX/clr/hipamd/src/hip_memory.cpp", line=2197, 
    function=0x7fff9a077550 "hipError_t ihipGetMemcpyParam3DCommand(amd::Command*&, const HIP_MEMCPY3D*, hip::Stream*)") at assert.c:101
#7  0x00007fff99d90199 in ihipGetMemcpyParam3DCommand (command=@0x7fffad5fbf70: 0x7fffad5fbfe0, pCopy=0x7fffad5fc040, stream=0x7fffaecbdd00) at /XXX/clr/hipamd/src/hip_memory.cpp:2197
#8  0x00007fff99d908c7 in ihipMemcpyParam3D (pCopy=0x7fffad5fc040, stream=0x0, isAsync=false) at /XXX/clr/hipamd/src/hip_memory.cpp:2298
#9  0x00007fff99dc9153 in hipDrvMemcpy2DUnaligned (pCopy=0x7fffad5fc850) at /XXX/clr/hipamd/src/hip_memory.cpp:3968
#10 0x00005555592b5ed9 in ccl::HIPDevice::tex_alloc (this=0x7fffa8e49300, mem=...) at /home/build/YPKG/root/blender/build/blender-4.0.0/intern/cycles/device/hip/device_impl.cpp:789
#11 0x00005555592b97ae in ccl::MultiDevice::mem_copy_to (this=0x7fffa30333c0, mem=...) at /home/build/YPKG/root/blender/build/blender-4.0.0/intern/cycles/device/multi/device.cpp:351
#12 0x00005555598c78b4 in ccl::ImageManager::device_load_image (this=0x7fffcb021800, device=0x7fffa30333c0, scene=0x7fffa3350800, slot=2, progress=<optimized out>)
    at /home/build/YPKG/root/blender/build/blender-4.0.0/intern/cycles/scene/image.cpp:820
#13 0x0000555559b04450 in std::function<void ()>::operator()() const (this=0x1a996) at /usr/bin/../lib64/gcc/x86_64-solus-linux/13/../../../../include/c++/13/bits/std_function.h:591
#14 tbb::detail::d2::(anonymous namespace)::task_ptr_or_nullptr<std::function<void ()> const&>(std::function<void ()> const&) (f=...) at /usr/include/tbb/../oneapi/tbb/task_group.h:135
#15 tbb::detail::d1::function_task<std::function<void ()> >::execute(tbb::detail::d1::execution_data&) (this=0x7fffbd1e5100, ed=...) at /usr/include/tbb/../oneapi/tbb/task_group.h:466
#16 0x00007ffff7f84d83 in ?? () from /usr/lib64/libtbb.so.12
#17 0x00007ffff7f86f62 in ?? () from /usr/lib64/libtbb.so.12
#18 0x00007ffff7a9a10a in start_thread (arg=<optimized out>) at pthread_create.c:444
#19 0x00007ffff7b27a8c in clone3 () at ../sysdeps/unix/sysv/linux/x86_64/clone3.S:78

[GZGavinZhao]

Actually, this is part of a known issue in https://github.com/ROCm/clr/blob/rocm-6.0.0/CHANGELOG.md#known-issues-1. hipPointerGetAttributes no longer fails, but any pointer you didn't obtain from HIP (i.e. malloc or stack) will be treated as memory type hipMemoryTypeUnregistered and remains unusable by HIP. Please ignore the Blender problem as that's a separate issue.

[abhimeda]

@ibaned and @GZGavinZhao can we close this issue as resolved?

[GZGavinZhao]

Yes, I think this can be closed because this is a known behavior now documented in the CHANGELOG.

[satyanveshd]

Closing this since the API doesn't fail now but it returns hipMemoryTypeUnregistered for the malloc and stack pointers.