rocm / rocclr

ROCclr development has been moved to https://github.com/ROCm/clr and this repo is not being used for active development since ROCm 5.6. Previously released branches are still available for reference.

rocclr's Issues

Intermittent Memory access fault by GPU - Reason: Unknown.

Hi,

2 out of 5 executions are failing on my program.

This seems like a better place than my original issue:
ROCm/HIP#2097

Any tips to debug would be appreciated.

System
Debian 10 buster-backports / kernel 5.5
GPU VEGA 64 RX Liquid
CPU - Ryzen 3950x @ stock
Mobo - Asrock Taichi x570 @ stock
RAM - 64GB DDR4 @ 3ghz XMPP auto

Multipe compilation errors

Ubuntu 20.04.2
gcc/g++ 9.3.0

cmake -DOPENCL_DIR="$OPENCL_DIR" -DCMAKE_INSTALL_PREFIX=/opt/rocm/rocclr ..
-- The C compiler identification is GNU 9.3.0
-- The CXX compiler identification is GNU 9.3.0
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Code Object Manager found at /opt/rocm/lib/cmake/amd_comgr.
-- HSA Runtime found at /opt/rocm/lib/cmake/hsa-runtime64.
-- Looking for pthread.h
-- Looking for pthread.h - found
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Failed
-- Check if compiler accepts -pthread
-- Check if compiler accepts -pthread - yes
-- Found Threads: TRUE
-- Found: /usr/lib/x86_64-linux-gnu/libnuma.so
-- Configuring done
-- Generating done
-- Build files have been written to: /home/ian/Documents/ROCclr/build
REDACTED:~/Documents/ROCclr/build$ make -j$(nproc)
Scanning dependencies of target oclrocm
[ 2%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocprintf.cpp.o
[ 4%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocprogram.cpp.o
[ 6%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocsettings.cpp.o
[ 8%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/roccounters.cpp.o
[ 15%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocmemory.cpp.o
[ 15%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocvirtual.cpp.o
[ 17%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocblit.cpp.o
[ 17%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rockernel.cpp.o
[ 19%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocdevice.cpp.o
[ 21%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocschedcl.cpp.o
[ 23%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rochostcall.cpp.o
[ 26%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocappprofile.cpp.o
[ 28%] Building CXX object device/rocm/CMakeFiles/oclrocm.dir/rocglinterop.cpp.o
/home/ian/Documents/ROCclr/device/rocm/rochostcall.cpp: In member function ‘void HostcallListener::consumePackets()’:
/home/ian/Documents/ROCclr/device/rocm/rochostcall.cpp:319:28: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
319 | uint64_t new_value = hsa_signal_wait_scacquire(doorbell_, HSA_SIGNAL_CONDITION_NE, signal_value, timeout,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
/home/ian/Documents/ROCclr/device/rocm/rochostcall.cpp: In member function ‘void HostcallListener::terminate()’:
/home/ian/Documents/ROCclr/device/rocm/rochostcall.cpp:347:3: error: ‘hsa_signal_store_screlease’ was not declared in this scope; did you mean ‘hsa_signal_store_release’?
347 | hsa_signal_store_screlease(doorbell_, SIGNAL_DONE);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_store_release
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:219: device/rocm/CMakeFiles/oclrocm.dir/rochostcall.cpp.o] Error 1
make[2]: *** Waiting for unfinished jobs....
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocmemory.hpp:28,
from /home/ian/Documents/ROCclr/device/rocm/rocmemory.cpp:31:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:30,
from /home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:23:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocmemory.hpp:28,
from /home/ian/Documents/ROCclr/device/rocm/rocprintf.cpp:25:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/roccounters.hpp:26,
from /home/ian/Documents/ROCclr/device/rocm/roccounters.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:35:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocblit.cpp:22:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocmemory.hpp:28,
from /home/ian/Documents/ROCclr/device/rocm/rocmemory.cpp:31:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:30,
from /home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:23:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:38,
from /home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:30,
from /home/ian/Documents/ROCclr/device/rocm/rockernel.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rockernel.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp: In function ‘bool roc::WaitForSignal(hsa_signal_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.hpp:54:7: error: ‘hsa_signal_wait_scacquire’ was not declared in this scope; did you mean ‘hsa_signal_wait_acquire’?
54 | if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_wait_acquire
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocmemory.hpp:28,
from /home/ian/Documents/ROCclr/device/rocm/rocprintf.cpp:25:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/roccounters.hpp:26,
from /home/ian/Documents/ROCclr/device/rocm/roccounters.cpp:21:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:21:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:35:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocblit.cpp:22:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:46,
from /home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:30,
from /home/ian/Documents/ROCclr/device/rocm/rockernel.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rockernel.cpp:21:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h: At global scope:
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:277:5: error: ‘hsa_loaded_code_object_t’ has not been declared
277 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t *code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t *code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t *code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:442:3: error: ‘hsa_loaded_code_object_t’ was not declared in this scope; did you mean ‘hsa_code_object_t’?
442 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:443:48: error: expected primary-expression before ‘attribute’
443 | hsa_ven_amd_loader_loaded_code_object_info_t attribute,
| ^~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:3: error: expected primary-expression before ‘void’
444 | void *value);
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:444:14: error: expression list treated as compound expression in initializer [-fpermissive]
444 | void *value);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t *code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:489:5: error: ‘hsa_file_t’ was not declared in this scope; did you mean ‘hsa_isa_t’?
489 | hsa_file_t file,
| ^~~~~~~~~~
| hsa_isa_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:490:12: error: expected primary-expression before ‘offset’
490 | size_t offset,
| ^~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:491:12: error: expected primary-expression before ‘size’
491 | size_t size,
| ^~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t *code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:5: error: ‘hsa_code_object_reader_t’ was not declared in this scope; did you mean ‘hsa_code_object_type_t’?
492 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:23:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:83:3: error: ‘hsa_code_object_reader_t’ does not name a type; did you mean ‘hsa_code_object_type_t’?
83 | hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:31: error: ‘code_object_reader’ was not declared in this scope
492 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:492:49: error: expression list treated as compound expression in initializer [-fpermissive]
492 | hsa_code_object_reader_t code_object_reader);
| ^
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:538:7: error: ‘hsa_loaded_code_object_t’ has not been declared
538 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:543:5: error: ‘hsa_loaded_code_object_t’ has not been declared
543 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:568:7: error: ‘hsa_loaded_code_object_t’ has not been declared
568 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:573:5: error: ‘hsa_loaded_code_object_t’ has not been declared
573 | hsa_loaded_code_object_t loaded_code_object,
| ^~~~~~~~~~~~~~~~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:579:7: error: ‘hsa_file_t’ has not been declared
579 | hsa_file_t file,
| ^~~~~~~~~~
/opt/rocm/include/hsa/hsa_ven_amd_loader.h:582:7: error: ‘hsa_code_object_reader_t’ has not been declared
582 | hsa_code_object_reader_t code_object_reader);
| ^~~~~~~~~~~~~~~~~~~~~~~~
In file included from /home/ian/Documents/ROCclr/device/rocm/rockernel.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rocprintf.cpp:26:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:83:3: error: ‘hsa_code_object_reader_t’ does not name a type; did you mean ‘hsa_code_object_type_t’?
83 | hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
In file included from /home/ian/Documents/ROCclr/device/rocm/roccounters.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/roccounters.hpp: In member function ‘bool roc::PerfCounterProfile::Create()’:
/home/ian/Documents/ROCclr/device/rocm/roccounters.hpp:121:36: error: ‘HSA_EXTENSION_AMD_AQLPROFILE’ was not declared in this scope; did you mean ‘HSA_EXTENSION_AMD_PROFILER’?
121 | hsa_system_extension_supported(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &system_support);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
| HSA_EXTENSION_AMD_PROFILER
In file included from /home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:38:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:83:3: error: ‘hsa_code_object_reader_t’ does not name a type; did you mean ‘hsa_code_object_type_t’?
83 | hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
In file included from /home/ian/Documents/ROCclr/device/rocm/rockernel.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:23:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:83:3: error: ‘hsa_code_object_reader_t’ does not name a type; did you mean ‘hsa_code_object_type_t’?
83 | hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp: In destructor ‘virtual roc::Program::~Program()’:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:78:7: error: ‘hsaCodeObjectReader_’ was not declared in this scope
78 | if (hsaCodeObjectReader_.handle != 0) {
| ^~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/roccounters.hpp:127:9: error: ‘hsa_system_get_major_extension_table’ was not declared in this scope; did you mean ‘hsa_system_get_extension_table’?
127 | if (hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_AQLPROFILE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_system_get_extension_table
In file included from /home/ian/Documents/ROCclr/device/rocm/rockernel.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rockernel.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:83:3: error: ‘hsa_code_object_reader_t’ does not name a type; did you mean ‘hsa_code_object_type_t’?
83 | hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
In file included from /home/ian/Documents/ROCclr/device/rocm/rockernel.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rocblit.cpp:25:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.hpp:83:3: error: ‘hsa_code_object_reader_t’ does not name a type; did you mean ‘hsa_code_object_type_t’?
83 | hsa_code_object_reader_t hsaCodeObjectReader_; //!< Handle to HSA code reader
| ^~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_type_t
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:79:5: error: ‘hsa_code_object_reader_destroy’ was not declared in this scope; did you mean ‘hsa_code_object_destroy’?
79 | hsa_code_object_reader_destroy(hsaCodeObjectReader_);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_code_object_destroy
In file included from /home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:26:
/home/ian/Documents/ROCclr/device/rocm/roccounters.hpp: In member function ‘bool roc::PerfCounterProfile::Create()’:
/home/ian/Documents/ROCclr/device/rocm/roccounters.hpp:121:36: error: ‘HSA_EXTENSION_AMD_AQLPROFILE’ was not declared in this scope; did you mean ‘HSA_EXTENSION_AMD_PROFILER’?
121 | hsa_system_extension_supported(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &system_support);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
| HSA_EXTENSION_AMD_PROFILER
/home/ian/Documents/ROCclr/device/rocm/rockernel.cpp: In member function ‘virtual bool roc::LightningKernel::init()’:
/home/ian/Documents/ROCclr/device/rocm/rockernel.cpp:75:15: error: ‘hsa_executable_get_symbol_by_name’ was not declared in this scope; did you mean ‘hsa_executable_get_symbol’?
75 | hsaStatus = hsa_executable_get_symbol_by_name(program()->hsaExecutable(),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_executable_get_symbol
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp: In static member function ‘static bool roc::Device::init()’:
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:470:40: error: ‘HSA_EXTENSION_AMD_LOADER’ was not declared in this scope; did you mean ‘HSA_EXTENSION_AMD_PROFILER’?
470 | hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(amd_loader_ext_table),
| ^~~~~~~~~~~~~~~~~~~~~~~~
| HSA_EXTENSION_AMD_PROFILER
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp: In constructor ‘roc::Program::Program(roc::NullDevice&, amd::Program&)’:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:86:3: error: ‘hsaCodeObjectReader_’ was not declared in this scope
86 | hsaCodeObjectReader_.handle = 0;
| ^~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocblit.cpp: In member function ‘bool roc::DmaBlitManager::hsaCopyStaged(const_address, address, size_t, address, bool) const’:
/home/ian/Documents/ROCclr/device/rocm/rocblit.cpp:682:5: error: ‘hsa_signal_silent_store_relaxed’ was not declared in this scope; did you mean ‘hsa_signal_store_relaxed’?
682 | hsa_signal_silent_store_relaxed(completion_signal_, kInitSignalValueOne);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_store_relaxed
/home/ian/Documents/ROCclr/device/rocm/roccounters.hpp:127:9: error: ‘hsa_system_get_major_extension_table’ was not declared in this scope; did you mean ‘hsa_system_get_extension_table’?
127 | if (hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_AQLPROFILE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_system_get_extension_table
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:470:3: error: ‘hsa_system_get_major_extension_table’ was not declared in this scope; did you mean ‘hsa_system_get_extension_table’?
470 | hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(amd_loader_ext_table),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_system_get_extension_table
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp: In member function ‘bool roc::VirtualGPU::dispatchGenericAqlPacket(AqlPacket, uint16_t, uint16_t, bool, size_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:452:20: error: there are no arguments to ‘hsa_queue_add_write_index_screlease’ that depend on a template parameter, so a declaration of ‘hsa_queue_add_write_index_screlease’ must be available [-fpermissive]
452 | uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:452:20: note: (if you use ‘-fpermissive’, G++ will accept your code, but allowing the use of an undeclared name is deprecated)
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:472:19: error: there are no arguments to ‘hsa_queue_load_read_index_scacquire’ that depend on a template parameter, so a declaration of ‘hsa_queue_load_read_index_scacquire’ must be available [-fpermissive]
472 | while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) {
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp: In member function ‘virtual bool roc::Program::createGlobalVarObj(amd::Memory, void, size_t, const char) const’:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:159:12: error: ‘hsa_executable_get_symbol_by_name’ was not declared in this scope; did you mean ‘hsa_executable_get_symbol’?
159 | status = hsa_executable_get_symbol_by_name(hsaExecutable_, global_name, &hsa_device,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_executable_get_symbol
In file included from /home/ian/Documents/ROCclr/include/top.hpp:101,
from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:513:36: error: ‘HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE’?
513 | extractAqlBits(header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:514:28: error: ‘HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE’?
514 | HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp: In member function ‘bool roc::Device::populateOCLDeviceConstants()’:
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:989:7: error: ‘hsa_isa_get_info_alt’ was not declared in this scope; did you mean ‘hsa_isa_get_info’?
989 | if (hsa_isa_get_info_alt(isa, HSA_ISA_INFO_NAME_LENGTH, &isaNameLength) != HSA_STATUS_SUCCESS) {
| ^~~~~~~~~~~~~~~~~~~~
| hsa_isa_get_info
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp: In member function ‘virtual bool roc::LightningProgram::setKernels(amd::option::Options, void, size_t, amd::Os::FileDesc, size_t, std::string)’:
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:484:12: error: ‘hsa_executable_create_alt’ was not declared in this scope; did you mean ‘hsa_executable_create’?
484 | status = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_executable_create
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:515:36: error: ‘HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE’?
515 | extractAqlBits(header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:997:7: error: ‘hsa_isa_get_info_alt’ was not declared in this scope; did you mean ‘hsa_isa_get_info’?
997 | if (hsa_isa_get_info_alt(isa, HSA_ISA_INFO_NAME, info_.targetId_) != HSA_STATUS_SUCCESS) {
| ^~~~~~~~~~~~~~~~~~~~
| hsa_isa_get_info
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:516:28: error: ‘HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE’?
516 | HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:496:72: error: ‘hsaCodeObjectReader_’ was not declared in this scope
496 | status = hsa_code_object_reader_create_from_memory(binary, binSize, &hsaCodeObjectReader_);
| ^~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:513:36: error: ‘HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE’?
513 | extractAqlBits(header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:514:28: error: ‘HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE’?
514 | HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:515:36: error: ‘HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE’?
515 | extractAqlBits(header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:496:12: error: ‘hsa_code_object_reader_create_from_memory’ was not declared in this scope
496 | status = hsa_code_object_reader_create_from_memory(binary, binSize, &hsaCodeObjectReader_);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp: In member function ‘hsa_queue_t roc::Device::acquireQueue(uint32_t, bool, const std::vector&, amd::CommandQueue::Priority)’:
/home/ian/Documents/ROCclr/device/rocm/rocdevice.cpp:2372:18: error: ‘HSA_QUEUE_TYPE_COOPERATIVE’ was not declared in this scope; did you mean ‘HSA_QUEUE_TYPE_MULTI’?
2372 | queue_type = HSA_QUEUE_TYPE_COOPERATIVE;
| ^~~~~~~~~~~~~~~~~~~~~~~~~~
| HSA_QUEUE_TYPE_MULTI
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:516:28: error: ‘HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE’?
516 | HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:531:3: error: there are no arguments to ‘hsa_signal_store_screlease’ that depend on a template parameter, so a declaration of ‘hsa_signal_store_screlease’ must be available [-fpermissive]
531 | hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocprogram.cpp:504:12: error: ‘hsa_executable_load_agent_code_object’ was not declared in this scope; did you mean ‘hsa_executable_load_code_object’?
504 | status = hsa_executable_load_agent_code_object(hsaExecutable_, agent, hsaCodeObjectReader_, nullptr,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_executable_load_code_object
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp: In member function ‘void roc::VirtualGPU::dispatchBarrierPacket(const hsa_barrier_and_packet_t)’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:591:20: error: ‘hsa_queue_add_write_index_screlease’ was not declared in this scope; did you mean ‘hsa_queue_add_write_index_release’?
591 | uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_queue_add_write_index_release
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:592:19: error: ‘hsa_queue_load_read_index_scacquire’ was not declared in this scope; did you mean ‘hsa_queue_load_read_index_acquire’?
592 | while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_queue_load_read_index_acquire
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:598:3: error: ‘hsa_signal_store_screlease’ was not declared in this scope; did you mean ‘hsa_signal_store_release’?
598 | hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
| ^~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_store_release
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:76: device/rocm/CMakeFiles/oclrocm.dir/rocprintf.cpp.o] Error 1
In file included from /home/ian/Documents/ROCclr/include/top.hpp:101,
from /home/ian/Documents/ROCclr/device/rocm/rocdevice.hpp:25,
from /home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:21:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:607:48: error: ‘HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE’?
607 | extractAqlBits(kBarrierPacketHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:102: device/rocm/CMakeFiles/oclrocm.dir/rocmemory.cpp.o] Error 1
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:608:27: error: ‘HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE’?
608 | HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:609:48: error: ‘HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE’?
609 | extractAqlBits(kBarrierPacketHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:610:27: error: ‘HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE’?
610 | HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:193:68: note: in definition of macro ‘ClPrint’
193 | amd::log_printf(level, FILENAME, LINE, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:607:48: error: ‘HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE’?
607 | extractAqlBits(kBarrierPacketHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:141: device/rocm/CMakeFiles/oclrocm.dir/rockernel.cpp.o] Error 1
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:608:27: error: ‘HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE’?
608 | HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:609:48: error: ‘HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE’?
609 | extractAqlBits(kBarrierPacketHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:610:27: error: ‘HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE’ was not declared in this scope; did you mean ‘HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE’?
610 | HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/ian/Documents/ROCclr/utils/debug.hpp:195:51: note: in definition of macro ‘ClPrint’
195 | amd::log_printf(level, "", 0, format, ##VA_ARGS);
| ^~~~~~~~~~~
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:63: device/rocm/CMakeFiles/oclrocm.dir/roccounters.cpp.o] Error 1
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:128: device/rocm/CMakeFiles/oclrocm.dir/rocblit.cpp.o] Error 1
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp: In instantiation of ‘bool roc::VirtualGPU::dispatchGenericAqlPacket(AqlPacket, uint16_t, uint16_t, bool, size_t) [with AqlPacket = hsa_kernel_dispatch_packet_s; uint16_t = short unsigned int; size_t = long unsigned int]’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:547:65: required from here
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:452:55: error: ‘hsa_queue_add_write_index_screlease’ was not declared in this scope; did you mean ‘hsa_queue_add_write_index_release’?
452 | uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~
| hsa_queue_add_write_index_release
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:472:54: error: ‘hsa_queue_load_read_index_scacquire’ was not declared in this scope; did you mean ‘hsa_queue_load_read_index_acquire’?
472 | while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) {
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~
| hsa_queue_load_read_index_acquire
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:531:29: error: ‘hsa_signal_store_screlease’ was not declared in this scope; did you mean ‘hsa_signal_store_release’?
531 | hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_store_release
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:89: device/rocm/CMakeFiles/oclrocm.dir/rocprogram.cpp.o] Error 1
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp: In instantiation of ‘bool roc::VirtualGPU::dispatchGenericAqlPacket(AqlPacket, uint16_t, uint16_t, bool, size_t) [with AqlPacket = hsa_barrier_and_packet_s; uint16_t = short unsigned int; size_t = long unsigned int]’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:553:65: required from here
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:452:55: error: ‘hsa_queue_add_write_index_screlease’ was not declared in this scope; did you mean ‘hsa_queue_add_write_index_release’?
452 | uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~
| hsa_queue_add_write_index_release
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:472:54: error: ‘hsa_queue_load_read_index_scacquire’ was not declared in this scope; did you mean ‘hsa_queue_load_read_index_acquire’?
472 | while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) {
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~
| hsa_queue_load_read_index_acquire
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:531:29: error: ‘hsa_signal_store_screlease’ was not declared in this scope; did you mean ‘hsa_signal_store_release’?
531 | hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_store_release
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp: In instantiation of ‘bool roc::VirtualGPU::dispatchGenericAqlPacket(AqlPacket, uint16_t, uint16_t, bool, size_t) [with AqlPacket = hsa_ext_amd_aql_pm4_packet_t; uint16_t = short unsigned int; size_t = long unsigned int]’:
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:570:90: required from here
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:452:55: error: ‘hsa_queue_add_write_index_screlease’ was not declared in this scope; did you mean ‘hsa_queue_add_write_index_release’?
452 | uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~
| hsa_queue_add_write_index_release
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:472:54: error: ‘hsa_queue_load_read_index_scacquire’ was not declared in this scope; did you mean ‘hsa_queue_load_read_index_acquire’?
472 | while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) {
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~
| hsa_queue_load_read_index_acquire
/home/ian/Documents/ROCclr/device/rocm/rocvirtual.cpp:531:29: error: ‘hsa_signal_store_screlease’ was not declared in this scope; did you mean ‘hsa_signal_store_release’?
531 | hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
| hsa_signal_store_release
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:115: device/rocm/CMakeFiles/oclrocm.dir/rocdevice.cpp.o] Error 1
make[2]: *** [device/rocm/CMakeFiles/oclrocm.dir/build.make:154: device/rocm/CMakeFiles/oclrocm.dir/rocvirtual.cpp.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:123: device/rocm/CMakeFiles/oclrocm.dir/all] Error 2
make: *** [Makefile:130: all] Error 2

(hipStreamDestory)Memory Leak Fix!!! Find by valgrind

bool HostQueue::terminate() {
.....
// Wake-up the command loop, so it can exit
    {
      ScopedLock sl(queueLock_);
      thread_.acceptingCommands_ = false;
      queueLock_.notify();
    }
// FIXME_lmoriche: fix termination handshake
    while (thread_.state() < Thread::FINISHED) {
      Os::yield();
    }

// Release last enque command !! Memory leak
if (lastEnqueueCommand_ != nullptr) {
      lastEnqueueCommand_->release();
      lastEnqueueCommand_ = nullptr;
    }

}

MLIR test crashes because getVQVirtualAddress() is called when virtualQueue_ is nullptr

The AMD MLIR group has a buildbot that tests the main-line MLIR repo with some AMD-specific options. One test has recently started failing, after a change in the codegen. The problem is not with that change, but rather with the cleanup after the test's kernel has run. The backtrace of the crash is

#1 0x7f53e6d63a01 in roc::Device::getRocMemory(amd::Memory*) const (/usr/home/pf/hipamd/build/lib/libamdhip64.so.5+0x53ea01)
#2 0x7f53e6da1a4b in roc::VirtualGPU::getVQVirtualAddress() (/usr/home/pf/hipamd/build/lib/libamdhip64.so.5+0x57ca4b)
#3 0x7f53e6daea68 in roc::VirtualGPU::submitKernelInternal(amd::NDRangeContainer const&, amd::Kernel const&, unsigned char const*, void*, unsigned int, amd::NDRangeKernelCommand*) (/usr/home/pf/hipamd/build/lib/libamdhip64.so.5+0x589a68)
#4 0x7f53e6daf591 in roc::VirtualGPU::submitKernel(amd::NDRangeKernelCommand&) (/usr/home/pf/hipamd/build/lib/libamdhip64.so.5+0x58a591)
#5 0x7f53e6d1d34d in amd::NDRangeKernelCommand::submit(device::VirtualDevice&) /home/pf/ROCclr/cmake/../platform/command.hpp:1059:63
#6 0x7f53e6d15995 in amd::Command::enqueue() /home/pf/ROCclr/platform/command.cpp:370:7
#7 0x7f53e6b84e2b in ihipModuleLaunchKernel(ihipModuleSymbol_t*, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, ihipStream_t*, void**, void**, ihipEvent_t*, ihipEvent_t*, unsigned int, unsigned int, unsigned int, unsigned int, unsigned long, unsigned long, unsigned int) /home/pf/hipamd/src/hip_module.cpp:394:12
#8 0x7f53e6b85bdf in hipModuleLaunchKernel /home/pf/hipamd/src/hip_module.cpp:423:3
#9 0x7f53e78ffd7f in mgpuLaunchKernel /usr/home/pf/llvm-project/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp:61:3
#10 0x7f53dded1089  (<unknown module>)
#11 0x7f53dded10dc  (<unknown module>)
#12 0x6a294a1 in compileAndExecute((anonymous namespace)::Options&, mlir::Operation*, llvm::StringRef, (anonymous namespace)::CompileAndExecuteConfig, void**, std::unique_ptr<llvm::TargetMachine, std::default_delete<llvm::TargetMachine> >) (/usr/home/pf/llvm-project/build/bin/mlir-cpu-runner+0x6a294a1)
#13 0x6a281c7 in compileAndExecuteVoidFunction((anonymous namespace)::Options&, mlir::Operation*, llvm::StringRef, (anonymous namespace)::CompileAndExecuteConfig, std::unique_ptr<llvm::TargetMachine, std::default_delete<llvm::TargetMachine> >) (/usr/home/pf/llvm-project/build/bin/mlir-cpu-runner+0x6a281c7)
#14 0x6a2487d in mlir::JitRunnerMain(int, char**, mlir::DialectRegistry const&, mlir::JitRunnerConfig) (/usr/home/pf/llvm-project/build/bin/mlir-cpu-runner+0x6a2487d)
#15 0x555a280 in main /usr/home/pf/llvm-project/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp:33:10
#16 0x7f53f6391082 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x24082) (BuildId: 1878e6b475720c7c51969e69ab2d276fae6d1dee)
#17 0x545e8ad in _start (/usr/home/pf/llvm-project/build/bin/mlir-cpu-runner+0x545e8ad)

Frames 17 up to 9 are from mlir-cpu-runner, which takes an MLIR file and runs it on the CPU. Then there are a couple of HIP frames, and then we're into ROCclr (via libamdhip64.so). We're in submitKernelInternal(), near the end, just after the printf buffer has been output. (The test happens to be just a printf, and I can see what it prints, before the crash.) It calls runScheduler(), and one argument is getVQVirtualAddress():

  if (gpuKernel.dynamicParallelism()) {
      dispatchBarrierPacket(kBarrierPacketHeader, true);
      static_cast<KernelBlitManager&>(blitMgr()).runScheduler(
          getVQVirtualAddress(), schedulerParam_, schedulerQueue_, schedulerSignal_, schedulerThreads_);
  }

The problem is that virtualQueue_ has never been set. Its default value is nullptr. There are two setup paths that give it a value, but neither is taken by this test. A call under getVQVirtualAddress() tries to read through virtualQueue_, and barfs because it's nullptr.

I do not know whether this code should guard against nullptr, or if there is a guarantee that one of the setup paths will be taken. For testing, I've added "&& virtualQueue_" to the IF quoted above, and that passes all the MLIR tests.

To reproduce, make sure that ROCm is installed (I was using 5.4.2), then compile and test as directed in https://mlir.llvm.org/getting_started/, with the additional cmake options "-DMLIR_INCLUDE_INTEGRATION_TESTS=ON -DMLIR_ENABLE_ROCM_RUNNER=ON -DMLIR_ENABLE_ROCM_CONVERSIONS=ON". I've been using Ubuntu 20.04, specifically a docker image based on the rocm/mlir:rocm5.4-latest image.

data race in memory update function

IMO the memory counting functions contain data races which lead to the possibility that the counter for free memory is underflowing.

https://github.com/ROCm-Developer-Tools/ROCclr/blob/90f1f61a9d6c28ffd2f844dc773e921444752e47/device/rocm/rocdevice.cpp#L2086-L2104

• Case 1
  • Two threads both allocating 2MiB going into line 2091 and the check is false because freeMem_ is 3MiB.
  • Both threads will decrement the counter freeMem_ in line 2101 which results in an underflow
• Case 2
  • Two threads where one is deallocating 2MiB and the other is allocating 5MiB and freeMem_ is 4MiB
  • The allocating thread (5MiB) is checking line 2091 and is going into the if body to line 2096
  • The second thread with the deallocation of 2MiB is exciting line 2088, freeMem_ is now 6MiB
  • The first thread is executing line 2098 which is setting freeMem_ to zero.
  • The result is that we lose 2MiB of memory we could potential allocate but due to the data race is not available anymore

Possible solution:

• When memory is freed freeMem_ is stored in a register and atomic CAS is used to reset the variable freeMem_
• line 2101 must be guarded by atomic Cas too to avoid that other threads reducing the value of freeMem_ in the same moment which results into a variant of Case 1

SIGSEGV in Device::setupCpuAgent

When running clinfo (system or ROCm version) I get a segmentation fault. This is on tag rocm-3.7.0. My system is OpenSuse Tumbleweed. This happens if I install from the official ROCm SUSE repository, or if I compile rocclr and rocm-opencl-runtime myself the issue persists (with other packages as the official packages). Compiling those two with debug is how I got the nice backtrace.

This is with AMD Ryzen 7 PRO 2700U Vega using the APU (integrated graphics).

Thread 1 "clinfo" received signal SIGSEGV, Segmentation fault.
0x00007ffff7923266 in roc::Device::setupCpuAgent (this=0x4f3870) at /home/samantha/git/ROCm/ROCclr/device/rocm/rocdevice.cpp:170
170       cpu_agent_ = cpu_agents_[index].agent;
(gdb) bt
#0  0x00007ffff7923266 in roc::Device::setupCpuAgent (this=0x4f3870) at /home/samantha/git/ROCm/ROCclr/device/rocm/rocdevice.cpp:170
#1  0x00007ffff7926840 in roc::Device::populateOCLDeviceConstants (this=0x4f3870) at /home/samantha/git/ROCm/ROCclr/device/rocm/rocdevice.cpp:1047
#2  0x00007ffff7924d08 in roc::Device::create (this=0x4f3870) at /home/samantha/git/ROCm/ROCclr/device/rocm/rocdevice.cpp:593
#3  0x00007ffff79245a1 in roc::Device::init () at /home/samantha/git/ROCm/ROCclr/device/rocm/rocdevice.cpp:489
#4  0x00007ffff78b7390 in amd::Device::init () at /home/samantha/git/ROCm/ROCclr/device/device.cpp:194
#5  0x00007ffff78d2a53 in amd::Runtime::init () at /home/samantha/git/ROCm/ROCclr/platform/runtime.cpp:74
#6  0x00007ffff78a318b in ShouldLoadPlatform () at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/amdocl/cl_icd.cpp:222
#7  0x00007ffff78a3287 in operator() (__closure=0x7fffffffcdaf) at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/amdocl/cl_icd.cpp:272
#8  0x00007ffff78a366b in std::__invoke_impl<void, clIcdGetPlatformIDsKHR(cl_uint, _cl_platform_id**, cl_uint*)::<lambda()> >(std::__invoke_other, struct {...} &&) (__f=...) at /usr/include/c++/10/bits/invoke.h:60
#9  0x00007ffff78a363a in std::__invoke<clIcdGetPlatformIDsKHR(cl_uint, _cl_platform_id**, cl_uint*)::<lambda()> >(struct {...} &&) (__fn=...) at /usr/include/c++/10/bits/invoke.h:95
#10 0x00007ffff78a3531 in operator() (this=0x7fffffffcd50) at /usr/include/c++/10/mutex:717
#11 0x00007ffff78a355b in operator() (this=0x0) at /usr/include/c++/10/mutex:722
#12 0x00007ffff78a356c in _FUN () at /usr/include/c++/10/mutex:722
#13 0x00007ffff7c2e36f in __pthread_once_slow () from /lib64/libpthread.so.0
#14 0x00007ffff78a3084 in __gthread_once (__once=0x7ffff7a4b050 <clIcdGetPlatformIDsKHR::initOnce>, __func=0x7ffff7e754e0 <std::__once_proxy()>) at /usr/include/c++/10/x86_64-suse-linux/bits/gthr-default.h:700
#15 0x00007ffff78a35f2 in std::call_once<clIcdGetPlatformIDsKHR(cl_uint, _cl_platform_id**, cl_uint*)::<lambda()> >(std::once_flag &, struct {...} &&) (__once=..., __f=...) at /usr/include/c++/10/mutex:729
#16 0x00007ffff78a32e1 in clIcdGetPlatformIDsKHR (num_entries=0, platforms=0x0, num_platforms=0x7fffffffcdec) at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/amdocl/cl_icd.cpp:272
#17 0x00007ffff7fc236d in khrIcdVendorAdd (libraryName=0x442110 "libamdocl64.so") at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/khronos/icd/loader/icd.c:87
#18 0x00007ffff7fc5d3e in khrIcdOsVendorsEnumerate () at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/khronos/icd/loader/linux/icd_linux.c:125
#19 0x00007ffff7c2e36f in __pthread_once_slow () from /lib64/libpthread.so.0
#20 0x00007ffff7fc5dcc in khrIcdOsVendorsEnumerateOnce () at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/khronos/icd/loader/linux/icd_linux.c:149
#21 0x00007ffff7fc2272 in khrIcdInitialize () at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/khronos/icd/loader/icd.c:31
#22 0x00007ffff7fc27c3 in clGetPlatformIDs (num_entries=0, platforms=0x0, num_platforms=0x7fffffffcfa0) at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/khronos/icd/loader/icd_dispatch.c:34
#23 0x0000000000408216 in cl::Platform::get (platforms=0x7fffffffd370) at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/khronos/headers/opencl2.2/CL/cl2.hpp:2482
#24 0x0000000000403813 in main (argc=1, argv=0x7fffffffda98) at /home/samantha/git/ROCm/ROCm-OpenCL-Runtime/tools/clinfo/clinfo.cpp:75

Using textures causes "Depth set on single layer image geometry."

See ROCm/ROCR-Runtime#113

According to HSA v1.1 http://www.hsafoundation.com/html_spec111/HSA_Library.htm#Runtime/Topics/03_Extensions/hsa_ext_image_descriptor_t.htm%3FTocPath%3DHSA%2520Runtime%2520Programmer%25E2%2580%2599s%2520Reference%2520Manual%2520Version%25201.1.1%2520|Chapter%25203.%2520HSA%2520Extensions%2520Programming%2520Guide|3.3%2520Images%2520and%2520samplers|3.3.1%2520Images%2520and%2520samplers%2520API|_____10

height, depth and array_size must be 0 if they are not used with the current geometry.

ROCclr sets them to 1: https://github.com/ROCm-Developer-Tools/ROCclr/blob/main/device/rocm/rocmemory.cpp#L959

Create (extra) 3.5.0 release with consistent download URL

Currently the download URL for the releases is inconsistent:

• https://github.com/ROCm-Developer-Tools/ROCclr/archive/roc-3.5.0.tar.gz
• https://github.com/ROCm-Developer-Tools/ROCclr/archive/rocm-3.7.0.tar.gz

Could you maybe create another Github release for 3.5.0 which reads rocm-3.5.0? This makes it easier for packages managers to automatically detect new releases/versions and have an educated guess about their download url.

hsa-runtime_DIR is not a valid identifier

If I attempt to set the hsa-runtime64_DIR environment variable, this fails in bash version 4.4.23 with the message that this is not a valid identifier. This post seems to suggest environment variables with dashes are not valid in POSIX.

`kind` may be used uninitialized if preceding `amd::Comgr::index_list_metadata` call fails.

kind may be used uninitialized if preceding amd::Comgr::index_list_metadata call fails.

https://github.com/ROCm-Developer-Tools/ROCclr/blob/ca35d38e98c3bf91da23707820060d1f68dd6e95/device/devkernel.cpp#L1313

  amd_comgr_metadata_kind_t kind;
  status = amd::Comgr::index_list_metadata(argsMeta, i, &argsNode);
  if (status == AMD_COMGR_STATUS_SUCCESS) {
    hsaArgsNode = true;
    // kind is set only if status==success.  ----->   vvvv
    status = amd::Comgr::get_metadata_kind(argsNode, &kind); 
  }
  if (kind != AMD_COMGR_METADATA_KIND_MAP) { // <-- kind may not be initialized.

MIssing licensing info

The project does not seem to include the top-level license file.
Some of the individual source files (e.g device/gpu/gslbe/**) also do not have any licensing info in them.

Please update the project with the relevant info. Without it it's unclear what licensing terms apply to the files without individual license header in them.

Hard GLX dependency prevents builds on systems without X

There seems to be a hard dependency on Unix platforms to use a GL interface provided by GLX. This causes problems when X is not available, such as when osmesa is the GL provider.

Is there a reason for including the GLX header explicitly here? Should this be guarded differently for cases when glx is not available?

build is ok but make install fails, it is not clear how to install onto deafult /opt/rocm directory as a result.

Using ubuntu18.04 rocm4.5.2 source to build rocclr:

readme file states:

Build ROCclr

Here is command to build ROCclr:

cd "$ROCclr_DIR"
mkdir -p build; cd build
cmake -DOPENCL_DIR="$OPENCL_DIR" -DCMAKE_INSTALL_PREFIX=/opt/rocm/rocclr ..
make -j$(nproc)
sudo make install

Everything until make -j8 is OK including make.
but make install target is not recognized. as a result, build files stuck in SOURCE build directory and can not install onto /opt/rocm because of it

cmake -DOPENCL_DIR="$OPENCL_DIR" -DCMAKE_INSTALL_PREFIX=/opt/rocm/rocclr ..
-- The C compiler identification is GNU 7.5.0
-- The CXX compiler identification is GNU 7.5.0
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
CMake Warning (dev) at CMakeLists.txt:26 (message):
  ROCclr is being built as a standalone project.  This isn't supported
  anymore.
This warning is for project developers.  Use -Wno-dev to suppress it.

-- Looking for pthread.h
-- Looking for pthread.h - found
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Failed
-- Check if compiler accepts -pthread
-- Check if compiler accepts -pthread - yes
-- Found Threads: TRUE  
-- Found AMD_OPENCL: /root/ROCm-4.5/ROCm-OpenCL-Runtime/khronos/headers/opencl2.2/CL  
-- Found NUMA: /usr/lib/x86_64-linux-gnu/libnuma.so  
-- Configuring done
-- Generating done
CMake Warning:
  Manually-specified variables were not used by the project:

    OPENCL_DIR


-- Build files have been written to: /root/ROCm-4.5/ROCclr/build
root@localhost:~/ROCm-4.5/ROCclr/build# ls
CMakeCache.txt  CMakeFiles  Makefile  cmake_install.cmake
root@localhost:~/ROCm-4.5/ROCclr/build# make -j8
Scanning dependencies of target rocclr
[  2%] Building CXX object CMakeFiles/rocclr.dir/compiler/lib/utils/options.cpp.o
[  4%] Building CXX object CMakeFiles/rocclr.dir/device/appprofile.cpp.o
[  6%] Building CXX object CMakeFiles/rocclr.dir/device/blit.cpp.o
[  8%] Building CXX object CMakeFiles/rocclr.dir/device/blitcl.cpp.o
[ 12%] Building CXX object CMakeFiles/rocclr.dir/device/devhcmessages.cpp.o
[ 12%] Building CXX object CMakeFiles/rocclr.dir/device/comgrctx.cpp.o
[ 14%] Building CXX object CMakeFiles/rocclr.dir/device/devhcprintf.cpp.o
[ 16%] Building CXX object CMakeFiles/rocclr.dir/device/devhostcall.cpp.o
[ 18%] Building CXX object CMakeFiles/rocclr.dir/device/device.cpp.o
[ 20%] Building CXX object CMakeFiles/rocclr.dir/device/devkernel.cpp.o
[ 22%] Building CXX object CMakeFiles/rocclr.dir/device/devprogram.cpp.o
[ 24%] Building CXX object CMakeFiles/rocclr.dir/device/devwavelimiter.cpp.o
[ 26%] Building CXX object CMakeFiles/rocclr.dir/device/hsailctx.cpp.o
[ 28%] Building CXX object CMakeFiles/rocclr.dir/device/hwdebug.cpp.o
[ 30%] Building CXX object CMakeFiles/rocclr.dir/elf/elf.cpp.o
[ 32%] Building CXX object CMakeFiles/rocclr.dir/os/alloc.cpp.o
[ 34%] Building CXX object CMakeFiles/rocclr.dir/os/os_posix.cpp.o
[ 36%] Building CXX object CMakeFiles/rocclr.dir/os/os_win32.cpp.o
[ 38%] Building CXX object CMakeFiles/rocclr.dir/os/os.cpp.o
[ 40%] Building CXX object CMakeFiles/rocclr.dir/platform/activity.cpp.o
[ 42%] Building CXX object CMakeFiles/rocclr.dir/platform/agent.cpp.o
[ 44%] Building CXX object CMakeFiles/rocclr.dir/platform/command.cpp.o
[ 46%] Building CXX object CMakeFiles/rocclr.dir/platform/commandqueue.cpp.o
[ 48%] Building CXX object CMakeFiles/rocclr.dir/platform/context.cpp.o
[ 51%] Building CXX object CMakeFiles/rocclr.dir/platform/kernel.cpp.o
[ 53%] Building CXX object CMakeFiles/rocclr.dir/platform/memory.cpp.o
[ 55%] Building CXX object CMakeFiles/rocclr.dir/platform/ndrange.cpp.o
[ 57%] Building CXX object CMakeFiles/rocclr.dir/platform/program.cpp.o
[ 59%] Building CXX object CMakeFiles/rocclr.dir/platform/runtime.cpp.o
[ 61%] Building CXX object CMakeFiles/rocclr.dir/thread/monitor.cpp.o
[ 63%] Building CXX object CMakeFiles/rocclr.dir/thread/semaphore.cpp.o
[ 65%] Building CXX object CMakeFiles/rocclr.dir/thread/thread.cpp.o
[ 67%] Building CXX object CMakeFiles/rocclr.dir/utils/debug.cpp.o
[ 71%] Building CXX object CMakeFiles/rocclr.dir/utils/flags.cpp.o
[ 71%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocappprofile.cpp.o
[ 73%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocblit.cpp.o
[ 75%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/roccounters.cpp.o
[ 77%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocdevice.cpp.o
[ 79%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocglinterop.cpp.o
[ 81%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rockernel.cpp.o
[ 83%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocmemory.cpp.o
[ 85%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocprintf.cpp.o
[ 87%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocprogram.cpp.o
[ 89%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocschedcl.cpp.o
[ 91%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocsettings.cpp.o
[ 93%] Building CXX object CMakeFiles
```/rocclr.dir/device/rocm/rocsignal.cpp.o
[ 95%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocvirtual.cpp.o
[ 97%] Building CXX object CMakeFiles/rocclr.dir/device/rocm/rocurilocator.cpp.o
[100%] Linking CXX static library librocclr.a
[100%] Built target rocclr
root@localhost:~/ROCm-4.5/ROCclr/build# make install
make: *** No rule to make target 'install'.  Stop.

Installation documentation appears to be out-of-date

The build documentation currently suggests that this package can be built in a standalone manner. However, since df14496 it appears that this configuration is no longer supported and that it should rather be built as a component of some other unspecified project.

This puts packagers in the difficult situation of knowing that ROCclr is needed yet having no idea of how to actually package it.

Do librocclr really need libamd_comgr?

Related to ROCm/hipamd#44

I'm maintaining hip and rocm-opencl-runtime package in Gentoo, and recently, I find although in wirtten cmake, rocclr is linked to libamd_comgr, in reality linker dropped it since there is not function from libamd_comgr called.

I browsed ROCclr source codes by ag amd::Comgr:: | grep -o -E "amd::Comgr::[^\(]*\(" | sort -u, finding these functions are defined in amd_comgr.h, not comgr .cpp files, so librocclr does not have to link libamd_comgr. Is this just a coincidence, or a design? For the later one, I can make hip and rocm-opencl-runtime independent of rocm-comgr, since some tests of rocm-comgr are only available when hip is installed.

Invalid buffer reallocation

It appears that hostAlloc() has been accidentally discarded here and we end up assigning sizeof(void*) to dbgBuffer_:
https://github.com/ROCm-Developer-Tools/ROCclr/blob/928f8b064ddbd3e24a9ed6973da1f4a6adefba1a/device/rocm/rocprintf.cpp#L50

4.0: basename_max

Running clinfo with version 4.0, after the clinfo output, I get a core dump:

LoadLib(libhsa-amd-aqlprofile64.so) failed: libhsa-amd-aqlprofile64.so: cannot open shared object file: No such file or directory
Number of platforms                               1
  Platform Name                                   AMD Accelerated Parallel Processing
  Platform Vendor                                 Advanced Micro Devices, Inc.
  Platform Version                                OpenCL 2.0 AMD-APP.dbg (3212.0)
  Platform Profile                                FULL_PROFILE
  Platform Extensions                             cl_khr_icd cl_amd_event_callback 
  Platform Extensions function suffix             AMD

  Platform Name                                   AMD Accelerated Parallel Processing
Number of devices                                 1
  Device Name                                     gfx803
  Device Vendor                                   Advanced Micro Devices, Inc.
  Device Vendor ID                                0x1002
  Device Version                                  OpenCL 1.2 
  Driver Version                                  3212.0 (HSA1.1,LC)
  Device OpenCL C Version                         OpenCL C 2.0 
  Device Type                                     GPU
  Device Board Name (AMD)                         Ellesmere [Radeon RX 470/480/570/570X/580/580X/590]
  Device PCI-e ID (AMD)                           0x67df
  Device Topology (AMD)                           PCI-E, 01:00.0
  Device Profile                                  FULL_PROFILE
  Device Available                                Yes
  Compiler Available                              Yes
  Linker Available                                Yes
  Max compute units                               36
  SIMD per compute unit (AMD)                     4
  SIMD width (AMD)                                16
  SIMD instruction width (AMD)                    1
  Max clock frequency                             1380MHz
  Graphics IP (AMD)                               8.0
  Device Partition                                (core)
    Max number of sub-devices                     36
    Supported partition types                     None
    Supported affinity domains                    (n/a)
  Max work item dimensions                        3
  Max work item sizes                             1024x1024x1024
  Max work group size                             256
  Preferred work group size (AMD)                 256
  Max work group size (AMD)                       1024
clinfo: /var/tmp/portage/dev-libs/rocclr-4.0.0/work/ROCclr-rocm-4.0.0/compiler/lib/utils/options.cpp:1568: void amd::option::Options::setDumpFileName(const char*): Assertion `basename_max != -1' failed.
Aborted (core dumped)

Here's the line with the failed assertion: https://github.com/ROCm-Developer-Tools/ROCclr/blob/rocm-4.0.0/compiler/lib/utils/options.cpp#L1568

Debian packaging: typo patch

Dear all, please accept this typo patch.
Best regards, Maxime

This is a matter in your clr directory of :
patch < file.patch

clEnqueueSVMMemcpy SegFault

Running clEnqueueSVMMemcpy(queue, CL_TRUE, dst, src, size, 0, NULL, NULL); with dst allocated with clSVMAlloc and src allocated by the system (e.g. posix_memalign) triggers a segmentation fault:

Thread 4 "Command Queue T" received signal SIGSEGV, Segmentation fault.

Backtrace:

#0  0x00007fffebed61b4 in amd::SharedReference<amd::Context>::operator() (this=0x68) at /space/rocm/ROCclr/platform/object.hpp:166
#1  0x00007fffebed1f1e in amd::Memory::getContext (this=0x0) at /space/rocm/ROCclr/platform/memory.hpp:302
#2  0x00007fffebfeb8e3 in roc::NullDevice::forceFineGrain (this=0x55555568f820, memory=0x0) at /space/rocm/ROCclr/device/rocm/rocdevice.hpp:194
#3  0x00007fffebfe04e0 in roc::VirtualGPU::submitSvmCopyMemory (this=0x7ffed8000b90, cmd=...) at /space/rocm/ROCclr/device/rocm/rocvirtual.cpp:1281
#4  0x00007fffebf00cd0 in amd::SvmCopyMemoryCommand::submit(device::VirtualDevice&) () from /space/rocm/ROCm-OpenCL-Runtime/build/lib/libamdocl64.so
#5  0x00007fffebf8f4e2 in amd::HostQueue::loop (this=0x5555555d1ce0, virtualDevice=0x7ffed8000b90) at /space/rocm/ROCclr/platform/commandqueue.cpp:167
#6  0x00007fffebf9251e in amd::HostQueue::Thread::run (this=0x5555555d1d88, data=0x5555555d1ce0) at /space/rocm/ROCclr/platform/commandqueue.hpp:161
#7  0x00007fffebf4a4bd in amd::Thread::main (this=0x5555555d1d88) at /space/rocm/ROCclr/thread/thread.cpp:93
#8  0x00007fffebf984a4 in amd::Thread::entry (thread=0x5555555d1d88) at /space/rocm/ROCclr/os/os_posix.cpp:318
#9  0x00007ffff6bc2609 in start_thread (arg=<optimized out>) at pthread_create.c:477
#10 0x00007ffff71cb103 in clone () at ../sysdeps/unix/sysv/linux/x86_64/clone.S:95

The problem are the checks in line 1281 and 1282, which will both trigger a segfault if srcMem or dstMem is a nullptr, that is, not present in the memory map (line 1259 and 1260):
https://github.com/ROCm-Developer-Tools/ROCclr/blob/roc-3.5.x/device/rocm/rocvirtual.cpp#L1281-L1282
https://github.com/ROCm-Developer-Tools/ROCclr/blob/roc-3.5.x/device/rocm/rocvirtual.cpp#L1259-L1260
This will call forceFineGrain with a nullptr instead of amd::Memory*:
https://github.com/ROCm-Developer-Tools/ROCclr/blob/roc-3.5.x/device/rocm/rocdevice.hpp#L193

Tested on lates ROCm 3.5.1 release with the coresponding roc-3.5.x or rocm-3.5.x branches.

OpenCL ICD Versions

I am having trouble building ROCm-OpenCL-Runtime due to what seems to be an inconsistency in the circular dependency between that package and this one. ROCclr forward declares KHRicdVendorDispatchRec here, but ROCclr relies upon the ROCm-OpenCL-Runtime headers which no longer provide that definition. Older Khronos ICDs do, but then it is not clear what ICD to build ROCm-OpenCL-Runtime against.

assert("string") never happens because its precondition is always true.

https://github.com/ROCm-Developer-Tools/ROCclr/blob/ca35d38e98c3bf91da23707820060d1f68dd6e95/device/devprogram.cpp#L2632

This should've been assert(false && "message")

Is it true that one stream can submit work to multiple different HSA queues?

In an article Exploring AMD GPU scheduling details by experimenting with “worst practices” (https://link.springer.com/article/10.1007/s11241-022-09381-y) there is a statement:
ROCclr’s software queues internally share a pool of HSA queues: one software queue may submit work to multiple diferent HSA queues, and each HSA queue may contain work from multiple software queues.
There is no doubt that one HSA queue can receive work from multiple different streams, but according to the relative source code, it seems that with the DIRECT DISPATCH feature since rocm 4.5, one stream may only submit to one HSA queue that it is bound to. I want to know if I understand the relative code correctly. Thanks in advance!

[Build error] Unable to find OpenCL includes

I'm using Arch Linux with cmake 3.18 and I'm unable to build ROCclr with the current build instructions / CMakeLists.txt. The problem seems to be that find_library also looks for a library in /opt/rocm which is not present of course as ROCm OpenCL Runtime needs ROCclr. Moreover, ROCclr only needs the include files from OpenCL, not the full package.
I successfully build ROCclr with the patched CMakeLists.txt from my PR #16 .

Compling ROCclr with PAL backend: Closed source elements

Hi,
I am trying to compile Roclr with PAL (BUILD_PAL cmake flag set) in order to get ROCm working on Azure instances.
Unfortunately, it seems not everything is publicly available: headers such as glATIInternal.h, a closed source version of PAL with extented structures/classes or additional enum values (Navi10_A0, Polaris22, ...).

Is there truly closed source elements or didn't I search enough ?

Anyway, is there a way to get a half-featured ROCM working on a VM with very limited PCI features ? I focused on the PAL way because it seemed to me it could have worked anywhere a Vulkan stack could run some shaders.

HSA_AMD_SVM_ATTRIB_READ_MOSTLY missing in ROCR-Runtime headers.

Hi,
I can't find the definition of HSA_AMD_SVM_ATTRIB_READ_MOSTLY in the current ROCR-Runtime repo.
Where is it ?
Thank you

hipMemcpy with hipamd pulls in full Comgr support

Even compiling the simple square example, this obviously needs to use hipMemcpy.

However, I was kind of surprised that internally in hipMemcpy the BlitProgram is used which (in an OpenCL fashion) seems to be built on first use using the compiler manager, and thus requiring effectively a full llvm dependency. In particular the device libs need to be present (which was what made me even aware of this issue)

Doesn't this undermine the whole point of doing binary builds for each ISA in hipamd in the first place?

To reproduce the issue run

AMD_COMGR_SAVE_TEMPS=1 AMD_COMGR_REDIRECT_LOGS=stdout AMD_COMGR_EMIT_VERBOSE_LOGS=1 AMD_LOG_LEVEL=7 ./square

Ideally this should only happen if hiprtc is used.

device memory accounting not correct

After I found #35 I checked where the memory update function is used.
I found places where the accounting is performed after the memory allocation happened.
This is introducing incorrect memory accounting.
If one thread is allocating memory and another thread is checking the free memory it is possible that the following data race is happening and the amount of free memory reported is wrong.

Case:

• Thread 0 is allocating memory but has not reached the function call updateFreeMemory() e.g. here
• Thread 1 is querying the amount of free memory, result will be the amount of free memory before Thread 0 allocated memory because the accounting was not performed
• Thread 0 is calling updateFreeMemory() -> now the internal free memory counter is correct but Thread 1 got the wrong amount of free memory

This issue is not very critical because in a parallel environment the amount of free memory could always be outdated of concurrent threads de/allocating memory.
Never the less IMO the account should be performed before the memory is allocated and in cases where the allocation failed the accounting should be rolled back.

Code snippets where the accounting is called after the real allocation:
https://github.com/ROCm-Developer-Tools/ROCclr/blob/6a62060c30cf910f07c2b98015def633daa7bf62/device/rocm/rocmemory.cpp#L809-L812
https://github.com/ROCm-Developer-Tools/ROCclr/blob/6a62060c30cf910f07c2b98015def633daa7bf62/device/rocm/rocmemory.cpp#L886-L892
https://github.com/ROCm-Developer-Tools/ROCclr/blob/6a62060c30cf910f07c2b98015def633daa7bf62/device/rocm/rocmemory.cpp#L1189-L1200

AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES is deprecated

...and a very long macro name 😂

This might have been overlooked previously since the warning wasn't properly printed when building with Clang before ROCm/ROCm-CompilerSupport@55deecf.

This also affects a few similar symbols. Warnings look like this:

external/rules_ll~override~rules_ll_dependencies~comgr/lib/comgr/src/comgr.cpp:209:8: warning: 'AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES' is deprecated: Will be removed in Comgr
 v3.0 (Rocm v6.0). Use AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC instead [-Wdeprecated-declarations]
  case AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES:
       ^
bazel-out/k8-fastbuild/bin/external/rules_ll~override~rules_ll_dependencies~comgr/comgr/amd_comgr/amd_comgr.h:1560:3: note: 'AMD_COMGR_ACTION_ADD_DEVICE_LIBRARIES' has been e
xplicitly marked deprecated here
  AMD_COMGR_DEPRECATED("Will be removed in Comgr v3.0 (Rocm v6.0). Use "
  ^
bazel-out/k8-fastbuild/bin/external/rules_ll~override~rules_ll_dependencies~comgr/comgr/amd_comgr/amd_comgr.h:55:50: note: expanded from macro 'AMD_COMGR_DEPRECATED'
#define AMD_COMGR_DEPRECATED(msg) __attribute__((deprecated(msg)))
                                                 ^
external/rules_ll~override~rules_ll_dependencies~comgr/lib/comgr/src/comgr.cpp:231:8: warning: 'AMD_COMGR_ACTION_COMPILE_SOURCE_TO_FATBIN' is deprecated: Will be removed in C
omgr v3.0 (Rocm v6.0). Use AMD_COMGR_ACTION_COMPILE_SOURCE_TO_BC, etc. instead [-Wdeprecated-declarations]
  case AMD_COMGR_ACTION_COMPILE_SOURCE_TO_FATBIN:
       ^
bazel-out/k8-fastbuild/bin/external/rules_ll~override~rules_ll_dependencies~comgr/comgr/amd_comgr/amd_comgr.h:1711:3: note: 'AMD_COMGR_ACTION_COMPILE_SOURCE_TO_FATBIN' has be
en explicitly marked deprecated here
  AMD_COMGR_DEPRECATED("Will be removed in Comgr v3.0 (Rocm v6.0). Use "
  ^

I tried naively updating the macros but that seems to break everything, so I guess we need the expertise of someone who actually knows what they are doing here 😇

cc @lamb-j