diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -276,15 +276,20 @@
     for (GlobalVariable &GV : M.getGlobalList())
       if (GV.getValueType() == OpenMPIRBuilder::Ident && GV.hasInitializer())
         if (GV.getInitializer() == Initializer)
-          return Ident = &GV;
-
-    auto *GV = new GlobalVariable(M, OpenMPIRBuilder::Ident,
-                                  /* isConstant = */ true,
-                                  GlobalValue::PrivateLinkage, Initializer);
-    GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
-    GV->setAlignment(Align(8));
-    Ident = GV;
+          Ident = &GV;
+
+    if (!Ident) {
+      auto *GV = new GlobalVariable(
+          M, OpenMPIRBuilder::Ident,
+          /* isConstant = */ true, GlobalValue::PrivateLinkage, Initializer, "",
+          nullptr, GlobalValue::NotThreadLocal,
+          M.getDataLayout().getDefaultGlobalsAddressSpace());
+      GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+      GV->setAlignment(Align(8));
+      Ident = GV;
+    }
   }
+
   return Builder.CreatePointerCast(Ident, IdentPtr);
 }
 
diff --git a/openmp/libomptarget/test/offloading/bug51982.c b/openmp/libomptarget/test/offloading/bug51982.c
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/bug51982.c
@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compile-generic -O1 && %libomptarget-run-generic
+// -O1 to run openmp-opt
+
+int main(void) {
+  long int aa = 0;
+
+  int ng = 12;
+  int nxyz = 5;
+
+  const long exp = ng * nxyz;
+
+#pragma omp target map(tofrom : aa)
+  for (int gid = 0; gid < nxyz; gid++) {
+#pragma omp parallel for
+    for (unsigned int g = 0; g < ng; g++) {
+#pragma omp atomic
+      aa += 1;
+    }
+  }
+  if (aa != exp) {
+    return 1;
+  }
+  return 0;
+}