Closed stuartarchibald closed 1 year ago
Do you have a minimal example? Currently the only output comes from tools which do not expose options to control when output is flushed, but as part of #2 we should replace much of this output to standard streams with LOG objects.
Unfortunately not. I have a tower of half working things whilst switching out the Numba ROC target compilation pipeline to use comgr
. The place I observe this is when I do the following:
Create a dataset with a single LLVM IR source and 9 bitcode sources:
Bitcodes are from the ROCm 1.9.x release:
AMD_COMGR_ACTION_ASSEMBLE_SOURCE_TO_RELOCATABLE
on the datasetAMD_COMGR_ACTION_LINK_RELOCATABLE_TO_RELOCATABLE
on the resulting datasetPerform an action kind of AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE
on the resulting dataset. At which point about 5 screens of blank flushing takes place. Scrolling up, there's loads of errors like:
: error: relocation R_AMDGPU_REL32_LO cannot be used against symbol get_local_id(unsigned int); recompile with -fPIC
>>> defined in /tmp/comgr-7102b3/input/numba_generated_ir.ll.o
>>> referenced by /tmp/comgr-7102b3/input/numba_generated_ir.ll.o:(hsapy_devfn_numba_2E_roc_2E_tests_2E_hsapy_2E_test_5F_scan_2E_shuffle_5F_up_24_2_2E_int32_2E_int64)
I'm guessing the error is because of the baked in -r
in the lld
binding there's no way to -shared
. Also can't find a way to add -fPIC
(which was a bit what this was about https://github.com/RadeonOpenCompute/ROCm-CompilerSupport/issues/5, in that I'm trying to relate the comgr
behaviours and source back to a familiar implementation). I can open the above link problem as a new issue if it's best resolved there.
Thanks for your help.
I'm guessing this is no longer an issue. I haven't seen or heard of this error recently. Follow up if you're still able to recreate it though!
If the library hits some problem during an action, e.g. cannot relocate symbols, it's fairly common for the output to get repeatedly flushed with large amounts of out empty output, technically this isn't a problem, but if preventable it'd make reading the output more convenient. Thanks.