[OpenMP][Clang] Add call to global data sharing stack initialization on the workers side

[gtbercea]

Summary

The workers also need to initialize the global stack. The call to the initialization function needs to happen after the kernel_init() function is called by the master. This ensures that the per-team data structures of the runtime have been initialized.

[ABataev]

LG

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

void CGOpenMPRuntimeNVPTX::emitWorkerLoop(CodeGenFunction &CGF,
 
   CGF.EmitBranch(AwaitBB);
 
   // Workers wait for work from master.
   CGF.EmitBlock(AwaitBB);
   // Wait for parallel work
   syncCTAThreads(CGF);
 
+  // For data sharing, we need to initialize the stack for workers.
+  CGF.EmitRuntimeCall(
+      createNVPTXRuntimeFunction(
+          OMPRTL_NVPTX__kmpc_data_sharing_init_stack));
+
   Address WorkFn =
       CGF.CreateDefaultAlignTempAlloca(CGF.Int8PtrTy, /*Name=*/"work_fn");
   Address ExecStatus =
       CGF.CreateDefaultAlignTempAlloca(CGF.Int8Ty, /*Name=*/"exec_status");
   CGF.InitTempAlloca(ExecStatus, Bld.getInt8(/*C=*/0));
   CGF.InitTempAlloca(WorkFn, llvm::Constant::getNullValue(CGF.Int8PtrTy));
 
   // TODO: Optimize runtime initialization and pass in correct value.

test/OpenMP/nvptx_data_sharing.cpp

#pragma omp target
     #pragma omp parallel private(c)
     {
       int *c1 = &c;
       b = a + 10000;
     }
   }
 }
 
+/// ========= In the worker function ========= ///
+// CK1: {{.*}}define internal void @__omp_offloading{{.*}}test_ds{{.*}}_worker()
+// CK1: call void @llvm.nvvm.barrier0()
+// CK1: call void @__kmpc_data_sharing_init_stack
+
 /// ========= In the kernel function ========= ///
 
 // CK1: {{.*}}define void @__omp_offloading{{.*}}test_ds{{.*}}()
 // CK1: [[SHAREDARGS1:%.+]] = alloca i8**
 // CK1: [[SHAREDARGS2:%.+]] = alloca i8**
 // CK1: call void @__kmpc_kernel_init
 // CK1: call void @__kmpc_data_sharing_init_stack
 // CK1: [[GLOBALSTACK:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 8, i16 0)