diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index fad50da8fbc4..30757877eafe 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3477,18 +3477,30 @@ annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees,
       /* Do not try to add automatic OpenACC annotations inside manually
         annotated loops.  Presumably, the user avoided doing it on
         purpose; for example, all available levels of parallelism may
-        have been used up.  */
-      {
-       struct annotation_info nested_info
-         = { NULL_TREE, NULL_TREE, false, as_explicit_annotation,
-             node, info };
-       if (info->state >= as_in_kernels_region)
-         do_not_annotate_loop_nest (info, as_explicit_annotation,
-                                    node);
-       walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions,
-                  (void *) &nested_info, NULL);
-       *walk_subtrees = 0;
-      }
+        have been used up.  However, assume that the combined construct
+        "#pragma acc kernels loop" means to try to process the whole
+        loop nest.
+        Note that a single OACC_LOOP construct represents an entire set
+        of collapsed loops so we do not have to deal explicitly with the
+        collapse clause here, as the Fortran front end does.  */
+      if (info->state == as_in_kernels_region && OACC_LOOP_COMBINED (node))
+       {
+         walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions,
+                    (void *) info, NULL);
+         *walk_subtrees = 0;
+       }
+      else
+       {
+         struct annotation_info nested_info
+           = { NULL_TREE, NULL_TREE, false, as_explicit_annotation,
+               node, info };
+         if (info->state >= as_in_kernels_region)
+           do_not_annotate_loop_nest (info, as_explicit_annotation,
+                                      node);
+         walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions,
+                    (void *) &nested_info, NULL);
+         *walk_subtrees = 0;
+       }
       break;

     case FOR_STMT:
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
index c2a3c57b48b8..2519f23d49f0 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-directives.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -110,7 +110,7 @@ test ()
 // { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } }
-// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop auto" 6 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c
new file mode 100644
index 000000000000..89ec6447625f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c
@@ -0,0 +1,18 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test that "acc kernels loop" directive causes annotation of the entire
+   loop nest.  */
+
+void f (float *a, float *b)
+{
+#pragma acc kernels loop
+  for (int k = 0; k < 20; k++)
+    for (int l = 0; l < 20; l++)
+      for (int m = 0; m < 20; m++)
+       b[m] = a[m];
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 2 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c
new file mode 100644
index 000000000000..77a3b7a9136d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c
@@ -0,0 +1,19 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test that "acc kernels loop" directive causes annotation of the entire
+   loop nest in the presence of a collapse clause.  */
+
+void f (float *a, float *b)
+{
+#pragma acc kernels loop collapse(2)
+  for (int k = 0; k < 20; k++)
+    for (int l = 0; l < 20; l++)
+      for (int m = 0; m < 20; m++)
+       b[m] = a[m];
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop collapse.2." 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */
