diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp
index 2f8c8a6..5f3b641 100644
--- gcc/c-family/ChangeLog.gomp
+++ gcc/c-family/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2014-11-05  James Norris  <jnorris@codesourcery.com>
+
+	* c-pragma.c (oacc_pragmas): Add "cache".
+
 2014-11-03  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* c-pragma.c (oacc_pragmas): Add entries for PRAGMA_OACC_ENTER_DATA
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index a28727e..e98b555 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1181,6 +1181,7 @@ static vec<pragma_ns_name> registered_pp_pragmas;
 
 struct omp_pragma_def { const char *name; unsigned int id; };
 static const struct omp_pragma_def oacc_pragmas[] = {
+  { "cache", PRAGMA_OACC_CACHE },
   { "data", PRAGMA_OACC_DATA },
   { "enter", PRAGMA_OACC_ENTER_DATA },
   { "exit", PRAGMA_OACC_EXIT_DATA },
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index f4e2010..7acd7b3 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2014-11-05  James Norris  <jnorris@codesourcery.com>
+
+	* c-parser.c (c_parser_omp_variable_list): Handle
+	OMP_NO_CLAUSE_CACHE.
+	(c_parser_oacc_cache): New function.
+	(c_parser_omp_construct): Use it for PRAGMA_OACC_CACHE.
+
 2014-11-03  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* c-parser.c (c_parser_oacc_enter_exit_data): New function.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 6ac1ace..410b19f 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -10053,6 +10053,14 @@ c_parser_omp_variable_list (c_parser *parser,
 	{
 	  switch (kind)
 	    {
+	    case OMP_NO_CLAUSE_CACHE:
+	      if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE)
+		{
+		  c_parser_error (parser, "expected %<[%>");
+		  t = error_mark_node;
+		  break;
+		}
+	      /* FALL THROUGH.  */
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    case OMP_CLAUSE_TO:
@@ -10091,6 +10099,29 @@ c_parser_omp_variable_list (c_parser *parser,
 		      t = error_mark_node;
 		      break;
 		    }
+
+		  if (kind == OMP_NO_CLAUSE_CACHE)
+		    {
+		      mark_exp_read (low_bound);
+		      mark_exp_read (length);
+
+		      if (TREE_CODE (low_bound) != INTEGER_CST
+			  && !TREE_READONLY (low_bound))
+			{
+			  error_at (clause_loc,
+					"%qD is not a constant", low_bound);
+			  t = error_mark_node;
+			}
+
+		      if (TREE_CODE (length) != INTEGER_CST
+			  && !TREE_READONLY (length))
+			{
+			  error_at (clause_loc,
+					"%qD is not a constant", length);
+			  t = error_mark_node;
+			}
+		    }
+
 		  t = tree_cons (low_bound, length, t);
 		}
 	      break;
@@ -11864,6 +11895,21 @@ c_parser_omp_structured_block (c_parser *parser)
 }
 
 /* OpenACC 2.0:
+   # pragma acc cache (variable-list) new-line
+
+   LOC is the location of the #pragma token.
+*/
+
+static tree
+c_parser_oacc_cache (location_t loc __attribute__((unused)), c_parser *parser)
+{
+  c_parser_omp_var_list_parens (parser, OMP_NO_CLAUSE_CACHE, NULL);
+  c_parser_skip_to_pragma_eol (parser);
+
+  return NULL_TREE;
+}
+
+/* OpenACC 2.0:
    # pragma acc data oacc-data-clause[optseq] new-line
      structured-block
 
@@ -14506,6 +14552,10 @@ c_parser_omp_construct (c_parser *parser)
 
   switch (p_kind)
     {
+    case PRAGMA_OACC_CACHE:
+      strcpy (p_name, "#pragma acc");
+      stmt = c_parser_oacc_cache (loc, parser);
+      break;
     case PRAGMA_OACC_DATA:
       stmt = c_parser_oacc_data (loc, parser);
       break;
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 8dc947d..4ac348d 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-11-05  James Norris  <jnorris@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c/cache-1.c: New file.
+	* testsuite/libgomp.oacc-c++/cache-1.C: Likewise.
+
 2014-11-05  Thomas Schwinge  <thomas@codesourcery.com>
 	    James Norris  <jnorris@codesourcery.com>
 
diff --git libgomp/testsuite/libgomp.oacc-c++/cache-1.C libgomp/testsuite/libgomp.oacc-c++/cache-1.C
new file mode 100644
index 0000000..eb94a68
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c++/cache-1.C
@@ -0,0 +1,50 @@
+/* { dg-do compile } */
+
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+#define N   2
+    int a[N], b[N];
+    int i;
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3;
+        b[i] = 0;
+    }
+
+#pragma acc parallel copyin (a[0:N]) copyout (b[0:N])
+{
+    int ii;
+
+    for (ii = 0; ii < N; ii++)
+    {
+        const int idx = ii;
+        int n = 1;
+        const int len = n;
+
+#pragma acc cache /* { dg-error "error: expected '\\(' before end of line" } */
+
+#pragma acc cache (a) /* { dg-error "error: expected '\\\['" } */
+
+#pragma acc cache (a[0:N]) copyin (a[0:N]) /* { dg-error "error: expected end of line before 'copyin'" } */
+
+#pragma acc cache () /* { dg-error "error: expected unqualified-id before '\\)' token" } */
+
+#pragma acc cache (a[0:N] b[0:N}) /* { dg-error "error: expected end of line before '\\\}' token|error: expected '\\)' before 'b'" } */
+
+        b[ii] = a[ii];
+    }
+}
+
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != b[i])
+            abort ();
+    }
+
+    return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c/cache-1.c libgomp/testsuite/libgomp.oacc-c/cache-1.c
new file mode 100644
index 0000000..5473475
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/cache-1.c
@@ -0,0 +1,69 @@
+/* { dg-do compile } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+#define N   2
+    int a[N], b[N];
+    int i;
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3;
+        b[i] = 0;
+    }
+
+#pragma acc parallel copyin (a[0:N]) copyout (b[0:N])
+{
+    int ii;
+
+    for (ii = 0; ii < N; ii++)
+    {
+        const int idx = ii;
+        int n = 1;
+        const int len = n;
+
+#pragma acc cache /* { dg-error "error: expected '\\(' before end of line" } */
+
+#pragma acc cache (a) /* { dg-error "error: expected '\\\[' before '\\)' token" } */
+
+#pragma acc cache (a[0:N]) copyin (a[0:N]) /* { dg-error "error: expected end of line before 'copyin'" } */
+
+#pragma acc cache () /* { dg-error "error: expected identifier before '\\)' token" } */
+
+#pragma acc cache (a[0:N] b[0:N}) /* { dg-error "error: expected '\\)' before 'b'" } */
+
+#pragma acc cache (a[0:N] /* { dg-error "error: expected '\\)' before end of line" } */
+
+#pragma acc cache (a[ii]) /* { dg-error "error: 'ii' is not a constant" } */
+
+#pragma acc cache (a[idx:n]) /* { dg-error "error: 'n' is not a constant" } */
+
+#pragma acc cache (a[0:N])
+
+#pragma acc cache (a[0:N], b[0:N])
+
+#pragma acc cache (a[0])
+
+#pragma acc cache (a[0], a[1], b[0:N])
+
+#pragma acc cache (a[idx])
+
+#pragma acc cache (a[idx:len])
+
+        b[ii] = a[ii];
+    }
+}
+
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != b[i])
+            abort ();
+    }
+
+    return 0;
+}
