[OPENMP] Remove unnecessary code generation of firstprivate clause for target directive in nvptx

[carlo.bertolli]

Summary

This patch removes code generation of firstprivate clause for target when compiling for nvptx. By using appropriate map types, the OpenMP offloading runtime already allocates a variable in gpu memory and copies the host value into it. This can be used directly by the target region as firstprivate value.

This modification is an optimization of the current support as it removes the need of allocating a further variable in the target region.

[ABataev]

Update patch to the latest version. It is not allowed to override CodeGen sequences in runtime. They must be universal and all must be defined in CGStmtOpenMP.cpp only

[carlo.bertolli]

Codegen of region bodies should not be specialized based on target.

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

void CGOpenMPRuntimeNVPTX::emitTargetOutlinedFunction(
   const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
 
   EntryFunctionState EST;
   WorkerFunctionState WST(CGM);
 
   // Emit target region as a standalone region.
   auto &&CodeGen = [&EST, &WST, &CS, &D, this](CodeGenFunction &CGF) {
     CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
-    (void)CGF.EmitOMPFirstprivateClause(D, PrivateScope);
     CGF.EmitOMPPrivateClause(D, PrivateScope);
     (void)PrivateScope.Privatize();
 
     emitEntryHeader(CGF, EST, WST);
     CGF.EmitStmt(CS.getCapturedStmt());
     emitEntryFooter(CGF, EST);
   };
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,

test/OpenMP/nvptx_target_firstprivate_codegen.cpp

int foo(int n, double *ptr) {
   TT<long long, char> d;
   
   #pragma omp target firstprivate(a)
   {
   }
   
   // TCHECK:  define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
   // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
-  // TCHECK:  [[A1:%.+]] = alloca i{{[0-9]+}},
   // TCHECK:  store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
   // TCHECK-64:  [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
-  // TCHECK-64:  [[A_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV]],
-  // TCHECK-32:  [[A_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
-  // TCHECK:  store i{{[0-9]+}} [[A_ADDR_VAL]], i{{[0-9]+}}* [[A1]],
   // TCHECK:  ret void  
 
 #pragma omp target firstprivate(aa,b,c,d)
   {
     aa += 1;
     b[2] = 1.0;
     c[1][2] = 1.0;
     d.X = 1;
     d.Y = 1;    
   }
   
-  // make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the
-  // target region
+  // the input parameters to the offloading function are already private variables to the target region
   // TCHECK:  define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A2_IN:%.+]], [10 x float]* {{.+}} [[B_IN:%.+]], [5 x [10 x double]]* {{.+}} [[C_IN:%.+]], [[TT]]* {{.+}} [[D_IN:%.+]])
   // TCHECK:  [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
   // TCHECK:  [[B_ADDR:%.+]] = alloca [10 x float]*,
   // TCHECK:  [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*,
   // TCHECK:  [[D_ADDR:%.+]] = alloca [[TT]]*,
-  // TCHECK:  [[A2_PRIV:%.+]] = alloca i{{[0-9]+}},
-  // TCHECK:  [[B_PRIV:%.+]] = alloca [10 x float],
-  // TCHECK:  [[C_PRIV:%.+]] = alloca [5 x [10 x double]],
-  // TCHECK:  [[D_PRIV:%.+]] = alloca [[TT]],
   // TCHECK:  store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]],
   // TCHECK:  store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]],
   // TCHECK:  store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]],
   // TCHECK:  store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]],
   // TCHECK:  [[CONV_A2ADDR:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
   // TCHECK:  [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]],
   // TCHECK:  [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]],
   // TCHECK:  [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]],
 
-  // firstprivate(aa): a_priv = a_in
-  // TCHECK:  [[A2_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV_A2ADDR]],
-  // TCHECK:  store i{{[0-9]+}} [[A2_CONV_VAL]], i{{[0-9]+}}* [[A2_PRIV]],
+  // aa += 1 (only check load and store are from and to right variable)
+  // TCHECK:  {{.+}} = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV_A2ADDR]],
+  // TCHECK:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[CONV_A2ADDR]],
   
-  //  firstprivate(b): memcpy(b_priv,b_in)
-  // TCHECK:  [[B_PRIV_BCAST:%.+]] = bitcast [10 x float]* [[B_PRIV]] to i8*
-  // TCHECK:  [[B_ADDR_REF_BCAST:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8*
-  // TCHECK:  call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_ADDR_REF_BCAST]], {{.+}})
-
-  // firstprivate(c)
-  // TCHECK:  [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8*
-  // TCHECK:  [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
-  // TCHECK:  call void @llvm.memcpy.{{.+}}(i8* [[C_PRIV_BCAST]], i8* [[C_IN_BCAST]],{{.+}})
-  
-  // firstprivate(d)
-  // TCHECK:  [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8*
-  // TCHECK:  [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
-  // TCHECK:  call void @llvm.memcpy.{{.+}}(i8* [[D_PRIV_BCAST]], i8* [[D_IN_BCAST]],{{.+}})
-
-  
-  #pragma omp target firstprivate(ptr)
+  //  b[2] = 1.0
+  // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[B_ADDR_REF]],
+  // TCHECK: store float {{.+}}, float* [[B_GEP]],
+
+  // c[1][2] = 1.0
+  // TCHECK:  [[B_GEP1:%.+]] = getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[C_ADDR_REF]],
+  // TCHECK:  [[B_GEP2:%.+]] = getelementptr inbounds [10 x double], [10 x double]* [[B_GEP1]],
+  // TCHECK:  store double {{.+}}, double* [[B_GEP2]],
+  
+  // d.X = 1; d.Y = 1
+  // TCHECK:  [[X:%.+]] = getelementptr inbounds [[TT]], [[TT]]* [[D_ADDR_REF]],
+  // TCHECK:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[X]]
+  // TCHECK:  [[Y:%.+]] = getelementptr inbounds [[TT]], [[TT]]* [[D_ADDR_REF]],
+  // TCHECK:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[Y]]
   {
     ptr[0]++;
   }
 
   // TCHECK:  define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
   // TCHECK:  [[PTR_ADDR:%.+]] = alloca double*,
-  // TCHECK:  [[PTR_PRIV:%.+]] = alloca double*,
   // TCHECK:  store double* [[PTR_IN]], double** [[PTR_ADDR]],
   // TCHECK:  [[PTR_IN_REF:%.+]] = load double*, double** [[PTR_ADDR]],
-  // TCHECK:  store double* [[PTR_IN_REF]], double** [[PTR_PRIV]],
+  // TCHECK:  [[PTR_GEP:%.+]] = getelementptr inbounds double, double* [[PTR_IN_REF]],
+  // TCHECK:  store double {{.+}}, double* [[PTR_GEP]],
 
   return a;
 }
 
 
 template<typename tx>
 tx ftemplate(int n) {
   tx a = 0;
#pragma omp target firstprivate(a,aaa,b)
 
   return a;
 }
 
 // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
 // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
-// TCHECK:  [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
-// TCHECK:  [[A3_PRIV:%.+]] = alloca i{{[0-9]+}},
-// TCHECK:  [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
 // TCHECK:  store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
 // TCHECK:  store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]],
 // TCHECK:  store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
 // TCHECK-64:  [[A_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
 // TCHECK:  [[A3_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A3_ADDR]] to i8*
 // TCHECK:  [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
 
-// firstprivate(a): a_priv = a_in
-// TCHECK-64:  [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_CONV]],
-// TCHECK-32:  [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
-// TCHECK:  store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A_PRIV]],
-
-// firstprivate(aaa)
-// TCHECK:  [[A3_IN_VAL:%.+]] = load i8, i8* [[A3_CONV]],
-// TCHECK:  store i{{[0-9]+}} [[A3_IN_VAL]], i{{[0-9]+}}* [[A3_PRIV]],
-
-// firstprivate(b)
-// TCHECK:  [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
-// TCHECK:  [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
-// TCHECK:  call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_IN_BCAST]],{{.+}})
+// a += 1
+// TCHECK-64:  {{.+}} = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_CONV]],
+// TCHECK-32:  {{.+}} = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK-64:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[A_CONV]],
+// TCHECK-32:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[A_ADDR]],
+
+// aaa += 1
+// TCHECK:  {{.+}} = load i{{[0-9]+}}, i{{[0-9]+}}* [[A3_CONV]],
+// TCHECK:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[A3_CONV]],
+
+// b[2] += 1
+// TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_ADDR_REF]],
+// TCHECK:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[B_GEP]],  
 
 // TCHECK:  ret void
 
 struct S1 {
   double a;
 
   int r1(int n){
     int b = n+1;
 
 #pragma omp target firstprivate(b)
     {
       this->a = (double)b + 1.5;
     }
 
     return (int)b;
   }
 
   // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]])
   // TCHECK:  [[TH_ADDR:%.+]] = alloca [[S1]]*,
   // TCHECK:  [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
-  // TCHECK:  [[B_PRIV:%.+]] = alloca i{{[0-9]+}},
 
   // TCHECK:  store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
   // TCHECK:  store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]],
   // TCHECK:  [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
   // TCHECK-64:  [[B_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}*
 
-  // firstprivate(b)
+  // this->a = (double)b + 1.5; (only check access to b)
   // TCHECK-64:  [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_ADDR_CONV]],
-  // TCHECK-32:  [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_ADDR]],
-  // TCHECK:  store i{{[0-9]+}} [[B_IN_VAL]], i{{[0-9]+}}* [[B_PRIV]], 
 
   // TCHECK: ret void
 };
 
 
 
 int bar(int n, double *ptr){
   int a = 0;
   a += foo(n, ptr);
   S1 S;
   a += S.r1(n);
   a += fstatic(n);
   a += ftemplate<int>(n);
 
   return a;
 }
 
 // template
 
 // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
 // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
-// TCHECK:  [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
-// TCHECK:  [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
 // TCHECK:  store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
 // TCHECK:  store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
 // TCHECK-64:  [[A_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
 // TCHECK:  [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
 
-// firstprivate(a)
+// a += 1
 // TCHECK-64:  [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR_CONV]]
 // TCHECK-32:  [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]]
-// TCHECK:  store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A_PRIV]],
+// TCHECK-64:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[A_ADDR_CONV]],
+// TCHECK-32:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[A_ADDR]],
 
-// firstprivate(b)
-// TCHECK:  [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
-// TCHECK:  [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
-// TCHECK:  call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_IN_BCAST]],{{.+}})
+// b[2] += 1
+// TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_ADDR_REF]],
+// TCHECK:  store i{{[0-9]+}} {{.+}}, i{{[0-9]+}}* [[B_GEP]],
 
 // TCHECK: ret void
 
 #endif