diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 90c40a20954b..cabe8b941faf 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,12 @@
+2007-06-21  Jakub Jelinek  <jakub@redhat.com>
+
+	PR middle-end/32362
+	* omp-low.c (lookup_decl_in_outer_ctx): Don't ICE if t is NULL,
+	but decl is a global var, instead return decl.
+	* gimplify.c (gimplify_adjust_omp_clauses_1): Add shared clauses
+	even for is_global_var decls, if they are private in some outer
+	context.
+
 2007-06-21  Richard Guenther  <rguenther@suse.de>
 
 	PR tree-optimization/32451
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 01ccaf02a823..c93bb4449d5b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -4882,7 +4882,20 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
   else if (flags & GOVD_SHARED)
     {
       if (is_global_var (decl))
-	return 0;
+	{
+	  struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp->outer_context;
+	  while (ctx != NULL)
+	    {
+	      splay_tree_node on
+		= splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+	      if (on && (on->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
+				      | GOVD_PRIVATE | GOVD_REDUCTION)) != 0)
+		break;
+	      ctx = ctx->outer_context;
+	    }
+	  if (ctx == NULL)
+	    return 0;
+	}
       code = OMP_CLAUSE_SHARED;
     }
   else if (flags & GOVD_PRIVATE)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 5b1f3c4c3d8c..700645f6551d 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1504,9 +1504,9 @@ lookup_decl_in_outer_ctx (tree decl, omp_context *ctx)
   for (up = ctx->outer, t = NULL; up && t == NULL; up = up->outer)
     t = maybe_lookup_decl (decl, up);
 
-  gcc_assert (t);
+  gcc_assert (t || is_global_var (decl));
 
-  return t;
+  return t ? t : decl;
 }
 
 
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 83c841c51250..a4abc445a879 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,10 @@
+2007-06-21  Jakub Jelinek  <jakub@redhat.com>
+
+	PR middle-end/32362
+	* testsuite/libgomp.c/pr32362-1.c: New test.
+	* testsuite/libgomp.c/pr32362-2.c: New test.
+	* testsuite/libgomp.c/pr32362-3.c: New test.
+
 2007-06-07  Jakub Jelinek  <jakub@redhat.com>
 
 	* team.c (gomp_team_start): Fix setting up thread_attr
diff --git a/libgomp/testsuite/libgomp.c/pr32362-1.c b/libgomp/testsuite/libgomp.c/pr32362-1.c
new file mode 100644
index 000000000000..3c62d4bdb269
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr32362-1.c
@@ -0,0 +1,32 @@
+/* PR middle-end/32362 */
+/* { dg-do run } */
+/* { dg-options "-O2" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int n[4] = { -1, -1, -1, -1 };
+  static int a = 2, b = 4;
+  omp_set_num_threads (4);
+  omp_set_dynamic (0);
+  omp_set_nested (1);
+#pragma omp parallel private(b)
+  {
+    b = omp_get_thread_num ();
+#pragma omp parallel firstprivate(a)
+    {
+      a = (omp_get_thread_num () + a) + 1;
+      if (b == omp_get_thread_num ())
+	n[omp_get_thread_num ()] = a + (b << 4);
+    }
+  }
+  if (n[0] != 3)
+    abort ();
+  if (n[3] != -1
+      && (n[1] != 0x14 || n[2] != 0x25 || n[3] != 0x36))
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr32362-2.c b/libgomp/testsuite/libgomp.c/pr32362-2.c
new file mode 100644
index 000000000000..43f36e0e98be
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr32362-2.c
@@ -0,0 +1,33 @@
+/* PR middle-end/32362 */
+/* { dg-do run } */
+/* { dg-options "-O2" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int a = 2, b = 4;
+
+int
+main ()
+{
+  int n[4] = { -1, -1, -1, -1 };
+  omp_set_num_threads (4);
+  omp_set_dynamic (0);
+  omp_set_nested (1);
+#pragma omp parallel private(b)
+  {
+    b = omp_get_thread_num ();
+#pragma omp parallel firstprivate(a)
+    {
+      a = (omp_get_thread_num () + a) + 1;
+      if (b == omp_get_thread_num ())
+	n[omp_get_thread_num ()] = a + (b << 4);
+    }
+  }
+  if (n[0] != 3)
+    abort ();
+  if (n[3] != -1
+      && (n[1] != 0x14 || n[2] != 0x25 || n[3] != 0x36))
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr32362-3.c b/libgomp/testsuite/libgomp.c/pr32362-3.c
new file mode 100644
index 000000000000..09a88f52a3b6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr32362-3.c
@@ -0,0 +1,34 @@
+/* PR middle-end/32362 */
+/* { dg-do run } */
+/* { dg-options "-O2" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int a = 2;
+
+int
+main ()
+{
+  int n[4] = { -1, -1, -1, -1 };
+  int b = 4;
+  omp_set_num_threads (4);
+  omp_set_dynamic (0);
+  omp_set_nested (1);
+#pragma omp parallel private(b)
+  {
+    b = omp_get_thread_num ();
+#pragma omp parallel firstprivate(a)
+    {
+      a = (omp_get_thread_num () + a) + 1;
+      if (b == omp_get_thread_num ())
+	n[omp_get_thread_num ()] = a + (b << 4);
+    }
+  }
+  if (n[0] != 3)
+    abort ();
+  if (n[3] != -1
+      && (n[1] != 0x14 || n[2] != 0x25 || n[3] != 0x36))
+    abort ();
+  return 0;
+}