Compile with -fPIC

[stuartarchibald]

Following on from: #4

I am trying to do the following:

• Create a dataset with a single LLVM IR source and 9 bitcode sources:

  • LLVM IR is generated by Numba

  • Bitcodes are from the ROCm 1.9.x release:

    • "opencl.amdgcn.bc",
    • "ocml.amdgcn.bc",
    • "ockl.amdgcn.bc",
    • "oclc_correctly_rounded_sqrt_off.amdgcn.bc",
    • "oclc_daz_opt_off.amdgcn.bc",
    • "oclc_finite_only_off.amdgcn.bc",
    • "oclc_isa_version_803.amdgcn.bc",
    • "oclc_unsafe_math_off.amdgcn.bc",
    • "irif.amdgcn.bc
• Perform an action kind of AMD_COMGR_ACTION_ASSEMBLE_SOURCE_TO_RELOCATABLE on the dataset
• Perform an action kind of AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_RELOCATABLE on the resulting dataset

• Perform an action kind of AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE on the resulting dataset. At which point 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)?

Thanks.

[scott-linder]

Are the actions you mention the only ones performed? The AMD_COMGR_ACTION_ASSEMBLE_SOURCE_TO_RELOCATABLE action operates on ISA-level assembly text (e.g. .S files) not on LLVM IR/BC. This is related to #7; "source" is not very well defined in the documentation, but it does not include LLVM IR.

If you are trying to create a relocatable from LLVM BC sources, you will want to use AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE.

[lamb-j]

Should be fixed by updating Comgr actions used. Let me know if you still have any issues with this!