diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index 1dec7caa0d16edad8352118fa3b376da237de8ff..5550ce255135760d50420352801ea228a74ab642 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -1881,6 +1881,10 @@
   ""
 {
   emit_insn (gen_omp_simt_exit (Pmode, operands[0]));
+  if (TARGET_PTX_6_0)
+    emit_insn (gen_nvptx_warpsync ());
+  else
+    emit_insn (gen_nvptx_uniform_warp_check ());
   DONE;
 })
 
diff --git a/libgomp/testsuite/libgomp.c/pr104783-2.c b/libgomp/testsuite/libgomp.c/pr104783-2.c
new file mode 100644
index 0000000000000000000000000000000000000000..8750d915d01e642d8a43687168ffdc93235d962e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr104783-2.c
@@ -0,0 +1,25 @@
+#define N (32 * 32)
+
+#define TYPE float
+#define VAR v
+#define INIT 0.0
+#define UPDATE + 1.0
+#define EXPECTED N
+
+int
+main (void)
+{
+  TYPE VAR = INIT;
+  #pragma omp target map(tofrom: VAR)
+  #pragma omp parallel for simd
+  for (int i = 0 ; i < N; i++)
+    {
+      #pragma omp atomic update
+      VAR = VAR UPDATE;
+    }
+
+  if (VAR != EXPECTED)
+    __builtin_abort ();
+
+  return 0;
+}