[OpenMP][AMDGCN] Generate global variables and attributes for AMDGCN

[saiislam]

Summary

Provide support for amdgcn specific global variables and attributes.
Generalize allocation of various common global variables and provide
their specialized implementations for nvptx and amdgcn.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.h

499

Cab this type and corresponding functions be made AMDGCN-specific only?

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.cpp

116

Is this possible?

119

FlatAttrEmitted

129

CompileTimeThreadLimit

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

1390–1392

Restore original formatting.

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

215–243

Make them protected, not public if possible. Try the same for other new functions.

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h

40–65

No need to add virtual, override is enough

[saiislam]

1. Moved amdgcn specific functions to CGOpenMPAMDGCN.cpp
2. Removed tautology condition
3. Corrected case of local variables
4. Restored original formatting
5. Changed back declaration of emit kernel methods as private
6. Added support of amdgcn specific PrePostActionTy implementation and its corresponding test cases
7. Changed static line numbers in new test cases with regex
8. Other small code corrections

[ABataev]

Reformat the code

clang/lib/CodeGen/CGOpenMPRuntime.h

498

Remove unnecessary formatting changes.

2489–2493

Better to make it a protected member function if you really require it. Plus, this function is very small and, I think, you simply create your own copy in CGOpenMPRuntimeAMDGCN

2497

Same here, make it protected or just create a copy, if it is small.

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.h

29–31

Add comments for all new members

85

1. Do you really need to make this class public?
2. final

[saiislam]

1. Reformarting
2. Comments
3. Reduced scope of specialized PrePostActionTy

[saiislam]

clang/lib/CodeGen/CGOpenMPRuntime.h

2489–2493

Not making it protected because it is used by various static functions. And don't want to create an object pointer of subclass of CGOpenMPRuntime in CGOpenMPRuntime.

2497

It calls static functions which in turn call other static functions, so it won't make sense to create a copy of whole function chain in amdgcn.

[saiislam]

Ping.

[saiislam]

Ping.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.h

498

Still not removed

684

Restore original formatting

2492–2497

Better to encapsulate these functions into a new utility class and make them public static.

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.h

34–67

Do you really need to expose all these new members as public?

35

Runtime does not support nested parallelism on GPU. Do you really need it?

93

It does not help to understand the functionality

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

1972

It leads to a mem leak.

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

37

1. Make it private or protected

2.Add default initializer

187–219

Are all these required to be public?

410–412

Make it private or protected

[saiislam]

1. Removed unnecessary formatting of untouched code.
2. Encapsulated addFieldToRecordDecl and createGlobalStruct methods in a class and made them static (triggered change at all calling sites).
3. Marked most of the member methods of CGOpenMPRuntimeAMDGCN as private (forgot to do same change in nvptx)
4. Fixed the memory leak
5. Marked appropriate member variables as protected in CGOpenMPRuntimeGPU

[JonChesterfield]

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.cpp

175

The nvptx emitSPMDKernelWrapper does nothing and the amdgcn one appends some metadata.

How about 'nvptx::generateMetadata(...)' that does nothing and 'amdgcn::generateMetadata(...)` that does this stuff, called from the end of emitSPMDKernel?

197

This metadata generation could be split out from the other changes.

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.h

43

I'm not convinced by this abstraction. It looks like amdgcn and nvptx want almost exactly the same variable in each case. The difference appears to be that nvptx uses internal linkage and amdgcn uses weak + externally initialized, in which case we're better off with

bool nvptx::needsExternalInitialization() {return false;}
bool amdgpu::needsExternalInitialization() {return true;}

Or, if the inline ternary is unappealing, amdgcn::NewGlobalVariable(...) that passes the arguments to llvm::GlobalVariable while setting the two fields that differ between the two.

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

170

Please put this back to the previous location so we can see whether it changed in the diff

[saiislam]

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.cpp

175

It will be then difficult to track what all things are being done differently in the two. So, the common code has been generalized and (no change in nvptx + some changes in amdgcn) has been used as specialization.

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.h

43

I understand what you are suggesting. But, there are multiple such variables where linkage between nvptx and amdgcn are different. Also current style gives flexibility to a future implementation to define these variables in their own way.
What do you think?

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

170

This movement changes them from private to protected.
I could have just added access specifiers and not move the definitions. It would have simplified the review, but it would have decreased the readability for future.

187–219

Yes, they are being called from outside class.

[JonChesterfield]

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

63

Perhaps (typed into browser):

llvm::GlobalVariable *CGOpenMPRuntimeNVPTX::createGlobal( CodeGenModule &CGM, llvm::ArrayType *Ty, StringRef Name) {
  return new llvm::GlobalVariable(
      CGM.getModule(), Ty, /*isConstant=*/false,
      llvm::GlobalVariable::CommonLinkage, llvm::Constant::getNullValue(Ty),
      Name,
      /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared),
/*isExternallyInitialized*/ true);
}

llvm::GlobalVariable *CGOpenMPRuntimeAMDGCN::createGlobal( CodeGenModule &CGM, llvm::ArrayType *Ty, StringRef Name) {
  return new llvm::GlobalVariable(
      CGM.getModule(), Ty, /*isConstant=*/false,
      llvm::GlobalVariable::WeakAnyLinkage, llvm::Constant::getNullValue(Ty),
      Name,
      /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared),
/*isExternallyInitialized*/ false);
}

[saiislam]

1. Simplifies overall patch after D90248.
2. Removes MaxParallelLevel and thus target specific PrePostActionTy.
3. Removes ExternallyInitialized qualifier from shared variables for AMDGCN.

[JonChesterfield]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

1355

This appears to be the same as the free function we had before, except now all the call sites are prefixed CodegenUtil. Is there a functional change I'm missing?

The rest of this patch would be easier to read with this part split off.

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.cpp

74

This is a very verbose way to say that amdgcn calls emitmetatdata at the end of emitkernel and nvptx doesn't. Suggest unconditionally calling emitmetatdata, and having emitmetatdata be a no-op for nvptx.

86

I think there's a credible chance this is useful to nvptx, so doesn't have to be amdgcn specific

105

I think this is about computing a maximum workgroup size which the runtime uses to limit the number of threads it launches. If so, this is probably useful for nvptx and amdgcn. I'm having trouble working out what the conditions are though. Maybe it's based on an openmp clause?

147

I think I remember seeing a diff that makes this attribute unconditionally emitted by some other part of the toolchain. If so, it may no longer be required

166

HostServices is unused. Mode is redundant with exec_mode. wg_size is redundant with the other wg_size symbol added above. This kern_desc object should be deleted, not upstreamed.

[saiislam]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

1355

addFieldToRecordDecl and createGlobalStruct methods had file static scope. To make them callable from other files, from amdgcn specific file in this case, they were put in this utility class.

D92167 puts this change into a separate patch. Will update this patch once D92167 gets accepted.

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.cpp

74

Won't the no-op approach be less extensible? Current way, though verbose, leaves scope for attaching prefix/suffix code as and when required around emitkernel. While in case of no-op, every implementing arch might have to use the exact same pattern of methods with and without code.

86

You are right, it can be useful for nvptx as well. May be we can club its generalization with the nvptx's use-case when it arrives in the future?

105

Yes, the if block in 111-147 corresponds to "number of threads" for thread_limit and num_threads clauses in teams and parallel directives.

166

Ok, thanks. Will update in next revision.

[JonChesterfield]

I don't believe the contents of this patch is necessary for codegen on amdgpu. One of the internal/weak distinctions works around a bug in the gfx800 toolchain, but we should root cause and fix that bug instead. The kern_desc object is redundant. I think amdgpu-flat-work-group-size is already emitted, but if not, we might want that.

The wg_size code is interesting but architecture independent, and it's probably more user friendly for nvptx and amdgcn to have the same handling of wg_size constraints.

clang/lib/CodeGen/CGOpenMPRuntime.h

private:
   ///   char      *name;       // Name of the function or global.
   ///   size_t     size;       // Size of the entry info (0 if it a function).
   ///   int32_t flags;
   ///   int32_t reserved;
   /// };
   QualType TgtOffloadEntryQTy;
   /// Entity that registers the offloading constants that were emitted so
   /// far.
+
+  QualType TgtAttributeStructQTy;
   class OffloadEntriesInfoManagerTy {
     CodeGenModule &CGM;
 
     /// Number of entries registered so far.
     unsigned OffloadingEntriesNum = 0;
 
   public:
     /// Base class of the entries info.
bool hasDeviceGlobalVarEntryInfo(StringRef VarName) const {
       return OffloadEntriesDeviceGlobalVar.count(VarName) > 0;
     }
     /// Applies action \a Action on all registered entries.
     typedef llvm::function_ref<void(StringRef,
                                     const OffloadEntryInfoDeviceGlobalVar &)>
         OffloadDeviceGlobalVarEntryInfoActTy;
     void actOnDeviceGlobalVarEntriesInfo(
         const OffloadDeviceGlobalVarEntryInfoActTy &Action);
 
   private:
     // Storage for target region entries kind. The storage is to be indexed by
     // file ID, device ID, parent function name and line number.
     typedef llvm::DenseMap<unsigned, OffloadEntryInfoTargetRegion>
         OffloadEntriesTargetRegionPerLine;
     typedef llvm::StringMap<OffloadEntriesTargetRegionPerLine>
         OffloadEntriesTargetRegionPerParentName;
     typedef llvm::DenseMap<unsigned, OffloadEntriesTargetRegionPerParentName>
public:
 
   /// Emits call of the outlined function with the provided arguments,
   /// translating these arguments to correct target-specific arguments.
   virtual void
   emitOutlinedFunctionCall(CodeGenFunction &CGF, SourceLocation Loc,
                            llvm::FunctionCallee OutlinedFn,
                            ArrayRef<llvm::Value *> Args = llvm::None) const;
 
+  /// Returns __tgt_attribute_struct type.
+  QualType getTgtAttributeStructQTy();
+
+  /// Emit structure descriptor for a kernel
+  void emitStructureKernelDesc(CodeGenModule &CGM, StringRef Name,
+                               int16_t WG_Size, int8_t Mode,
+                               int8_t HostServices, int8_t MaxParallelLevel);
+
   /// Emits OpenMP-specific function prolog.
   /// Required for device constructs.
   virtual void emitFunctionProlog(CodeGenFunction &CGF, const Decl *D);
 
   /// Gets the OpenMP-specific address of the local variable.
   virtual Address getAddressOfLocalVariable(CodeGenFunction &CGF,
                                             const VarDecl *VD);
 
public:
 
   /// Gets the OpenMP-specific address of the local variable.
   Address getAddressOfLocalVariable(CodeGenFunction &CGF,
                                     const VarDecl *VD) override {
     return Address::invalid();
   }
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif

clang/lib/CodeGen/CGOpenMPRuntime.cpp

for (const FieldDecl *FD : RD->fields()) {
     for (unsigned I = PrevIdx; I < Idx; ++I)
       Fields.add(llvm::Constant::getNullValue(StructTy->getElementType(I)));
     PrevIdx = Idx + 1;
     Fields.add(*DI);
     ++DI;
   }
 }
 
 template <class... As>
 static llvm::GlobalVariable *
 createGlobalStruct(CodeGenModule &CGM, QualType Ty, bool IsConstant,
                    ArrayRef<llvm::Constant *> Data, const Twine &Name,
                    As &&... Args) {
   const auto *RD = cast<RecordDecl>(Ty->getAsTagDecl());
   const CGRecordLayout &RL = CGM.getTypes().getCGRecordLayout(RD);
   ConstantInitBuilder CIBuilder(CGM);
   ConstantStructBuilder Fields = CIBuilder.beginStruct(RL.getLLVMType());
if (!KmpRoutineEntryPtrTy) {
     QualType KmpRoutineEntryTyArgs[] = {KmpInt32Ty, C.VoidPtrTy};
     FunctionProtoType::ExtProtoInfo EPI;
     KmpRoutineEntryPtrQTy = C.getPointerType(
         C.getFunctionType(KmpInt32Ty, KmpRoutineEntryTyArgs, EPI));
     KmpRoutineEntryPtrTy = CGM.getTypes().ConvertType(KmpRoutineEntryPtrQTy);
   }
 }
 
+/// Emit structure descriptor for a kernel
+void CGOpenMPRuntime::emitStructureKernelDesc(CodeGenModule &CGM,
+                                              StringRef Name, int16_t WG_Size,
+                                              int8_t Mode, int8_t HostServices,
+                                              int8_t MaxParallelLevel) {
+
+  // Create all device images
+  llvm::Constant *AttrData[] = {
+      llvm::ConstantInt::get(CGM.Int16Ty, 2), // Version
+      llvm::ConstantInt::get(CGM.Int16Ty, 9), // Size in bytes
+      llvm::ConstantInt::get(CGM.Int16Ty, WG_Size),
+      llvm::ConstantInt::get(CGM.Int8Ty, Mode), // 0 => SPMD, 1 => GENERIC
+      llvm::ConstantInt::get(CGM.Int8Ty, HostServices), // 1 => use HostServices
+      llvm::ConstantInt::get(CGM.Int8Ty, MaxParallelLevel)}; // number of nests
+
+  llvm::GlobalVariable *AttrImages = createGlobalStruct(
+      CGM, getTgtAttributeStructQTy(), isDefaultLocationConstant(), AttrData,
+      Name + Twine("_kern_desc"), llvm::GlobalValue::WeakAnyLinkage);
+  CGM.addCompilerUsedGlobal(AttrImages);
+}
+
+// Create Tgt Attribute Sruct type.
+QualType CGOpenMPRuntime::getTgtAttributeStructQTy() {
+  ASTContext &C = CGM.getContext();
+  QualType KmpInt8Ty = C.getIntTypeForBitwidth(/*Width=*/8, /*Signed=*/1);
+  QualType KmpInt16Ty = C.getIntTypeForBitwidth(/*Width=*/16, /*Signed=*/1);
+  if (TgtAttributeStructQTy.isNull()) {
+    RecordDecl *RD = C.buildImplicitRecord("__tgt_attribute_struct");
+    RD->startDefinition();
+    addFieldToRecordDecl(C, RD, KmpInt16Ty); // Version
+    addFieldToRecordDecl(C, RD, KmpInt16Ty); // Struct Size in bytes.
+    addFieldToRecordDecl(C, RD, KmpInt16Ty); // WG_size
+    addFieldToRecordDecl(C, RD, KmpInt8Ty);  // Mode
+    addFieldToRecordDecl(C, RD, KmpInt8Ty);  // HostServices
+    addFieldToRecordDecl(C, RD, KmpInt8Ty);  // MaxParallelLevel
+    RD->completeDefinition();
+    TgtAttributeStructQTy = C.getRecordType(RD);
+  }
+  return TgtAttributeStructQTy;
+}
+
 QualType CGOpenMPRuntime::getTgtOffloadEntryQTy() {
   // Make sure the type of the entry is already created. This is the type we
   // have to create:
   // struct __tgt_offload_entry{
   //   void      *addr;       // Pointer to the offload entry info.
   //                          // (function or global)
   //   char      *name;       // Name of the function or global.
   //   size_t     size;       // Size of the entry info (0 if it a function).

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.h

 #include "clang/AST/StmtOpenMP.h"
 
 namespace clang {
 namespace CodeGen {
 
 class CGOpenMPRuntimeAMDGCN final : public CGOpenMPRuntimeGPU {
 
 public:
+  /// Nesting level of parallel region.
+  int ParallelLevel = 0;
+  int MaxParallelLevel = 0;
+
   explicit CGOpenMPRuntimeAMDGCN(CodeGenModule &CGM);
 
   /// Get the GPU warp size.
   llvm::Value *getGPUWarpSize(CodeGenFunction &CGF) override;
 
   /// Get the id of the current thread on the GPU.
   llvm::Value *getGPUThreadID(CodeGenFunction &CGF) override;
 
   /// Get the maximum number of threads in a block of the GPU.
   llvm::Value *getGPUNumThreads(CodeGenFunction &CGF) override;
+
+  /// Allocate global variable for TransferMedium
+  virtual llvm::GlobalVariable *
+  allocateTransferMediumGlobal(CodeGenModule &CGM, llvm::ArrayType *Ty,
+                               StringRef TransferMediumName) override;
+
+  /// Allocate global variable for SharedStaticRD
+  virtual llvm::GlobalVariable *
+  allocateSharedStaticRDGlobal(CodeGenModule &CGM,
+                               llvm::Type *LLVMStaticTy) override;
+
+  /// Get global variable KernelStaticGlobalized which is a shared pointer for
+  /// the global memory in the global memory buffer used for the given kernel
+  virtual llvm::GlobalVariable *
+  allocateKernelStaticGlobalized(CodeGenModule &CGM) override;
+
+  /// Emit target specifc SPMD kernel
+  virtual void emitSPMDKernelWrapper(const OMPExecutableDirective &D,
+                                     StringRef ParentName,
+                                     llvm::Function *&OutlinedFn,
+                                     llvm::Constant *&OutlinedFnID,
+                                     bool IsOffloadEntry,
+                                     const RegionCodeGenTy &CodeGen) override;
+
+  /// Emit target specific Non-SPMD kernel
+  virtual void
+  emitNonSPMDKernelWrapper(const OMPExecutableDirective &D,
+                           StringRef ParentName, llvm::Function *&OutlinedFn,
+                           llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
+                           const RegionCodeGenTy &CodeGen) override;
+
+  /// Create a unique global variable to indicate the flat-work-group-size
+  /// for this region. Values are [256..1024].
+  static void setPropertyWorkGroupSize(CodeGenModule &CGM, StringRef Name,
+                                       unsigned WGSize);
+
+  /// Generate global variables _wg_size, kern_desc, __tgt_attribute_struct.
+  /// Also generate appropriate value of attribute amdgpu-flat-work-group-size
+  void generateMetaData(CodeGenModule &CGM, const OMPExecutableDirective &D,
+                        llvm::Function *&OutlinedFn, bool IsGeneric);
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif // LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMEAMDGCN_H

clang/lib/CodeGen/CGOpenMPRuntimeAMDGCN.cpp

 //===-- CGOpenMPRuntimeAMDGCN.cpp - Interface to OpenMP AMDGCN Runtimes --===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This provides a class for OpenMP runtime code generation specialized to
 // AMDGCN targets from generalized CGOpenMPRuntimeGPU class.
 //
 //===----------------------------------------------------------------------===//
 
 #include "CGOpenMPRuntimeAMDGCN.h"
+#include "CGOpenMPRuntime.h"
 #include "CGOpenMPRuntimeGPU.h"
 #include "CodeGenFunction.h"
 #include "clang/AST/Attr.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/StmtVisitor.h"
 #include "clang/Basic/Cuda.h"
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/IR/IntrinsicsAMDGPU.h"
 
 using namespace clang;
 using namespace CodeGen;
 using namespace llvm::omp;
 
 CGOpenMPRuntimeAMDGCN::CGOpenMPRuntimeAMDGCN(CodeGenModule &CGM)
     : CGOpenMPRuntimeGPU(CGM) {
   if (!CGM.getLangOpts().OpenMPIsDevice)
     llvm_unreachable("OpenMP AMDGCN can only handle device code.");
+  StaticRDLinkage = llvm::GlobalValue::PrivateLinkage;
 }
 
 llvm::Value *CGOpenMPRuntimeAMDGCN::getGPUWarpSize(CodeGenFunction &CGF) {
   CGBuilderTy &Bld = CGF.Builder;
   // return constant compile-time target-specific warp size
   unsigned WarpSize = CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size);
   return Bld.getInt32(WarpSize);
 }
llvm::Value *CGOpenMPRuntimeAMDGCN::getGPUNumThreads(CodeGenFunction &CGF) {
   if (!F) {
     F = llvm::Function::Create(
         llvm::FunctionType::get(CGF.Int64Ty, {CGF.Int32Ty}, false),
         llvm::GlobalVariable::ExternalLinkage, LocSize, &CGF.CGM.getModule());
   }
   return Bld.CreateTrunc(
       Bld.CreateCall(F, {Bld.getInt32(0)}, "nvptx_num_threads"), CGF.Int32Ty);
 }
+
+llvm::GlobalVariable *CGOpenMPRuntimeAMDGCN::allocateTransferMediumGlobal(
+    CodeGenModule &CGM, llvm::ArrayType *Ty, StringRef TransferMediumName) {
+  return new llvm::GlobalVariable(
+      CGM.getModule(), Ty, /*isConstant=*/false,
+      llvm::GlobalVariable::WeakAnyLinkage, llvm::UndefValue::get(Ty),
+      TransferMediumName,
+      /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared),
+      /*isExternallyInitialized*/ true);
+}
+
+llvm::GlobalVariable *
+CGOpenMPRuntimeAMDGCN::allocateSharedStaticRDGlobal(CodeGenModule &CGM,
+                                                    llvm::Type *LLVMStaticTy) {
+  return new llvm::GlobalVariable(
+      CGM.getModule(), LLVMStaticTy,
+      /*isConstant=*/false, llvm::GlobalValue::WeakAnyLinkage,
+      llvm::UndefValue::get(LLVMStaticTy), "_openmp_shared_static_glob_rd_$_",
+      /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared),
+      /*isExternallyInitialized*/ true);
+}
+
+llvm::GlobalVariable *
+CGOpenMPRuntimeAMDGCN::allocateKernelStaticGlobalized(CodeGenModule &CGM) {
+  return new llvm::GlobalVariable(
+      CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
+      llvm::GlobalValue::WeakAnyLinkage, llvm::UndefValue::get(CGM.VoidPtrTy),
+      "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
+      llvm::GlobalValue::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared),
+      /*isExternallyInitialized*/ true);
+}
+
+void CGOpenMPRuntimeAMDGCN::setPropertyWorkGroupSize(CodeGenModule &CGM,
+                                                     StringRef Name,
+                                                     unsigned WGSize) {
+  auto *GVMode = new llvm::GlobalVariable(
+      CGM.getModule(), CGM.Int16Ty, /*isConstant=*/true,
+      llvm::GlobalValue::WeakAnyLinkage,
+      llvm::ConstantInt::get(CGM.Int16Ty, WGSize), Name + Twine("_wg_size"),
+      /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_device),
+      /*isExternallyInitialized*/ false);
+  CGM.addCompilerUsedGlobal(GVMode);
+}
+
+void CGOpenMPRuntimeAMDGCN::generateMetaData(CodeGenModule &CGM,
+                                             const OMPExecutableDirective &D,
+                                             llvm::Function *&OutlinedFn,
+                                             bool IsGeneric) {
+  if (!CGM.getTriple().isAMDGCN())
+    return;
+  int FlatAttr = 0;
+  bool flatAttrEmitted = false;
+  unsigned DefaultWorkGroupSz =
+      CGM.getTarget().getGridValue(llvm::omp::GVIDX::GV_Default_WG_Size);
+
+  if (isOpenMPTeamsDirective(D.getDirectiveKind()) ||
+      isOpenMPParallelDirective(D.getDirectiveKind())) {
+    const auto *ThreadLimitClause = D.getSingleClause<OMPThreadLimitClause>();
+    const auto *NumThreadsClause = D.getSingleClause<OMPNumThreadsClause>();
+    unsigned MaxWorkGroupSz =
+        CGM.getTarget().getGridValue(llvm::omp::GVIDX::GV_Max_WG_Size);
+    unsigned compileTimeThreadLimit = 0;
+    // Only one of thread_limit or num_threads is used, cant do it for both
+    if (ThreadLimitClause && !NumThreadsClause) {
+      Expr *ThreadLimitExpr = ThreadLimitClause->getThreadLimit();
+      clang::Expr::EvalResult Result;
+      if (ThreadLimitExpr->EvaluateAsInt(Result, CGM.getContext()))
+        compileTimeThreadLimit = Result.Val.getInt().getExtValue();
+    } else if (!ThreadLimitClause && NumThreadsClause) {
+      Expr *NumThreadsExpr = NumThreadsClause->getNumThreads();
+      clang::Expr::EvalResult Result;
+      if (NumThreadsExpr->EvaluateAsInt(Result, CGM.getContext()))
+        compileTimeThreadLimit = Result.Val.getInt().getExtValue();
+    }
+
+    // Add kernel metadata if ThreadLimit Clause is compile time constant > 0
+    if (compileTimeThreadLimit > 0) {
+      // Add the WarpSize to generic, to reflect what runtime dispatch does.
+      if (IsGeneric)
+        compileTimeThreadLimit +=
+            CGM.getTarget().getGridValue(llvm::omp::GVIDX::GV_Warp_Size);
+      if (compileTimeThreadLimit > MaxWorkGroupSz)
+        compileTimeThreadLimit = MaxWorkGroupSz;
+      std::string AttrVal = llvm::utostr(compileTimeThreadLimit);
+      FlatAttr = compileTimeThreadLimit;
+      OutlinedFn->addFnAttr("amdgpu-flat-work-group-size",
+                            AttrVal + "," + AttrVal);
+      setPropertyWorkGroupSize(CGM, OutlinedFn->getName(),
+                               compileTimeThreadLimit);
+    }
+    flatAttrEmitted = true;
+  } // end of amdgcn teams or parallel directive
+
+  // emit amdgpu-flat-work-group-size if not emitted already.
+  if (!flatAttrEmitted) {
+    std::string FlatAttrVal = llvm::utostr(DefaultWorkGroupSz);
+    OutlinedFn->addFnAttr("amdgpu-flat-work-group-size",
+                          FlatAttrVal + "," + FlatAttrVal);
+  }
+  // Emit a kernel descriptor for runtime.
+  StringRef KernDescName = OutlinedFn->getName();
+  CGOpenMPRuntime::emitStructureKernelDesc(CGM, KernDescName, FlatAttr,
+                                           IsGeneric,
+                                           1, // Uses HostServices
+                                           MaxParallelLevel);
+  // Reset it to zero for any subsequent kernel
+  MaxParallelLevel = 0;
+}
+
+void CGOpenMPRuntimeAMDGCN::emitSPMDKernelWrapper(
+    const OMPExecutableDirective &D, StringRef ParentName,
+    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
+    bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+  emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
+                 CodeGen);
+  generateMetaData(CGM, D, OutlinedFn, /*SPMD*/ false);
+}
+
+void CGOpenMPRuntimeAMDGCN::emitNonSPMDKernelWrapper(
+    const OMPExecutableDirective &D, StringRef ParentName,
+    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
+    bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+  emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
+                    CodeGen);
+  generateMetaData(CGM, D, OutlinedFn, /*Generic*/ true);
+}

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

public:
   enum ExecutionMode {
     /// SPMD execution mode (all threads are worker threads).
     EM_SPMD,
     /// Non-SPMD execution mode (1 master thread, others are workers).
     EM_NonSPMD,
     /// Unknown execution mode (orphaned directive).
     EM_Unknown,
   };
+  /// Linkage type of StaticRD Global variable
+  llvm::GlobalValue::LinkageTypes StaticRDLinkage;
+
 private:
   /// Parallel outlined function work for workers to execute.
   llvm::SmallVector<llvm::Function *, 16> Work;
 
   struct EntryFunctionState {
     llvm::BasicBlock *ExitBB = nullptr;
   };
 
private:
   //
 
   /// Creates offloading entry for the provided entry ID \a ID,
   /// address \a Addr, size \a Size, and flags \a Flags.
   void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr,
                           uint64_t Size, int32_t Flags,
                           llvm::GlobalValue::LinkageTypes Linkage) override;
 
-  /// Emit outlined function specialized for the Fork-Join
-  /// programming model for applicable target directives on the NVPTX device.
-  /// \param D Directive to emit.
-  /// \param ParentName Name of the function that encloses the target region.
-  /// \param OutlinedFn Outlined function value to be defined by this call.
-  /// \param OutlinedFnID Outlined function ID value to be defined by this call.
-  /// \param IsOffloadEntry True if the outlined function is an offload entry.
-  /// An outlined function may not be an entry if, e.g. the if clause always
-  /// evaluates to false.
-  void emitNonSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName,
-                         llvm::Function *&OutlinedFn,
-                         llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
-                         const RegionCodeGenTy &CodeGen);
-
-  /// Emit outlined function specialized for the Single Program
-  /// Multiple Data programming model for applicable target directives on the
-  /// NVPTX device.
-  /// \param D Directive to emit.
-  /// \param ParentName Name of the function that encloses the target region.
-  /// \param OutlinedFn Outlined function value to be defined by this call.
-  /// \param OutlinedFnID Outlined function ID value to be defined by this call.
-  /// \param IsOffloadEntry True if the outlined function is an offload entry.
-  /// \param CodeGen Object containing the target statements.
-  /// An outlined function may not be an entry if, e.g. the if clause always
-  /// evaluates to false.
-  void emitSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName,
-                      llvm::Function *&OutlinedFn,
-                      llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
-                      const RegionCodeGenTy &CodeGen);
-
   /// Emit outlined function for 'target' directive on the NVPTX
   /// device.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \param OutlinedFn Outlined function value to be defined by this call.
   /// \param OutlinedFnID Outlined function ID value to be defined by this call.
   /// \param IsOffloadEntry True if the outlined function is an offload entry.
   /// An outlined function may not be an entry if, e.g. the if clause always
protected:
   bool isDefaultLocationConstant() const override { return true; }
 
   /// Returns additional flags that can be stored in reserved_2 field of the
   /// default location.
   /// For NVPTX target contains data about SPMD/Non-SPMD execution mode +
   /// Full/Lightweight runtime mode. Used for better optimization.
   unsigned getDefaultLocationReserved2Flags() const override;
 
 public:
   explicit CGOpenMPRuntimeGPU(CodeGenModule &CGM);
   void clear() override;
 
   /// Declare generalized virtual functions which need to be defined
   /// by all specializations of OpenMPGPURuntime Targets like AMDGCN
   /// and NVPTX.
 
   /// Get the GPU warp size.
   virtual llvm::Value *getGPUWarpSize(CodeGenFunction &CGF) = 0;
 
   /// Get the id of the current thread on the GPU.
   virtual llvm::Value *getGPUThreadID(CodeGenFunction &CGF) = 0;
 
   /// Get the maximum number of threads in a block of the GPU.
   virtual llvm::Value *getGPUNumThreads(CodeGenFunction &CGF) = 0;
 
+  /// Allocate global variable for TransferMedium
+  virtual llvm::GlobalVariable *
+  allocateTransferMediumGlobal(CodeGenModule &CGM, llvm::ArrayType *Ty,
+                               StringRef TransferMediumName) = 0;
+
+  /// Allocate global variable for SharedStaticRD
+  virtual llvm::GlobalVariable *
+  allocateSharedStaticRDGlobal(CodeGenModule &CGM,
+                               llvm::Type *LLVMStaticTy) = 0;
+
+  /// Allocate global variable for KernelStaticGlobalized
+  virtual llvm::GlobalVariable *
+  allocateKernelStaticGlobalized(CodeGenModule &CGM) = 0;
+
+  virtual void emitSPMDKernelWrapper(const OMPExecutableDirective &D,
+                                     StringRef ParentName,
+                                     llvm::Function *&OutlinedFn,
+                                     llvm::Constant *&OutlinedFnID,
+                                     bool IsOffloadEntry,
+                                     const RegionCodeGenTy &CodeGen) = 0;
+
+  virtual void emitNonSPMDKernelWrapper(const OMPExecutableDirective &D,
+                                        StringRef ParentName,
+                                        llvm::Function *&OutlinedFn,
+                                        llvm::Constant *&OutlinedFnID,
+                                        bool IsOffloadEntry,
+                                        const RegionCodeGenTy &CodeGen) = 0;
+
+  /// Emit outlined function specialized for the Single Program
+  /// Multiple Data programming model for applicable target directives on the
+  /// NVPTX device.
+  /// \param D Directive to emit.
+  /// \param ParentName Name of the function that encloses the target region.
+  /// \param OutlinedFn Outlined function value to be defined by this call.
+  /// \param OutlinedFnID Outlined function ID value to be defined by this call.
+  /// \param IsOffloadEntry True if the outlined function is an offload entry.
+  /// \param CodeGen Object containing the target statements.
+  /// An outlined function may not be an entry if, e.g. the if clause always
+  /// evaluates to false.
+  void emitSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName,
+                      llvm::Function *&OutlinedFn,
+                      llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
+                      const RegionCodeGenTy &CodeGen);
+
+  /// Emit outlined function specialized for the Fork-Join
+  /// programming model for applicable target directives on the NVPTX device.
+  /// \param D Directive to emit.
+  /// \param ParentName Name of the function that encloses the target region.
+  /// \param OutlinedFn Outlined function value to be defined by this call.
+  /// \param OutlinedFnID Outlined function ID value to be defined by this call.
+  /// \param IsOffloadEntry True if the outlined function is an offload entry.
+  /// An outlined function may not be an entry if, e.g. the if clause always
+  /// evaluates to false.
+  void emitNonSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName,
+                         llvm::Function *&OutlinedFn,
+                         llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
+                         const RegionCodeGenTy &CodeGen);
+
   /// Emit call to void __kmpc_push_proc_bind(ident_t *loc, kmp_int32
   /// global_tid, int proc_bind) to generate code for 'proc_bind' clause.
   virtual void emitProcBindClause(CodeGenFunction &CGF,
                                   llvm::omp::ProcBindKind ProcBind,
                                   SourceLocation Loc) override;
 
   /// Emits call to void __kmpc_push_num_threads(ident_t *loc, kmp_int32
   /// global_tid, kmp_int32 num_threads) to generate code for 'num_threads'
public:
   /// their declaration context.
   enum DataSharingMode {
     /// CUDA data sharing mode.
     CUDA,
     /// Generic data-sharing mode.
     Generic,
   };
 
   /// Cleans up references to the objects in finished function.
   ///
   void functionFinished(CodeGenFunction &CGF) override;
 
   /// Choose a default value for the dist_schedule clause.
   void getDefaultDistScheduleAndChunk(CodeGenFunction &CGF,
       const OMPLoopDirective &S, OpenMPDistScheduleClauseKind &ScheduleKind,
       llvm::Value *&Chunk) const override;
 
   /// Choose a default value for the schedule clause.
   void getDefaultScheduleAndChunk(CodeGenFunction &CGF,

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

 //===---- CGOpenMPRuntimeGPU.cpp - Interface to OpenMP GPU Runtimes ----===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This provides a generalized class for OpenMP runtime code generation
 // specialized by GPU targets NVPTX and AMDGCN.
 //
 //===----------------------------------------------------------------------===//
 
 #include "CGOpenMPRuntimeGPU.h"
+#include "CGOpenMPRuntimeAMDGCN.h"
 #include "CGOpenMPRuntimeNVPTX.h"
 #include "CodeGenFunction.h"
 #include "clang/AST/Attr.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/StmtVisitor.h"
 #include "clang/Basic/Cuda.h"
 #include "llvm/ADT/SmallPtrSet.h"
void Exit(CodeGenFunction &CGF) override {
       RT.emitNonSPMDEntryFooter(CGF, EST);
     }
   } Action(EST, WST);
   CodeGen.setAction(Action);
   IsInTTDRegion = true;
   // Reserve place for the globalized memory.
   GlobalizedRecords.emplace_back();
   if (!KernelStaticGlobalized) {
-    KernelStaticGlobalized = new llvm::GlobalVariable(
-        CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
-        llvm::GlobalValue::InternalLinkage,
-        llvm::ConstantPointerNull::get(CGM.VoidPtrTy),
-        "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
-        llvm::GlobalValue::NotThreadLocal,
-        CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
+    auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGM.getOpenMPRuntime());
+    KernelStaticGlobalized = RT.allocateKernelStaticGlobalized(CGM);
   }
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
                                    IsOffloadEntry, CodeGen);
   IsInTTDRegion = false;
 
   // Now change the name of the worker function to correspond to this target
   // region's entry function.
   WST.WorkerFn->setName(Twine(OutlinedFn->getName(), "_worker"));
void Exit(CodeGenFunction &CGF) override {
       RT.emitSPMDEntryFooter(CGF, EST);
     }
   } Action(*this, EST, D);
   CodeGen.setAction(Action);
   IsInTTDRegion = true;
   // Reserve place for the globalized memory.
   GlobalizedRecords.emplace_back();
   if (!KernelStaticGlobalized) {
-    KernelStaticGlobalized = new llvm::GlobalVariable(
-        CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
-        llvm::GlobalValue::InternalLinkage,
-        llvm::ConstantPointerNull::get(CGM.VoidPtrTy),
-        "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
-        llvm::GlobalValue::NotThreadLocal,
-        CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
+    auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGM.getOpenMPRuntime());
+    KernelStaticGlobalized = RT.allocateKernelStaticGlobalized(CGM);
   }
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
                                    IsOffloadEntry, CodeGen);
   IsInTTDRegion = false;
 }
 
 void CGOpenMPRuntimeGPU::emitSPMDEntryHeader(
     CodeGenFunction &CGF, EntryFunctionState &EST,

 // Create a unique global variable to indicate the execution mode of this target
 // region. The execution mode is either 'generic', or 'spmd' depending on the
 // target directive. This variable is picked up by the offload library to setup
 // the device appropriately before kernel launch. If the execution mode is
 // 'generic', the runtime reserves one warp for the master, otherwise, all
 // warps participate in parallel work.
 static void setPropertyExecutionMode(CodeGenModule &CGM, StringRef Name,
                                      bool Mode) {
-  auto *GVMode =
-      new llvm::GlobalVariable(CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
-                               llvm::GlobalValue::WeakAnyLinkage,
-                               llvm::ConstantInt::get(CGM.Int8Ty, Mode ? 0 : 1),
-                               Twine(Name, "_exec_mode"));
+  auto *GVMode = new llvm::GlobalVariable(
+      CGM.getModule(), CGM.Int8Ty,
+      /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage,
   CGM.addCompilerUsedGlobal(GVMode);
 }
 
 void CGOpenMPRuntimeGPU::emitWorkerFunction(WorkerFunctionState &WST) {
   ASTContext &Ctx = CGM.getContext();
 
   CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
   CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, WST.WorkerFn, WST.CGFI, {},
void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
     bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
   if (!IsOffloadEntry) // Nothing to do.
     return;
 
   assert(!ParentName.empty() && "Invalid target region parent name!");
 
   bool Mode = supportsSPMDExecutionMode(CGM.getContext(), D);
   if (Mode)
-    emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
-                   CodeGen);
+    emitSPMDKernelWrapper(D, ParentName, OutlinedFn, OutlinedFnID,
+                          IsOffloadEntry, CodeGen);
   else
-    emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
-                      CodeGen);
-
+    emitNonSPMDKernelWrapper(D, ParentName, OutlinedFn, OutlinedFnID,
+                             IsOffloadEntry, CodeGen);
   setPropertyExecutionMode(CGM, OutlinedFn->getName(), Mode);
 }
 
 namespace {
 LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
 /// Enum for accesseing the reserved_2 field of the ident_t struct.
 enum ModeFlagsTy : unsigned {
   /// Bit set to 1 when in SPMD mode.
void Enter(CodeGenFunction &CGF) override {
       PrevIsInParallelRegion = IsInParallelRegion;
       IsInParallelRegion = true;
     }
     void Exit(CodeGenFunction &CGF) override {
       IsInParallelRegion = PrevIsInParallelRegion;
     }
   } Action(IsInParallelRegion);
   CodeGen.setAction(Action);
   bool PrevIsInTTDRegion = IsInTTDRegion;
   IsInTTDRegion = false;
   bool PrevIsInTargetMasterThreadRegion = IsInTargetMasterThreadRegion;
   IsInTargetMasterThreadRegion = false;
   auto *OutlinedFun =
       cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(
           D, ThreadIDVar, InnermostKind, CodeGen));
   if (CGM.getLangOpts().Optimize) {
     OutlinedFun->removeFnAttr(llvm::Attribute::NoInline);
auto *Fn = llvm::Function::Create(CGM.getTypes().GetFunctionType(CGFI),
                                     llvm::GlobalValue::InternalLinkage,
                                     "_omp_reduction_inter_warp_copy_func", &M);
   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
   Fn->setDoesNotRecurse();
   CodeGenFunction CGF(CGM);
   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
 
   CGBuilderTy &Bld = CGF.Builder;
+  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
 
   // This array is used as a medium to transfer, one reduce element at a time,
   // the data from the first lane of every warp to lanes in the first warp
   // in order to perform the final step of a reduction in a parallel region
   // (reduction across warps).  The array is placed in NVPTX __shared__ memory
   // for reduced latency, as well as to have a distinct copy for concurrently
   // executing target regions.  The array is declared with common linkage so
   // as to be shared across compilation units.
   StringRef TransferMediumName =
       "__openmp_nvptx_data_transfer_temporary_storage";
   llvm::GlobalVariable *TransferMedium =
       M.getGlobalVariable(TransferMediumName);
   unsigned WarpSize = CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size);
   if (!TransferMedium) {
     auto *Ty = llvm::ArrayType::get(CGM.Int32Ty, WarpSize);
-    unsigned SharedAddressSpace = C.getTargetAddressSpace(LangAS::cuda_shared);
-    TransferMedium = new llvm::GlobalVariable(
-        M, Ty, /*isConstant=*/false, llvm::GlobalVariable::CommonLinkage,
-        llvm::Constant::getNullValue(Ty), TransferMediumName,
-        /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
-        SharedAddressSpace);
+    TransferMedium =
+        RT.allocateTransferMediumGlobal(CGM, Ty, TransferMediumName);
     CGM.addCompilerUsedGlobal(TransferMedium);
   }
 
-  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
   // Get the CUDA thread id of the current OpenMP thread on the GPU.
   llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
   // nvptx_lane_id = nvptx_id % warpsize
   llvm::Value *LaneID = getNVPTXLaneID(CGF);
   // nvptx_warp_id = nvptx_id / warpsize
   llvm::Value *WarpID = getNVPTXWarpID(CGF);
 
   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
case CudaArch::UNKNOWN:
     break;
   case CudaArch::LAST:
     llvm_unreachable("Unexpected Cuda arch.");
   }
   llvm_unreachable("Unexpected NVPTX target without ptx feature.");
 }
 
 void CGOpenMPRuntimeGPU::clear() {
+  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGM.getOpenMPRuntime());
   if (!GlobalizedRecords.empty() &&
       !CGM.getLangOpts().OpenMPCUDATargetParallel) {
     ASTContext &C = CGM.getContext();
     llvm::SmallVector<const GlobalPtrSizeRecsTy *, 4> GlobalRecs;
     llvm::SmallVector<const GlobalPtrSizeRecsTy *, 4> SharedRecs;
     RecordDecl *StaticRD = C.buildImplicitRecord(
         "_openmp_static_memory_type_$_", RecordDecl::TagKind::TTK_Union);
     StaticRD->startDefinition();
for (const GlobalPtrSizeRecsTy &Records : GlobalizedRecords) {
         StaticRD->addDecl(Field);
         GlobalRecs.push_back(&Records);
       }
       Records.RecSize->setInitializer(llvm::ConstantInt::get(CGM.SizeTy, Size));
       Records.UseSharedMemory->setInitializer(
           llvm::ConstantInt::get(CGM.Int16Ty, UseSharedMemory ? 1 : 0));
     }
     // Allocate SharedMemorySize buffer for the shared memory.
-    // FIXME: nvlink does not handle weak linkage correctly (object with the
-    // different size are reported as erroneous).
-    // Restore this code as sson as nvlink is fixed.
     if (!SharedStaticRD->field_empty()) {
       llvm::APInt ArySize(/*numBits=*/64, SharedMemorySize);
       QualType SubTy = C.getConstantArrayType(
           C.CharTy, ArySize, nullptr, ArrayType::Normal, /*IndexTypeQuals=*/0);
       auto *Field = FieldDecl::Create(
           C, SharedStaticRD, SourceLocation(), SourceLocation(), nullptr, SubTy,
           C.getTrivialTypeSourceInfo(SubTy, SourceLocation()),
           /*BW=*/nullptr, /*Mutable=*/false,
           /*InitStyle=*/ICIS_NoInit);
       Field->setAccess(AS_public);
       SharedStaticRD->addDecl(Field);
     }
     SharedStaticRD->completeDefinition();
     if (!SharedStaticRD->field_empty()) {
       QualType StaticTy = C.getRecordType(SharedStaticRD);
       llvm::Type *LLVMStaticTy = CGM.getTypes().ConvertTypeForMem(StaticTy);
-      auto *GV = new llvm::GlobalVariable(
-          CGM.getModule(), LLVMStaticTy,
-          /*isConstant=*/false, llvm::GlobalValue::CommonLinkage,
-          llvm::Constant::getNullValue(LLVMStaticTy),
-          "_openmp_shared_static_glob_rd_$_", /*InsertBefore=*/nullptr,
-          llvm::GlobalValue::NotThreadLocal,
-          C.getTargetAddressSpace(LangAS::cuda_shared));
+      auto *GV = RT.allocateSharedStaticRDGlobal(CGM, LLVMStaticTy);
       auto *Replacement = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
           GV, CGM.VoidPtrTy);
       for (const GlobalPtrSizeRecsTy *Rec : SharedRecs) {
         Rec->Buffer->replaceAllUsesWith(Replacement);
         Rec->Buffer->eraseFromParent();
       }
     }
     StaticRD->completeDefinition();
     if (!StaticRD->field_empty()) {
       QualType StaticTy = C.getRecordType(StaticRD);
       std::pair<unsigned, unsigned> SMsBlockPerSM = getSMsBlocksPerSM(CGM);
       llvm::APInt Size1(32, SMsBlockPerSM.second);
       QualType Arr1Ty =
           C.getConstantArrayType(StaticTy, Size1, nullptr, ArrayType::Normal,
                                  /*IndexTypeQuals=*/0);
       llvm::APInt Size2(32, SMsBlockPerSM.first);
       QualType Arr2Ty =
           C.getConstantArrayType(Arr1Ty, Size2, nullptr, ArrayType::Normal,
                                  /*IndexTypeQuals=*/0);
       llvm::Type *LLVMArr2Ty = CGM.getTypes().ConvertTypeForMem(Arr2Ty);
-      // FIXME: nvlink does not handle weak linkage correctly (object with the
-      // different size are reported as erroneous).
-      // Restore CommonLinkage as soon as nvlink is fixed.
-      auto *GV = new llvm::GlobalVariable(
-          CGM.getModule(), LLVMArr2Ty,
-          /*isConstant=*/false, llvm::GlobalValue::InternalLinkage,
+      auto *GV =
+          new llvm::GlobalVariable(CGM.getModule(), LLVMArr2Ty,
+                                   /*isConstant=*/false, RT.StaticRDLinkage,
                                    llvm::Constant::getNullValue(LLVMArr2Ty),
                                    "_openmp_static_glob_rd_$_");
       auto *Replacement = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
           GV, CGM.VoidPtrTy);
       for (const GlobalPtrSizeRecsTy *Rec : GlobalRecs) {
         Rec->Buffer->replaceAllUsesWith(Replacement);
         Rec->Buffer->eraseFromParent();
       }
     }
   }
if (!TeamsReductions.empty()) {
     QualType StaticTy = C.getRecordType(StaticRD);
     llvm::Type *LLVMReductionsBufferTy =
         CGM.getTypes().ConvertTypeForMem(StaticTy);
     // FIXME: nvlink does not handle weak linkage correctly (object with the
     // different size are reported as erroneous).
     // Restore CommonLinkage as soon as nvlink is fixed.
     auto *GV = new llvm::GlobalVariable(
         CGM.getModule(), LLVMReductionsBufferTy,
-        /*isConstant=*/false, llvm::GlobalValue::InternalLinkage,
+        /*isConstant=*/false, RT.StaticRDLinkage,
         llvm::Constant::getNullValue(LLVMReductionsBufferTy),
         "_openmp_teams_reductions_buffer_$_");
     KernelTeamsReductionPtr->setInitializer(
         llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(GV,
                                                              CGM.VoidPtrTy));
   }
   CGOpenMPRuntime::clear();
 }

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h

public:
   /// Get the GPU warp size.
   llvm::Value *getGPUWarpSize(CodeGenFunction &CGF) override;
 
   /// Get the id of the current thread on the GPU.
   llvm::Value *getGPUThreadID(CodeGenFunction &CGF) override;
 
   /// Get the maximum number of threads in a block of the GPU.
   llvm::Value *getGPUNumThreads(CodeGenFunction &CGF) override;
+
+  /// Allocate global variable for TransferMedium
+  virtual llvm::GlobalVariable *
+  allocateTransferMediumGlobal(CodeGenModule &CGM, llvm::ArrayType *Ty,
+                               StringRef Name) override;
+
+  /// Allocate global variable for SharedStaticRD
+  virtual llvm::GlobalVariable *
+  allocateSharedStaticRDGlobal(CodeGenModule &CGM,
+                               llvm::Type *LLVMStaticTy) override;
+
+  /// Allocate global variable for KernelStaticGlobalized
+  virtual llvm::GlobalVariable *
+  allocateKernelStaticGlobalized(CodeGenModule &CGM) override;
+
+  /// Emit target specific SPMD kernel
+  virtual void emitSPMDKernelWrapper(const OMPExecutableDirective &D,
+                                     StringRef ParentName,
+                                     llvm::Function *&OutlinedFn,
+                                     llvm::Constant *&OutlinedFnID,
+                                     bool IsOffloadEntry,
+                                     const RegionCodeGenTy &CodeGen) override;
+  /// Emit target specific Non-SPMD kernel
+  virtual void
+  emitNonSPMDKernelWrapper(const OMPExecutableDirective &D,
+                           StringRef ParentName, llvm::Function *&OutlinedFn,
+                           llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
+                           const RegionCodeGenTy &CodeGen) override;
 };
 
 } // CodeGen namespace.
 } // clang namespace.
 
 #endif // LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMENVPTX_H

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

 using namespace clang;
 using namespace CodeGen;
 using namespace llvm::omp;
 
 CGOpenMPRuntimeNVPTX::CGOpenMPRuntimeNVPTX(CodeGenModule &CGM)
     : CGOpenMPRuntimeGPU(CGM) {
   if (!CGM.getLangOpts().OpenMPIsDevice)
     llvm_unreachable("OpenMP NVPTX can only handle device code.");
+
+  // FIXME: nvlink does not handle weak linkage correctly (object with the
+  // different size are reported as erroneous).
+  // Restore CommonLinkage as soon as nvlink is fixed.
+  StaticRDLinkage = llvm::GlobalValue::InternalLinkage;
 }
 
 llvm::Value *CGOpenMPRuntimeNVPTX::getGPUWarpSize(CodeGenFunction &CGF) {
   return CGF.EmitRuntimeCall(
       llvm::Intrinsic::getDeclaration(
           &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_warpsize),
       "nvptx_warp_size");
 }
 
 llvm::Value *CGOpenMPRuntimeNVPTX::getGPUThreadID(CodeGenFunction &CGF) {
   CGBuilderTy &Bld = CGF.Builder;
   llvm::Function *F;
   F = llvm::Intrinsic::getDeclaration(
       &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_tid_x);
   return Bld.CreateCall(F, llvm::None, "nvptx_tid");
 }
 
 llvm::Value *CGOpenMPRuntimeNVPTX::getGPUNumThreads(CodeGenFunction &CGF) {
   CGBuilderTy &Bld = CGF.Builder;
   llvm::Function *F;
   F = llvm::Intrinsic::getDeclaration(
       &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_ntid_x);
   return Bld.CreateCall(F, llvm::None, "nvptx_num_threads");
 }
+
+llvm::GlobalVariable *CGOpenMPRuntimeNVPTX::allocateTransferMediumGlobal(
+    CodeGenModule &CGM, llvm::ArrayType *Ty, StringRef TransferMediumName) {
+  return new llvm::GlobalVariable(
+      CGM.getModule(), Ty, /*isConstant=*/false,
+      llvm::GlobalVariable::CommonLinkage, llvm::Constant::getNullValue(Ty),
+      TransferMediumName,
+      /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
+}
+
+llvm::GlobalVariable *
+CGOpenMPRuntimeNVPTX::allocateSharedStaticRDGlobal(CodeGenModule &CGM,
+                                                   llvm::Type *LLVMStaticTy) {
+  return new llvm::GlobalVariable(
+      CGM.getModule(), LLVMStaticTy,
+      /*isConstant=*/false, llvm::GlobalValue::CommonLinkage,
+      llvm::Constant::getNullValue(LLVMStaticTy),
+      "_openmp_shared_static_glob_rd_$_", /*InsertBefore=*/nullptr,
+      llvm::GlobalValue::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
+}
+
+llvm::GlobalVariable *
+CGOpenMPRuntimeNVPTX::allocateKernelStaticGlobalized(CodeGenModule &CGM) {
+  return new llvm::GlobalVariable(
+      CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
+      llvm::GlobalValue::InternalLinkage,
+      llvm::ConstantPointerNull::get(CGM.VoidPtrTy),
+      "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
+      llvm::GlobalValue::NotThreadLocal,
+      CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
+}
+
+void CGOpenMPRuntimeNVPTX::emitSPMDKernelWrapper(
+    const OMPExecutableDirective &D, StringRef ParentName,
+    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
+    bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+  emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
+                 CodeGen);
+}
+
+void CGOpenMPRuntimeNVPTX::emitNonSPMDKernelWrapper(
+    const OMPExecutableDirective &D, StringRef ParentName,
+    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
+    bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+  emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
+                    CodeGen);
+}

clang/test/OpenMP/amdgcn_target_codegen.cpp

 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 #define N 1000
 
+// CHECK: @"_openmp_kernel_static_glob_rd$ptr" = weak addrspace(3) externally_initialized global i8* undef
+
+// CHECK: @__omp_offloading_{{.*}}test_amdgcn_target_tid_threadsv_l36_kern_desc = weak constant %struct.__tgt_attribute_struct { i16 2, i16 9, i16 0, i8 1, i8 1, i8 0 }, align 2
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_tid_threadsv_l36_exec_mode = weak constant i8 1
+
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_tid_threads_simdv_l52_kern_desc = weak constant %struct.__tgt_attribute_struct { i16 2, i16 9, i16 0, i8 0, i8 1, i8 0 }, align 2
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_tid_threads_simdv_l52_exec_mode = weak constant i8 0
+
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_attributes_spmdv_l63_wg_size = weak addrspace(1) constant i16 10
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_attributes_spmdv_l63_kern_desc = weak constant %struct.__tgt_attribute_struct { i16 2, i16 9, i16 10, i8 0, i8 1, i8 0 }, align 2
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_attributes_spmdv_l63_exec_mode = weak constant i8 0
+
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_attributes_non_spmdv_l75_wg_size = weak addrspace(1) constant i16 74
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_attributes_non_spmdv_l75_kern_desc = weak constant %struct.__tgt_attribute_struct { i16 2, i16 9, i16 74, i8 1, i8 1, i8 0 }, align 2
+// CHECK : @__omp_offloading_{{.*}}test_amdgcn_target_attributes_non_spmdv_l75_exec_mode = weak constant i8 1
+
 int test_amdgcn_target_tid_threads() {
 // CHECK-LABEL: define weak void @{{.*}}test_amdgcn_target_tid_threads
 
   int arr[N];
 
 // CHECK: [[NUM_THREADS:%.+]] = call i64 @__ockl_get_local_size(i32 0)
 // CHECK-NEXT: [[VAR:%.+]] = trunc i64 [[NUM_THREADS]] to i32
 // CHECK-NEXT: sub nuw i32 [[VAR]], 64

 // CHECK-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[VAR]], i16 0, i16 0)
 #pragma omp target simd
   for (int i = 0; i < N; i++) {
     arr[i] = 1;
   }
   return arr[0];
 }
 
+int test_amdgcn_target_attributes_spmd() {
+  int arr[N];
+
+// CHECK: {{.*}}"amdgpu-flat-work-group-size"="10,10"
+#pragma omp target parallel num_threads(10)
+  for (int i = 0; i < N; i++) {
+    arr[i] = 1;
+  }
+
+  return arr[0];
+}
+
+int test_amdgcn_target_attributes_non_spmd() {
+  int arr[N];
+
+// CHECK: {{.*}}"amdgpu-flat-work-group-size"="74,74"
+#pragma omp target teams thread_limit(10)
+  for (int i = 0; i < N; i++) {
+    arr[i] = 1;
+  }
+
+  return arr[0];
+}
+
+int test_amdgcn_target_attributes_max_work_group_size() {
+  int arr[N];
+
+// CHECK: {{.*}}"amdgpu-flat-work-group-size"="1024,1024"
+#pragma omp target teams thread_limit(1500)
+  for (int i = 0; i < N; i++) {
+    arr[i] = 1;
+  }
+
+  return arr[0];
+}
+
 #endif