rocclr's Issues
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?
(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;
}
}
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.
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.
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
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.
assert("string") never happens because its precondition is always true.
This should've been assert(false && "message")
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
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
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.
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
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
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.
- 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
- Two threads both allocating 2MiB going into line 2091 and the check is false because
- 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
- Two threads where one is deallocating 2MiB and the other is allocating 5MiB and
Possible solution:
- When memory is freed
freeMem_
is stored in a register and atomic CAS is used to reset the variablefreeMem_
- 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 ofCase 1
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
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
`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.
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.
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
Using textures causes "Depth set on single layer image geometry."
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
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.
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
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 .
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.
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.
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.
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.
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.
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.
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.