[WIP] Generate LLVM IR for target parallel kernels

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -96,7 +96,7 @@ CodeGenFunction::~CodeGenFunction() {
   // time of the CodeGenModule, because we have to ensure the IR has not yet
   // been "emitted" to the outside, thus, modifications are still sensible.
   if (CGM.getLangOpts().OpenMPIRBuilder && CurFn)
-    CGM.getOpenMPRuntime().getOMPBuilder().finalize(CurFn);
+    CGM.getOpenMPRuntime().getOMPBuilder().finalizeFunction(CurFn);
 }
 
 // Map the LangOption for exception behavior into

clang/test/OpenMP/irbuilder_omp_offload_metadata.c

@@ -0,0 +1,16 @@
+// This test checks if OpenMPIRBuilder generates the same number of omp offload
+// info nodes as Clang does. The wrong number of metadata nodes can provide
+// miscompilation of the device code for enabled OpenMPIRBuilder
+// RUN: %clang_cc1 -triple x86_64--unknown-linux-gnu -emit-llvm -fopenmp -fopenmp-enable-irbuilder -fopenmp-targets=amdgcn-amd-amdhsa -faddrsig  %s -o - | FileCheck --check-prefix BUILDER %s
+// RUN: %clang_cc1 -triple x86_64--unknown-linux-gnu -emit-llvm -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -faddrsig  %s -o - | FileCheck --check-prefix NOBUILDER %s
+
+void omp_offload_metadata_irbuilder_test() {
+int a[256];
+#pragma omp target parallel for
+  for (int i = 0; i < 256; i++) {
+    a[i] = i;
+  }
+}
+
+//BUILDER: !omp_offload.info = !{!{{[0-9]+}}}
+//NOBUILDER: !omp_offload.info = !{!{{[0-9]+}}}

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

@@ -448,10 +448,14 @@ class OpenMPIRBuilder {
 
   void setConfig(OpenMPIRBuilderConfig C) { Config = C; }
 
-  /// Finalize the underlying module, e.g., by outlining regions.
+  /// Finalize the underlying function, e.g., by outlining regions.
   /// \param Fn                    The function to be finalized. If not used,
   ///                              all functions are finalized.
-  void finalize(Function *Fn = nullptr);
+  void finalizeFunction(Function *Fn = nullptr);
+
+  /// Finalize the underlying module. Finalize all functions and create
+  /// offload metadata for the module
+  void finalizeModule();
 
   /// Add attributes known for \p FnID to \p Fn.
   void addAttributes(omp::RuntimeFunction FnID, Function &Fn);
@@ -870,6 +874,22 @@ class OpenMPIRBuilder {
       Type *LlvmPtrTy, Constant *Addr);
 
 private:
+  /// Modifies the canonical loop to be a statically-scheduled workshare loop
+  /// which is executed on the device
+  ///
+  /// This takes a \p LoopInfo representing a canonical loop, such as the one
+  /// created by \p createCanonicalLoop and emits additional instructions to
+  /// turn it into a workshare loop. In particular, it calls to an OpenMP
+  /// runtime function in the preheader to call OpenMP device rtl function
+  /// which handles worksharing of loop body interations.
+  ///
+  /// \param DL       Debug location for instructions added for the
+  ///                 workshare-loop construct itself.
+  /// \param CLI      A descriptor of the canonical loop to workshare.
+  ///
+  /// \returns Point where to insert code after the workshare construct.
+  InsertPointTy applyWorkshareLoopDevice(DebugLoc DL, CanonicalLoopInfo *CLI);
+
   /// Modifies the canonical loop to be a statically-scheduled workshare loop.
   ///
   /// This takes a \p LoopInfo representing a canonical loop, such as the one
@@ -2144,7 +2164,7 @@ class OpenMPIRBuilder {
                              int32_t NumThreads,
                              SmallVectorImpl<Value *> &Inputs,
                              GenMapInfoCallbackTy GenMapInfoCB,
-                             TargetBodyGenCallbackTy BodyGenCB);
+                             TargetBodyGenCallbackTy BodyGenCB, bool IsSPMD);
 
   /// Returns __kmpc_for_static_init_* runtime function for the specified
   /// size \a IVSize and sign \a IVSigned. Will create a distribute call

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

@@ -461,6 +461,10 @@ __OMP_RTL(__kmpc_target_deinit, false, Void,)
 __OMP_RTL(__kmpc_kernel_prepare_parallel, false, Void, VoidPtr)
 __OMP_RTL(__kmpc_parallel_51, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
           VoidPtr, VoidPtr, VoidPtrPtr, SizeTy)
+__OMP_RTL(__kmpc_for_static_loop_4, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32)
+__OMP_RTL(__kmpc_for_static_loop_4u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32)
+__OMP_RTL(__kmpc_for_static_loop_8, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64)
+__OMP_RTL(__kmpc_for_static_loop_8u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64)
 __OMP_RTL(__kmpc_kernel_parallel, false, Int1, VoidPtrPtr)
 __OMP_RTL(__kmpc_kernel_end_parallel, false, Void, )
 __OMP_RTL(__kmpc_serialized_parallel, false, Void, IdentPtr, Int32)