From f616c3eeb43f3732f53f81d291723a6a34af2de1 Mon Sep 17 00:00:00 2001 From: Saiyedul Islam Date: Thu, 17 Aug 2023 07:45:40 -0500 Subject: [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 --- clang/lib/CodeGen/CodeGenModule.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) (limited to 'clang/lib/CodeGen/CodeGenModule.h') diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index ce14a6c..d4032aa 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1571,6 +1571,11 @@ public: void handleAMDGPUWavesPerEUAttr(llvm::Function *F, const AMDGPUWavesPerEUAttr *A); + llvm::Constant * + GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty, LangAS AddrSpace, + const VarDecl *D, + ForDefinition_t IsForDefinition = NotForDefinition); + private: llvm::Constant *GetOrCreateLLVMFunction( StringRef MangledName, llvm::Type *Ty, GlobalDecl D, bool ForVTable, @@ -1593,11 +1598,6 @@ private: void UpdateMultiVersionNames(GlobalDecl GD, const FunctionDecl *FD, StringRef &CurName); - llvm::Constant * - GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty, LangAS AddrSpace, - const VarDecl *D, - ForDefinition_t IsForDefinition = NotForDefinition); - bool GetCPUAndFeaturesAttributes(GlobalDecl GD, llvm::AttrBuilder &AttrBuilder, bool SetTargetFeatures = true); -- cgit v1.1