Generate similiar *.ll code as Clang does for #pragma omp target parallel

tools/flang1/flang1exe/symtab.c

@@ -479,7 +479,8 @@ get_ieee_arith_intrin(char *nm)
 int
 getsymbol(const char *name)
 {
-  return getsym(name, strlen(name));
+  int sym = getsym(name, strlen(name));
+  return sym;
 }
 
 /** \brief Enter symbol with indicated name into symbol table, initialize

tools/flang2/flang2exe/cgmain.cpp

@@ -4301,7 +4301,7 @@ make_stmt(STMT_Type stmt_type, int ilix, bool deletable, SPTR next_bih_label,
   int alignment;
   INSTR_LIST *Curr_Instr;
 
-  DBGTRACEIN2(" type: %s ilix: %d", stmt_names[stmt_type], ilix)
+  DBGTRACEIN2(" type: %s ilix: %d", stmt_names[stmt_type], ilix);
 
   curr_stmt_type = stmt_type;
   if (last_stmt_is_branch && stmt_type != STMT_LABEL) {
@@ -12267,7 +12267,7 @@ process_sptr_offset(SPTR sptr, ISZ_T off)
     }
     if ((flg.smp || (XBIT(34, 0x200) || gbl.usekmpc)) &&
         (gbl.outlined || ISTASKDUPG(GBL_CURRFUNC))) {
-      if (sptr == ll_get_shared_arg(gbl.currsub)) {
+      if (sptr == ll_get_shared_arg(gbl.currsub) && !gbl.is_init_spmd_kernel) {
         LLTYPE(sptr) = make_ptr_lltype(make_lltype_from_dtype(DT_INT8));
       }
     }
@@ -14063,7 +14063,6 @@ process_formal_arguments(LL_ABI_Info *abi)
       /* Other by-value kinds. */
       break;
     }
-
     /* This op represents the real LLVM argument, not the local variable. */
     arg_op = make_operand();
     arg_op->ot_type = OT_VAR;
@@ -14407,7 +14406,8 @@ INLINE void static add_property_struct(char *func_name,
   print_token("@");
   print_token(func_name);
 
-  if (is_SPMD_mode(mode)) {
+  if (mode >= mode_target_teams_distribute_parallel_for
+	        && mode <= mode_target_parallel_for_simd) {
     print_token("__exec_mode = weak constant i8 2\n");
   }
   else {

tools/flang2/flang2exe/exp_ftn.cpp

@@ -4147,7 +4147,7 @@ exp_bran(ILM_OP opc, ILM *ilmp, int curilm)
 /***************************************************************/
 
 void
-exp_misc(ILM_OP opc, ILM *ilmp, int curilm)
+exp_misc(ILM_OP opc, ILM *ilmp, int curilm, bool process_expanded)
 {
   int tmp;
   int ilix, listilix;
@@ -4289,11 +4289,11 @@ exp_misc(ILM_OP opc, ILM *ilmp, int curilm)
     break;
 
   case IM_ENDF:
-    exp_end(ilmp, curilm, true);
+    exp_end(ilmp, curilm, true, process_expanded);
     break;
 
   case IM_END:
-    exp_end(ilmp, curilm, false);
+    exp_end(ilmp, curilm, false, process_expanded);
     break;
 
   case IM_BYVAL:

tools/flang2/flang2exe/exp_ftn.h

@@ -53,7 +53,7 @@ void exp_bran(ILM_OP opc, ILM *ilmp, int curilm);
 /**
    \brief ...
  */
-void exp_misc(ILM_OP opc, ILM *ilmp, int curilm);
+void exp_misc(ILM_OP opc, ILM *ilmp, int curilm, bool process_expanded = false);
 
 /**
    \brief ...

tools/flang2/flang2exe/exp_rte.cpp

@@ -2139,7 +2139,7 @@ exp_alloca(ILM *ilmp)
 static void gen_funcret(finfo_t *);
 
 void
-exp_end(ILM *ilmp, int curilm, bool is_func)
+exp_end(ILM *ilmp, int curilm, bool is_func, bool process_expanded)
 {
   int tmp;
   int op1;
@@ -2158,10 +2158,12 @@ exp_end(ILM *ilmp, int curilm, bool is_func)
   int ilix;
   if (flg.omptarget && !is_func) {
     if (XBIT(232, 0x40) && gbl.ompaccel_intarget && !OMPACCFUNCDEVG(gbl.currsub) /*is_gpu_output_file() */  ) {
-      ilix = ll_make_kmpc_target_deinit(
-        ompaccel_tinfo_get(gbl.currsub)->mode);
-      iltb.callfg = 1;
-      chk_block(ilix);
+      OMP_TARGET_MODE mode = ompaccel_tinfo_get(gbl.currsub)->mode;
+      if (!is_SPMD_mode(mode) && !process_expanded) {
+        ilix = ll_make_kmpc_target_deinit(mode);
+        iltb.callfg = 1;
+        chk_block(ilix);
+      }
     }
   }
 #endif

tools/flang2/flang2exe/exp_rte.h

@@ -100,7 +100,7 @@ void exp_cgoto(ILM *ilmp, int curilm);
 /**
    \brief ...
  */
-void exp_end(ILM *ilmp, int curilm, bool is_func);
+void exp_end(ILM *ilmp, int curilm, bool is_func, bool process_expanded = false);
 
 /**
    \brief ...

tools/flang2/flang2exe/expand.cpp

@@ -54,6 +54,8 @@
 #ifdef OMP_OFFLOAD_LLVM
 #include "tgtutil.h"
 #include "kmpcutil.h"
+#include <vector>
+#include <map>
 #endif
 extern int in_extract_inline; /* Bottom-up auto-inlining */
 
@@ -62,6 +64,8 @@ static int create_ref(SPTR sym, int *pnmex, int basenm, int baseilix,
                       int *pclen, int *pmxlen, int *prestype);
 static int jsr2qjsr(int);
 
+SPTR
+eval_ilm_check_if_skip(int ilmx, int *skip_expand = nullptr, int *process_expanded = nullptr);
 #define DO_PFO ((XBIT(148, 0x1000) && !XBIT(148, 0x4000)) || XBIT(148, 1))
 
 /***************************************************************/
@@ -215,7 +219,6 @@ parse_im_file(const ILM *ilmp, int *lineno_out, int *findex_out, int *ftag_out)
 }
 
 /***************************************************************/
-
 /** \brief Expand ILMs to ILIs */
 int
 expand(void)
@@ -229,12 +232,36 @@ expand(void)
   int last_ftag = 0;
   int nextftag = 0, nextfindex = 0;
   int last_cpp_branch = 0;
-
+  static int skip_expand;
+  static int skip_expand_sptr;
+  static std::map<int, int> process_expanded_map = std::map<int,int>();
+  auto it = process_expanded_map.find(gbl.currsub);
+  int process_expanded = 0;
+
+  // we reset flag because we do not know if we generate initialization
+  // function for SPMD kernel (the function with kmpc_parallel_51 call)
+  // or the proper kernel code (the function which is passed as an argument
+  // to kmpc_parallel_51 call or generic kernel
+  gbl.is_init_spmd_kernel = false;
+  if (it != process_expanded_map.end())
+  {
+	  process_expanded = it->second;
+  }
+  else
+  {
+	 process_expanded = 0;
+  }
   /*
    * NOTE, for an ILM: ilmx is needed to access the ILM_AUX area, ilmp is
    * needed to access the ILM area
    */
   exp_init();
+
+  //set current target info if given target region was already processed
+  if(ompaccel_tinfo_get(gbl.currsub))
+  {
+	  ompaccel_tinfo_current_set(ompaccel_tinfo_get(gbl.currsub));
+  }
   /* During expand, we want to generate unique proc ili each time a
    * proc ILM is processed.  The assumption is that the scheduler will
    * cse a proc ili if it appears multiple times in a block. E.g.,
@@ -299,7 +326,13 @@ expand(void)
 
       ilmp = (ILM *)(ilmb.ilm_base + ilmx);
       opc = ILM_OPC(ilmp);
-
+      /* Do not expand map statements for helper function for kmpc_parallel_51 */
+      if ((opc == IM_MP_MAP || opc == IM_MP_EMAP) && process_expanded)
+	      continue;
+      if (process_expanded)
+      {
+	      gbl.ompoutlinedfunc = gbl.currsub;
+      }
       if (opc == IM_BR) {
         last_cpp_branch = ILM_OPND(ilmp, 1);
       } else if (opc == IM_LABEL) {
@@ -318,8 +351,17 @@ expand(void)
                                    * variable operands */
       if (IM_TRM(opc)) {
         int cur_label = BIH_LABEL(expb.curbih);
-        eval_ilm(ilmx);
-      }
+	if (!skip_expand){
+          SPTR sptr1 = eval_ilm_check_if_skip(ilmx, &skip_expand, &process_expanded);
+	if (skip_expand) {
+          skip_expand_sptr = sptr1;
+          process_expanded_map[skip_expand_sptr] = 1;
+          ll_write_ilm_header((int)sptr1, ilmx);
+          restartRewritingILM(ilmx);
+	}
+	} else {
+	  ll_rewrite_ilms(-1, ilmx, len);
+	}}
       else if (flg.smp && len) {
         ll_rewrite_ilms(-1, ilmx, len);
       }
@@ -366,7 +408,6 @@ expand(void)
     new_callee_scope = 0;
   }
   while (opc != IM_END && opc != IM_ENDF);
-
   if (DBGBIT(10, 2) && (bihb.stg_avail != 1)) {
     int bih;
     for (bih = 1; bih != 0; bih = BIH_NEXT(bih)) {
@@ -423,6 +464,13 @@ expand(void)
   } else {
     fihb.nextfindex = fihb.currfindex = 1;
   }
+  if (skip_expand && !process_expanded)
+  {
+	process_expanded = 1;
+	unsetRewritingILM();
+  }
+  skip_expand = 0;
+
   return expb.nilms;
 }
 
@@ -451,10 +499,56 @@ eval_ilm_argument1(int opr, ILM *ilmpx, int ilmx)
   }
 } /* eval_ilm_argument1 */
 
-void
-eval_ilm(int ilmx)
+static std::vector<int> get_allocated_symbols(OMPACCEL_TINFO *orig_symbols)
+{
+  int num_of_symbols = orig_symbols->n_symbols;
+  char allocated_symbol_name[128];
+  SPTR allocated_symbol;
+  std::vector<int> init_symbols{};
+  int store_instr;
+  int load_instr;
+  for (unsigned i = 0; i < num_of_symbols; ++i) {
+    if (PASSBYVALG(orig_symbols->symbols[i].device_sym) &&
+              !PASSBYREFG(orig_symbols->symbols[i].device_sym))
+	    continue;
+    if (!DT_ISSCALAR(DTYPEG(orig_symbols->symbols[i].device_sym))
+         && STYPEG(orig_symbols->symbols[i].host_sym) != ST_STRUCT) {
+	    continue;
+    }
+    snprintf(allocated_symbol_name, sizeof(allocated_symbol_name),
+            ".allocated_symbol_%d", i);
+    allocated_symbol = getsymbol(allocated_symbol_name);
+    STYPEP(allocated_symbol, ST_VAR);
+    if (STYPEG(orig_symbols->symbols[i].host_sym) == ST_STRUCT)
+      DTYPEP(allocated_symbol,DT_CPTR);
+    else
+      DTYPEP(allocated_symbol,
+           get_type(2,TY_PTR,DTYPEG(orig_symbols->symbols[i].device_sym)));
+    SCP(allocated_symbol, SC_AUTO);
+    store_instr = ad4ili(IL_ST,
+                         ad_acon(orig_symbols->symbols[i].device_sym,0),
+                         ad_acon(allocated_symbol,0),
+                         addnme(NT_VAR, allocated_symbol, 0,0),
+                         MSZ_I8);
+    chk_block(store_instr);
+    load_instr = mk_ompaccel_ldsptr(allocated_symbol);
+    chk_block(load_instr);
+
+    init_symbols.push_back(load_instr);
+  }
+  return init_symbols;
+
+}
+void eval_ilm(int ilmx)
+{
+  eval_ilm_check_if_skip(ilmx, nullptr, nullptr);
+}
+
+SPTR
+eval_ilm_check_if_skip(int ilmx, int *skip_expand, int *process_expanded)
 {
 
+		SPTR sptr1 = SPTR_NULL;
   ILM *ilmpx;
   int noprs,   /* number of operands in the ILM	 */
       ilix,    /* ili index				 */
@@ -478,7 +572,7 @@ eval_ilm(int ilmx)
         /* Set line no for EPARx */
         gbl.lineno = ILM_OPND(ilmpx, 1);
       }
-      return;
+      return sptr1;
     }
   }
 
@@ -510,12 +604,12 @@ eval_ilm(int ilmx)
           }
         } else if (opcx == IM_MP_EREDUCTION) {
           ompaccel_notify_reduction(false);
-          return;
+          return sptr1;
         }
       }
 
       if (ompaccel_is_reduction_region())
-        return;
+        return sptr1;
     }
 #endif
     /*-
@@ -614,7 +708,7 @@ eval_ilm(int ilmx)
     if (IM_I8(opcx))
       ILM_RESTYPE(ilmx) = ILM_ISI8;
 
-    return;
+    return sptr1;
   }
   switch (IM_TYPE(opcx)) { /* special-cased ILM		 */
 
@@ -645,7 +739,10 @@ eval_ilm(int ilmx)
     break;
 
   case IMTY_MISC: /* miscellaneous  */
-    exp_misc(opcx, ilmpx, ilmx);
+    if (process_expanded && *process_expanded)
+      exp_misc(opcx, ilmpx, ilmx, true);
+    else
+      exp_misc(opcx, ilmpx, ilmx);
     break;
 
   case IMTY_FSTR: /* fortran string */
@@ -687,7 +784,12 @@ eval_ilm(int ilmx)
     /* We do not initialize spmd kernel library since we do not use spmd data
      * sharing model. It does extra work and allocates device on-chip memory.
      * */
-    if (XBIT(232, 0x40) && gbl.ompaccel_intarget) {
+    if (XBIT(232, 0x40) && gbl.ompaccel_intarget && !*process_expanded) {
+      //TODO move initialization to separate function
+      std::vector<int> allocated_symbols;
+      if (is_SPMD_mode(ompaccel_tinfo_get(gbl.currsub)->mode)) {
+	  allocated_symbols = get_allocated_symbols(ompaccel_tinfo_get(gbl.currsub));
+      }
       ilix = ll_make_kmpc_target_init(ompaccel_tinfo_get(gbl.currsub)->mode);
 
       /* Generate new control flow for generic kernel */
@@ -714,9 +816,24 @@ eval_ilm(int ilmx)
       exp_label(target_code_lab);
 
       if (is_SPMD_mode(ompaccel_tinfo_get(gbl.currsub)->mode)) {
-        iltb.callfg = 1;
         ilix = ll_make_kmpc_global_thread_num();
+        iltb.callfg = 1;
         chk_block(ilix);
+        gbl.is_init_spmd_kernel = true;
+        sptr1 = ll_make_helper_function_for_kmpc_parallel_51((SPTR)0, ompaccel_tinfo_get(gbl.currsub));
+        ilix = ll_make_kmpc_parallel_51(ilix, allocated_symbols, sptr1);
+        iltb.callfg = 1;
+        chk_block(ilix);
+        ilix = ll_make_kmpc_target_deinit(ompaccel_tinfo_get(gbl.currsub)->mode);
+        iltb.callfg = 1;
+        chk_block(ilix);
+        expb.curilt = addilt(expb.curilt, ad1ili(IL_EXIT, gbl.currsub));
+        BIH_XT(expb.curbih) = 1;
+        BIH_LAST(expb.curbih) = 1;
+        wr_block();
+	if (skip_expand && process_expanded && (*process_expanded == 0)){
+	  *skip_expand = 1;
+	} 
       }
 
       iltb.callfg = 1;
@@ -727,6 +844,7 @@ eval_ilm(int ilmx)
 #endif
   if (IM_I8(opcx))
     ILM_RESTYPE(ilmx) = ILM_ISI8;
+  return sptr1;
 }
 
 /***************************************************************/