Mi250x: pointer is unknown by backend (error code = HIP:1)

[ravil-mobile]

Hi all,

I am running my application on multiple nodes of LUMI (4 Mi250x per node). The program terminates with a bunch of hipSYCL errors i.e.,

[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
[hipSYCL Error] from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
============== hipSYCL error report ============== 
hipSYCL has caught the following undhandled asynchronous errors: 

   0. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
   1. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
   2. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
   3. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
The application will now be terminated.
============== hipSYCL error report ============== 
hipSYCL has caught the following undhandled asynchronous errors: 

   0. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
   1. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
   2. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
   3. from /scratch/project_465000427/download/sycl-2/hipsycl/src/runtime/hip/hip_allocator.cpp:131 @ query_pointer(): hip_allocator: query_pointer(): pointer is unknown by backend (error code = HIP:1)
The application will now be terminated.

The problem occurs only on a multi-node setup which is weird i.e., no issues while running the up on single node with 8 MPI processes.

my current openSYCL configuration

$ opensycl-info
=================Backend information===================
Loaded backend 0: HIP
Found device:
Loaded backend 1: OpenMP
Found device: hipSYCL OpenMP host device

=================Device information===================
***************** Devices for backend HIP *****************
Device 0:
General device information:
Name:
Backend: HIP
Vendor: AMD
Arch: gfx90a:sramecc+:xnack+
Driver version: 50422804
Is CPU: 0
Is GPU: 1
Default executor information:
Is in-order queue: 1
Is out-of-order queue: 0
Is task graph: 0
Device support queries:
images: 0
error_correction: 0
host_unified_memory: 0
little_endian: 1
global_mem_cache: 1
global_mem_cache_read_only: 0
global_mem_cache_read_write: 1
emulated_local_memory: 0
sub_group_independent_forward_progress: 1
usm_device_allocations: 1
usm_host_allocations: 1
usm_atomic_host_allocations: 0
usm_shared_allocations: 1
usm_atomic_shared_allocations: 0
usm_system_allocations: 0
execution_timestamps: 1
sscp_kernels: 0
Device properties:
max_compute_units: 110
max_global_size0: 2199023254528
max_global_size1: 2199023254528
max_global_size2: 2199023254528
max_group_size: 1024
max_num_sub_groups: 16
preferred_vector_width_char: 4
preferred_vector_width_double: 1
preferred_vector_width_float: 1
preferred_vector_width_half: 2
preferred_vector_width_int: 1
preferred_vector_width_long: 1
preferred_vector_width_short: 2
native_vector_width_char: 4
native_vector_width_double: 1
native_vector_width_float: 1
native_vector_width_half: 2
native_vector_width_int: 1
native_vector_width_long: 1
native_vector_width_short: 2
max_clock_speed: 1700
max_malloc_size: 68702699520
address_bits: 64
max_read_image_args: 0
max_write_image_args: 0
image2d_max_width: 0
image2d_max_height: 0
image3d_max_width: 0
image3d_max_height: 0
image3d_max_depth: 0
image_max_buffer_size: 0
image_max_array_size: 0
max_samplers: 0
max_parameter_size: 18446744073709551615
mem_base_addr_align: 8
global_mem_cache_line_size: 128
global_mem_cache_size: 8388608
global_mem_size: 68702699520
max_constant_buffer_size: 2147483647
max_constant_args: 18446744073709551615
local_mem_size: 65536
printf_buffer_size: 18446744073709551615
partition_max_sub_devices: 0
vendor_id: 1022
sub_group_sizes: 64

***************** Devices for backend OpenMP *****************
Device 0:
General device information:
Name: hipSYCL OpenMP host device
Backend: OpenMP
Vendor: the hipSYCL project
Arch:
Driver version: 1.2
Is CPU: 1
Is GPU: 0
Default executor information:
Is in-order queue: 1
Is out-of-order queue: 0
Is task graph: 0
Device support queries:
images: 0
error_correction: 0
host_unified_memory: 1
little_endian: 1
global_mem_cache: 1
global_mem_cache_read_only: 0
global_mem_cache_read_write: 1
emulated_local_memory: 1
sub_group_independent_forward_progress: 0
usm_device_allocations: 1
usm_host_allocations: 1
usm_atomic_host_allocations: 1
usm_shared_allocations: 1
usm_atomic_shared_allocations: 1
usm_system_allocations: 1
execution_timestamps: 1
sscp_kernels: 0
Device properties:
max_compute_units: 8
max_global_size0: 18446744073709551615
max_global_size1: 18446744073709551615
max_global_size2: 18446744073709551615
max_group_size: 1024
max_num_sub_groups: 18446744073709551615
preferred_vector_width_char: 4
preferred_vector_width_double: 1
preferred_vector_width_float: 1
preferred_vector_width_half: 2
preferred_vector_width_int: 1
preferred_vector_width_long: 1
preferred_vector_width_short: 2
native_vector_width_char: 4
native_vector_width_double: 1
native_vector_width_float: 1
native_vector_width_half: 2
native_vector_width_int: 1
native_vector_width_long: 1
native_vector_width_short: 2
max_clock_speed: 0
max_malloc_size: 18446744073709551615
address_bits: 64
max_read_image_args: 0
max_write_image_args: 0
image2d_max_width: 0
image2d_max_height: 0
image3d_max_width: 0
image3d_max_height: 0
image3d_max_depth: 0
image_max_buffer_size: 0
image_max_array_size: 0
max_samplers: 0
max_parameter_size: 18446744073709551615
mem_base_addr_align: 8
global_mem_cache_line_size: 64
global_mem_cache_size: 1
global_mem_size: 18446744073709551615
max_constant_buffer_size: 18446744073709551615
max_constant_args: 18446744073709551615
local_mem_size: 18446744073709551615
printf_buffer_size: 18446744073709551615
partition_max_sub_devices: 0
vendor_id: 18446744073709551615
sub_group_sizes: 1

openSYCL version

syclcc [Open SYCL compilation driver], Copyright (C) 2018-2023 Aksel Alpay and the Open SYCL project
  Open SYCL version: 0.9.4
  Installation root: /pfs/lustrep4/projappl/project_465000427/usr/hipsycl-0.9.3-rocm-5.4
  Plugin LLVM version: 15, can accelerate CPU: False
  Available runtime backends:
     librt-backend-hip.so
     librt-backend-omp.so


Full configuration [can be overriden using environment variables or command line arguments]:
    version-major: 0
    version-minor: 9
    version-patch: 4
    plugin-llvm-version-major: 15
    plugin-with-cpu-acceleration: false
    default-clang: /pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3/llvm/bin/clang++
    default-nvcxx: (unconfigured)
    default-platform: rocm
    default-cuda-path: (unconfigured)
    default-gpu-arch: gfx90a
    default-cpu-cxx: /opt/cray/pe/craype/2.7.19/bin/CC
    default-rocm-path: /pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3
    default-use-bootstrap-mode: false
    default-is-dryrun: false
    default-use-accelerated-cpu: false
    default-clang-include-path: /pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3/llvm/lib/clang/15.0.0/include/..
    default-sequential-link-line: -L/project/project_465000427/usr/hipsycl-0.9.3-rocm-5.4/lib -lboost_context -lboost_fiber -Wl,-rpath=/project/project_465000427/usr/hipsycl-0.9.3-rocm-5.4/lib
    default-sequential-cxx-flags: -I/project/project_465000427/usr/hipsycl-0.9.3-rocm-5.4/include -D_ENABLE_EXTENDED_ALIGNED_STORAGE -I/opt/cray/pe/mpich/8.1.23/ofi/amd/5.0/include -I/opt/cray/pe/libsci/22.12.1.1/AMD/4.0/x86_64/include -I/opt/cray/pe/dsmml/0.2.2/dsmml//include -I/opt/cray/xpmem/2.5.2-2.4_3.20__gd0f7936.shasta/include
    default-omp-link-line: -L/project/project_465000427/usr/hipsycl-0.9.3-rocm-5.4/lib -lboost_context -lboost_fiber -Wl,-rpath=/project/project_465000427/usr/hipsycl-0.9.3-rocm-5.4/lib -fopenmp
    default-omp-cxx-flags: -I/project/project_465000427/usr/hipsycl-0.9.3-rocm-5.4/include -fopenmp -D_ENABLE_EXTENDED_ALIGNED_STORAGE
    default-rocm-link-line: -Wl,-rpath=$HIPSYCL_ROCM_PATH/lib -Wl,-rpath=$HIPSYCL_ROCM_PATH/hip/lib -L/pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3/lib -L/pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3/hip/lib -lamdhip64
    default-rocm-cxx-flags: -isystem $HIPSYCL_PATH/include/hipSYCL/std/hiplike -isystem /pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3/llvm/lib/clang/15.0.0/include/.. -U__FLOAT128__ -U__SIZEOF_FLOAT128__ -I$HIPSYCL_ROCM_PATH/include -I$HIPSYCL_ROCM_PATH/include --rocm-device-lib-path=$HIPSYCL_ROCM_PATH/amdgcn/bitcode --rocm-path=$HIPSYCL_ROCM_PATH -fhip-new-launch-api -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -D__HIP_ROCclr__
    default-cuda-link-line: -Wl,-rpath=$HIPSYCL_CUDA_LIB_PATH -L$HIPSYCL_CUDA_LIB_PATH -lcudart
    default-cuda-cxx-flags: -U__FLOAT128__ -U__SIZEOF_FLOAT128__ -isystem $HIPSYCL_PATH/include/hipSYCL/std/hiplike
    default-is-explicit-multipass: false
    default-save-temps: false

CMake Configuration

cmake .. -DCMAKE_BUILD_TYPE:STRING=Release \
-DCMAKE_C_COMPILER=cc \
-DCMAKE_CXX_COMPILER=CC \
-DCMAKE_INSTALL_PREFIX=${HIPSYCL_HOME} \
-DWITH_CPU_BACKEND:Bool=FALSE \
-DWITH_ROCM_BACKEND:Bool=TRUE \
-DWITH_CUDA_BACKEND:Bool=FALSE \
-DGPU_TARGETS="gfx90a" \
-DAMDGPU_TARGETS="gfx90a" \
-DDEFAULT_GPU_ARCH="gfx90a" \
-DWITH_ACCELERATED_CPU=OFF \
-DLLVM_DIR:String=${CLANG_DIR}/lib/cmake/llvm \
-DCLANG_INCLUDE_PATH:String=${CLANG_DIR}/lib/clang/${CLANG_VERSION}/include \
-DCLANG_EXECUTABLE_PATH:String=${CLANG_EXE} \
-DROCM_PATH:String=${ROCM_PATH} \
-DBoost_NO_BOOST_CMAKE=TRUE \
-DBoost_NO_SYSTEM_PATHS=TRUE \
-DBOOST_ROOT:PATHNAME=${BOOST_ROOT} \
-DWITH_SSCP_COMPILER=OFF \
-DBoost_LIBRARY_DIRS:FILEPATH=${BOOST_ROOT}/lib

SLURM Batch script

...

cat << EOF > select_gpu
#!/bin/bash

export ROCR_VISIBLE_DEVICES=\$SLURM_LOCALID
exec \$*
EOF

chmod +x ./select_gpu

srun  ./select_gpu <exec> <args>

So far, I haven't found a way to write a small reproducible example yet. Therefore, I just want to start a discussion.

Here is my `hipcc` version

HIP version: 5.4.22804-474e8620
AMD clang version 15.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-5.4.3 23045 a29fe425c7b0e5aba97ed2f95f61fd5ecba68aed)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /pfs/lustrep2/projappl/project_462000125/samantao-public/rocm/rocm-5.4.3/llvm/bin

The following code is called while calling hip_allocator::free(void *mem) (Regardless whether the memory was allocated with malloc_device or malloc_shared)
https://github.com/OpenSYCL/OpenSYCL/blob/12fdcaedfa990ab58ddf8bce304fa8cf917e6182/src/runtime/hip/hip_allocator.cpp#L122-L156

Does anybody have an idea why it is happening?

Another problem occurs with allocate_usm. hipMallocManaged does not return hipSuccess.

https://github.com/OpenSYCL/OpenSYCL/blob/12fdcaedfa990ab58ddf8bce304fa8cf917e6182/src/runtime/hip/hip_allocator.cpp#L98-L113

[ravil-mobile]

I guess I figured out what the culprit is. The application which I am working on is two explicitly coupled solvers. It has some complex mesh partitioning. It can so happen that one partition of the second solver won't get any elements to process (it becomes relevant during strong scaling when we substantially increase the number of MPI processes). Therefore, sycl::malloc_shared will be called requesting 0 bytes.

I guess CUDA manages to handle this special case gracefully but it is not the case for ROCm.

Does anybody know what SYCL standard tells about it? In my opinion, it is essential to establish the same behavior for all backends. In theory, it can be easily done. For example, when a user calls sycl::malloc_* requesting 0 bytes 1) either throw an exception (in all backends) or 2) return (let's say) a predefined pointer which can be later checked during sycl::free

[fodinabor]

There's some disagreement about what a SYCL impl should be doing when allocating 0 bytes...
See KhronosGroup/SYCL-Docs#355 and KhronosGroup/SYCL-Docs#356
I think the consensus is: don't throw an exception and as long as it's not dereferenced don't complain..?

[illuhad]

Ahahaha, this is currently the most contentious issue in the SYCL world :D The current conclusion is to not require consistent behavior and not throw an exception. (I was of the opinion that malloc(0) should return nullptr in any case).

However, malloc(0) should not fail.

Let me link this discussion for reference of an actual user code with this issue in the Khronos issue.

[ravil-mobile]

Hi @fodinabor and @illuhad, many thanks for answering.

I was of the opinion that malloc(0) should return nullptr in any case

Makes sense.

[keryell]

The funny part is that instead of adding this handling in the spec to solve this forever the committee preferred to spread the problem across all the users around...