[OPENMP]Add option -fopenmp-cuda-const-firstprivate to control address space of the corresponding global.

[ABataev]

Summary

Added options -f-[no]openmp-cuda-const-firstprivate to control the
address space of the the corresponding global variable, created for the
target firstprivate variable. Disabled by default.

[jdoerfert]

Can you please show me a test case or explain to me when/how this global is actually used.

[ABataev]

In D99689#2662848, @jdoerfert wrote:

Can you please show me a test case or explain to me when/how this global is actually used.

It is passed as an argument to the target region. When libomptarget requests the memory for the firstprivate, it returns the pointer to this const global, which then passed as argument to the kernel.

[jdoerfert]

In D99689#2662852, @ABataev wrote:

In D99689#2662848, @jdoerfert wrote:

Can you please show me a test case or explain to me when/how this global is actually used.

It is passed as an argument to the target region. When libomptarget requests the memory for the firstprivate, it returns the pointer to this const global, which then passed as argument to the kernel.

So if we use it, why would we disable it?

[ABataev]

In D99689#2662856, @jdoerfert wrote:

In D99689#2662852, @ABataev wrote:

In D99689#2662848, @jdoerfert wrote:

Can you please show me a test case or explain to me when/how this global is actually used.

It is passed as an argument to the target region. When libomptarget requests the memory for the firstprivate, it returns the pointer to this const global, which then passed as argument to the kernel.

So if we use it, why would we disable it?

With this new option you can control how to handle it. You can either dynamically allocate memory using libomptarget memmanager (default for this option) or use preallocated constant memory, if you're not going to remove the var constantness.

[jdoerfert]

In D99689#2662860, @ABataev wrote:

In D99689#2662856, @jdoerfert wrote:

In D99689#2662852, @ABataev wrote:

In D99689#2662848, @jdoerfert wrote:

Can you please show me a test case or explain to me when/how this global is actually used.

It is passed as an argument to the target region. When libomptarget requests the memory for the firstprivate, it returns the pointer to this const global, which then passed as argument to the kernel.

So if we use it, why would we disable it?

With this new option you can control how to handle it. You can either dynamically allocate memory using libomptarget memmanager (default for this option) or use preallocated constant memory, if you're not going to remove the var constantness.

I get what the new option does, what I want to know is why we would ever want to disable the constant memory usage. Is it potentially slower or otherwise problematic?
Also, I am not sure but I imagine the generated code would be better if we would use the constant global directly, or, add the address space to the corresponding kernel argument.

[ABataev]

In D99689#2662875, @jdoerfert wrote:

In D99689#2662860, @ABataev wrote:

In D99689#2662856, @jdoerfert wrote:

In D99689#2662852, @ABataev wrote:

In D99689#2662848, @jdoerfert wrote:

Can you please show me a test case or explain to me when/how this global is actually used.

It is passed as an argument to the target region. When libomptarget requests the memory for the firstprivate, it returns the pointer to this const global, which then passed as argument to the kernel.

So if we use it, why would we disable it?

With this new option you can control how to handle it. You can either dynamically allocate memory using libomptarget memmanager (default for this option) or use preallocated constant memory, if you're not going to remove the var constantness.

I get what the new option does, what I want to know is why we would ever want to disable the constant memory usage. Is it potentially slower or otherwise problematic?

Because you mentioned that it would not work if we would like to use const_cast to remove constness. It should be faster than the global memory, that's why I introduced this feature.

Also, I am not sure but I imagine the generated code would be better if we would use the constant global directly, or, add the address space to the corresponding kernel argument.

Ye, but it will require some rework in the libomptarget and the compiler itself. Instead, I just reused the existing mechanism to allocate the memory in the constant pool and use it instead of the global.

[ABataev]

Rebase + update

[ABataev]

Rebase

clang/include/clang/Basic/LangOptions.def

 LANGOPT(OpenMPCUDAMode    , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode")
 LANGOPT(OpenMPIRBuilder   , 1, 0, "Use the experimental OpenMP-IR-Builder codegen path.")
 LANGOPT(OpenMPCUDAForceFullRuntime , 1, 0, "Force to use full runtime in all constructs when offloading to CUDA devices")
 LANGOPT(OpenMPCUDANumSMs  , 32, 0, "Number of SMs for CUDA devices.")
 LANGOPT(OpenMPCUDABlocksPerSM  , 32, 0, "Number of blocks per SM for CUDA devices.")
 LANGOPT(OpenMPCUDAReductionBufNum , 32, 1024, "Number of the reduction records in the intermediate reduction buffer used for the teams reductions.")
 LANGOPT(OpenMPOptimisticCollapse  , 1, 0, "Use at most 32 bits to represent the collapsed loop nest counter.")
 LANGOPT(OpenMPCUDATargetParallel, 1, 0, "Support parallel execution of target region on Cuda-based devices.")
+LANGOPT(OpenMPCUDAConstFirstprivate, 1, 0, "Place firstprivate const variables in the constant address space.")
 LANGOPT(RenderScript      , 1, 0, "RenderScript")
 
 LANGOPT(CUDAIsDevice      , 1, 0, "compiling for CUDA device")
 LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "allowing variadic functions in CUDA device code")
 LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
 LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental functions")
 LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code")
 LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP")

clang/include/clang/Driver/Options.td

defm openmp_optimistic_collapse : BoolFOption<"openmp-optimistic-collapse",
   LangOpts<"OpenMPOptimisticCollapse">, DefaultFalse,
   PosFlag<SetTrue, [CC1Option]>, NegFlag<SetFalse>, BothFlags<[NoArgumentUnused, HelpHidden]>>;
 def fopenmp_cuda_parallel_target_regions : Flag<["-"], "fopenmp-cuda-parallel-target-regions">, Group<f_Group>,
   Flags<[CC1Option, NoArgumentUnused, HelpHidden]>,
   HelpText<"Support parallel execution of target regions on Cuda-based devices.">;
 def fno_openmp_cuda_parallel_target_regions : Flag<["-"], "fno-openmp-cuda-parallel-target-regions">, Group<f_Group>,
   Flags<[NoArgumentUnused, HelpHidden]>,
   HelpText<"Support only serial execution of target regions on Cuda-based devices.">;
+def fopenmp_cuda_const_firstprivate : Flag<["-"], "fopenmp-cuda-const-firstprivate">, Group<f_Group>,
+  Flags<[CC1Option, NoArgumentUnused, HelpHidden]>,
+  HelpText<"Allow placement of firstprivate const variables in the constant address space.">;
+def fno_openmp_cuda_const_firstprivate : Flag<["-"], "fno-openmp-cuda-const-firstprivate">, Group<f_Group>,
+  Flags<[NoArgumentUnused, HelpHidden]>,
+  HelpText<"Do not allow placement of firstprivate const variables in the constant address space.">;
 def static_openmp: Flag<["-"], "static-openmp">,
   HelpText<"Use the static host OpenMP runtime while linking.">;
 def fno_optimize_sibling_calls : Flag<["-"], "fno-optimize-sibling-calls">, Group<f_Group>;
 def foptimize_sibling_calls : Flag<["-"], "foptimize-sibling-calls">, Group<f_Group>;
 defm escaping_block_tail_calls : BoolFOption<"escaping-block-tail-calls",
   CodeGenOpts<"NoEscapingBlockTailCalls">, DefaultFalse,
   NegFlag<SetTrue, [CC1Option]>, PosFlag<SetFalse>>;
 def force__cpusubtype__ALL : Flag<["-"], "force_cpusubtype_ALL">;

clang/lib/CodeGen/CGOpenMPRuntime.cpp

if (CI.capturesThis()) {
       CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
           CGF.getTypeSize(ElementType), CGF.Int64Ty, /*isSigned=*/true));
       // The default map type for a scalar/complex type is 'to' because by
       // default the value doesn't have to be retrieved. For an aggregate
       // type, the default is 'tofrom'.
       CombinedInfo.Types.push_back(getMapModifiersForPrivateClauses(CI));
       const VarDecl *VD = CI.getCapturedVar();
       auto I = FirstPrivateDecls.find(VD);
-      if (I != FirstPrivateDecls.end() &&
+      if (CGF.getLangOpts().OpenMPCUDAConstFirstprivate &&
+          I != FirstPrivateDecls.end() &&
           VD->getType().isConstant(CGF.getContext())) {
         llvm::Constant *Addr =
             CGF.CGM.getOpenMPRuntime().registerTargetFirstprivateCopy(CGF, VD);
         // Copy the value of the original variable to the new global copy.
         CGF.Builder.CreateMemCpy(
             CGF.MakeNaturalAlignAddrLValue(Addr, ElementType).getAddress(CGF),
             Address(CV, CGF.getContext().getTypeAlignInChars(ElementType)),
             CombinedInfo.Sizes.back(), /*IsVolatile=*/false);

clang/lib/CodeGen/CGStmtOpenMP.cpp

 }
 
 bool CodeGenFunction::EmitOMPFirstprivateClause(const OMPExecutableDirective &D,
                                                 OMPPrivateScope &PrivateScope) {
   if (!HaveInsertPoint())
     return false;
   bool DeviceConstTarget =
       getLangOpts().OpenMPIsDevice &&
+      getLangOpts().OpenMPCUDAConstFirstprivate &&
       isOpenMPTargetExecutionDirective(D.getDirectiveKind());
   bool FirstprivateIsLastprivate = false;
   llvm::DenseMap<const VarDecl *, OpenMPLastprivateModifier> Lastprivates;
   for (const auto *C : D.getClausesOfKind<OMPLastprivateClause>()) {
     for (const auto *D : C->varlists())
       Lastprivates.try_emplace(
           cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl(),
           C->getKind());

clang/lib/Driver/ToolChains/Clang.cpp

case Driver::OMPRT_IOMP5:
 
       // When in OpenMP offloading mode with NVPTX target, forward
       // cuda-mode flag
       if (Args.hasFlag(options::OPT_fopenmp_cuda_mode,
                        options::OPT_fno_openmp_cuda_mode, /*Default=*/false))
         CmdArgs.push_back("-fopenmp-cuda-mode");
 
       // When in OpenMP offloading mode with NVPTX target, forward
+      // cuda-const-firstprivate flag
+      if (Args.hasFlag(options::OPT_fopenmp_cuda_const_firstprivate,
+                       options::OPT_fno_openmp_cuda_const_firstprivate,
+                       /*Default=*/false))
+        CmdArgs.push_back("-fopenmp-cuda-const-firstprivate");
+
+      // When in OpenMP offloading mode with NVPTX target, forward
       // cuda-parallel-target-regions flag
       if (Args.hasFlag(options::OPT_fopenmp_cuda_parallel_target_regions,
                        options::OPT_fno_openmp_cuda_parallel_target_regions,
                        /*Default=*/true))
         CmdArgs.push_back("-fopenmp-cuda-parallel-target-regions");
 
       // When in OpenMP offloading mode with NVPTX target, check if full runtime
       // is required.

clang/lib/Frontend/CompilerInvocation.cpp

#undef LANG_OPTION_WITH_MARSHALLING
   }
 
   if (!Opts.OMPHostIRFile.empty())
     GenerateArg(Args, OPT_fopenmp_host_ir_file_path, Opts.OMPHostIRFile, SA);
 
   if (Opts.OpenMPCUDAMode)
     GenerateArg(Args, OPT_fopenmp_cuda_mode, SA);
 
+  if (Opts.OpenMPCUDAConstFirstprivate)
+    GenerateArg(Args, OPT_fopenmp_cuda_const_firstprivate, SA);
+
   if (Opts.OpenMPCUDATargetParallel)
     GenerateArg(Args, OPT_fopenmp_cuda_parallel_target_regions, SA);
 
   if (Opts.OpenMPCUDAForceFullRuntime)
     GenerateArg(Args, OPT_fopenmp_cuda_force_full_runtime, SA);
 
   // The arguments used to set Optimize, OptimizeSize and NoInlineDefine are
   // generated from CodeGenOptions.
if (Arg *A = Args.getLastArg(options::OPT_fopenmp_host_ir_file_path)) {
     if (!llvm::sys::fs::exists(Opts.OMPHostIRFile))
       Diags.Report(diag::err_drv_omp_host_ir_file_not_found)
           << Opts.OMPHostIRFile;
   }
 
   // Set CUDA mode for OpenMP target NVPTX/AMDGCN if specified in options
   Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) &&
                         Args.hasArg(options::OPT_fopenmp_cuda_mode);
+  Opts.OpenMPCUDAConstFirstprivate =
+      Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) &&
+      Args.hasArg(options::OPT_fopenmp_cuda_const_firstprivate);
 
   // Set CUDA support for parallel execution of target regions for OpenMP target
   // NVPTX/AMDGCN if specified in options.
   Opts.OpenMPCUDATargetParallel =
       Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) &&
       Args.hasArg(options::OPT_fopenmp_cuda_parallel_target_regions);
 
   // Set CUDA mode for OpenMP target NVPTX/AMDGCN if specified in options

clang/test/Driver/openmp-offload-gpu.c

 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-mode -fno-openmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-mode -fno-openmp-cuda-mode 2>&1 \
 // RUN:   | FileCheck -check-prefix=NO_CUDA_MODE %s
 // NO_CUDA_MODE-NOT: "-{{fno-|f}}openmp-cuda-mode"
 
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=CUDA_CONST_FP %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-const-firstprivate -fopenmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=CUDA_CONST_FP %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=CUDA_CONST_FP %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-const-firstprivate -fopenmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=CUDA_CONST_FP %s
+// CUDA_CONST_FP: clang{{.*}}"-cc1"{{.*}}"-triple" "{{nvptx64-nvidia-cuda|amdgcn-amd-amdhsa}}"
+// CUDA_CONST_FP-SAME: "-fopenmp-cuda-const-firstprivate"
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_CUDA_CONST_FP %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-const-firstprivate -fno-openmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_CUDA_CONST_FP %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_CUDA_CONST_FP %s
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-const-firstprivate -fno-openmp-cuda-const-firstprivate 2>&1 \
+// RUN:   | FileCheck -check-prefix=NO_CUDA_CONST_FP %s
+// NO_CUDA_CONST_FP-NOT: "-{{fno-|f}}openmp-cuda-const-firstprivate"
+
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-force-full-runtime -fopenmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s
 // RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-force-full-runtime -fopenmp-cuda-force-full-runtime 2>&1 \
 // RUN:   | FileCheck -check-prefix=FULL_RUNTIME %s

clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp

 // Test target codegen - host bc file has to be created first.
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 --check-prefix TCHECK-NONCONST
+// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-cuda-const-firstprivate -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 --check-prefix TCHECK-CONST
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
-// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 --check-prefix TCHECK-NONCONST
+// RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-cuda-const-firstprivate -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 --check-prefix TCHECK-CONST
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 template <typename tx, typename ty>
 struct TT {
   tx X;
   ty Y;
 };
 
 // TCHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // TCHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
 // TCHECK-DAG:  [[S1:%.+]] = type { double }
 
-// TCHECK: @__omp_offloading_firstprivate__{{.+}}_e_l27 = internal addrspace(4) global [[TTII]] zeroinitializer
+// TCHECK-NONCONST-NOT: @__omp_offloading_firstprivate__{{.+}}_e_l30
+// TCHECK-CONST:    @__omp_offloading_firstprivate__{{.+}}_e_l30 = internal addrspace(4) global [[TTII]] zeroinitializer
 int foo(int n, double *ptr) {
   int a = 0;
   short aa = 0;
   float b[10];
   double c[5][10];
   TT<long long, char> d;
   const TT<int, int> e = {n, n};
 
 #pragma omp target firstprivate(a, e) map(tofrom \
                                           : b)
   {
     b[a] = a;
     b[a] += e.X;
   }
 
   // TCHECK:  define {{.*}}void @__omp_offloading_{{.+}}([10 x float] addrspace(1)* noalias [[B_IN:%.+]], i{{[0-9]+}} [[A_IN:%.+]], [[TTII]]* noalias [[E_IN:%.+]])
   // TCHECK-NOT: alloca [[TTII]],
   // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
-  // TCHECK-NOT: alloca [[TTII]],
+  // TCHECK-NONCONST: alloca [[TTII]],
+  // TCHECK-CONST-NOT: alloca [[TTII]],
   // TCHECK-NOT: alloca i{{[0-9]+}},
   // TCHECK-64:  call void @llvm.dbg.declare(metadata [10 x float] addrspace(1)** %{{.+}}, metadata !{{[0-9]+}}, metadata !DIExpression())
   // TCHECK:  store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
   // TCHECK:  ret void
 
 #pragma omp target firstprivate(aa, b, c, d)
   {
     aa += 1;

clang/test/OpenMP/target_firstprivate_codegen.cpp

 // CHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // CHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
 // CHECK-DAG:  [[S1:%.+]] = type { double }
 
 // TCHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // TCHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
 // TCHECK-DAG:  [[S1:%.+]] = type { double }
 
-// CHECK-DAG:  [[FP_E:@__omp_offloading_firstprivate_.+_e_l76]] = internal global [[TTII]] zeroinitializer
 // CHECK-DAG:  [[SIZET:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ:32|64]] 4, i{{64|32}} {{8|4}}]
 // CHECK-DAG:  [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 288, i64 49]
 // CHECK-DAG:  [[MAPT2:@.+]] = private unnamed_addr constant [9 x i64] [i64 288, i64 161, i64 800, i64 161, i64 161, i64 800, i64 800, i64 161, i64 161]
 // CHECK-DAG:  [[SIZET3:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i{{32|64}} 0, i{{32|64}} 8]
 // CHECK-DAG:  [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 37]
 // CHECK-DAG:  [[MAPT4:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 800, i64 800, i64 161]
 // CHECK-DAG:  [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG:  [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 161]

 
 #pragma omp target firstprivate(ptr, e)
   {
     ptr[0] = e.X;
     ptr[0]++;
   }
   // CHECK:  [[PTR_ADDR_REF:%.+]] = load double*, double** [[PTR_ADDR]],
 
-  // CHECK:  [[E_BC:%.+]] = bitcast [[TTII]]* [[E:%.+]] to i8*
-  // CHECK:  call void @llvm.memcpy.p0i8.p0i8.i{{64|32}}(i8* {{.*}} bitcast ([[TTII]]* [[FP_E]] to i8*), i8* {{.*}} [[E_BC]], i{{64|32}} 8, i1 false)
   // CHECK:  [[BASE_PTR_GEP3_0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP3_0]] to double**
   // CHECK:  store double* [[PTR_ADDR_REF]], double** [[BCAST_TOPTR]],
   // CHECK:  [[PTR_GEP3_0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTR_GEP3_0]] to double**
   // CHECK:  store double* [[PTR_ADDR_REF]], double** [[BCAST_TOPTR]],
   // CHECK:  [[BASE_PTR_GEP3_1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP3_1]] to [[TTII]]**
-  // CHECK:  store [[TTII]]* [[FP_E]], [[TTII]]** [[BCAST_TOPTR]],
+  // CHECK:  store [[TTII]]* [[FP_E:%.+]], [[TTII]]** [[BCAST_TOPTR]],
   // CHECK:  [[PTR_GEP3_1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
   // CHECK:  [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTR_GEP3_1]] to [[TTII]]**
   // CHECK:  store [[TTII]]* [[FP_E]], [[TTII]]** [[BCAST_TOPTR]],
 
   // CHECK:  [[BASE_PTR_GEP_ARG3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK:  [[PTR_GEP_ARG3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
   // CHECK: {{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, {{.+}}, i32 2, i8** [[BASE_PTR_GEP_ARG3]], i8** [[PTR_GEP_ARG3]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET3]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT3]], i32 0, i32 0), i8** null, i8** null)
 
   // TCHECK:  define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]], [[TTII]]* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[E:%.+]])
-  // TCHECK-NOT: alloca [[TTII]],
   // TCHECK:  [[PTR_ADDR:%.+]] = alloca double*,
-  // TCHECK-NOT: alloca [[TTII]],
+  // TCHECK:  alloca [[TTII]],
   // TCHECK-NOT: alloca double*,
   // TCHECK:  store double* [[PTR_IN]], double** [[PTR_ADDR]],
   // TCHECK-NOT: store double* %
 
   return a;
 }
 
 template <typename tx>