[all-commits] [llvm/llvm-project] f616c3: [OpenMP][DeviceRTL][AMDGPU] Support code object ve...

[Saiyedul Islam via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: f616c3eeb43f3732f53f81d291723a6a34af2de1
      https://github.com/llvm/llvm-project/commit/f616c3eeb43f3732f53f81d291723a6a34af2de1
  Author: Saiyedul Islam <Saiyedul.Islam at amd.com>
  Date:   2023-08-29 (Tue, 29 Aug 2023)

  Changed paths:
    M clang/lib/CodeGen/CGBuiltin.cpp
    M clang/lib/CodeGen/CodeGenModule.cpp
    M clang/lib/CodeGen/CodeGenModule.h
    M clang/lib/CodeGen/TargetInfo.h
    M clang/lib/CodeGen/Targets/AMDGPU.cpp
    M clang/lib/Driver/ToolChain.cpp
    M clang/lib/Driver/ToolChains/Clang.cpp
    A clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu
    M clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
    M clang/test/CodeGenOpenCL/opencl_types.cl
    M clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
    M openmp/libomptarget/DeviceRTL/CMakeLists.txt
    M openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp
    M openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

  Log Message:
  -----------
  [OpenMP][DeviceRTL][AMDGPU] Support code object version 5

Update DeviceRTL and the AMDGPU plugin to support code
object version 5. Default is code object version 4.

CodeGen for __builtin_amdgpu_workgroup_size generates code
for cov4 as well as cov5 if -mcode-object-version=none
is specified. DeviceRTL compilation passes this argument
via Xclang option to generate abi-agnostic code.

Generated code for the above builtin uses a clang
control constant "llvm.amdgcn.abi.version" to branch on
the abi version, which is available during linking of
user's OpenMP code. Load of this constant gets eliminated
during linking.

AMDGPU plugin queries the ELF for code object version
and then prepares various implicitargs accordingly.

Differential Revision: https://reviews.llvm.org/D139730

Reviewed By: jhuber6, yaxunl