diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 938de5859513f..b2ea52258fb46 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12217,8 +12217,8 @@ def err_acc_construct_appertainment
     : Error<"OpenACC construct '%0' cannot be used here; it can only "
             "be used in a statement context">;
 def err_acc_branch_in_out_compute_construct
-    : Error<"invalid %select{branch|return}0 %select{out of|into}1 OpenACC "
-            "Compute Construct">;
+    : Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 "
+            "OpenACC Compute Construct">;
 def note_acc_branch_into_compute_construct
     : Note<"invalid branch into OpenACC Compute Construct">;
 def note_acc_branch_out_of_compute_construct
diff --git a/clang/include/clang/Sema/Scope.h b/clang/include/clang/Sema/Scope.h
index 1cb2fa83e0bb3..536c12cb9d64e 100644
--- a/clang/include/clang/Sema/Scope.h
+++ b/clang/include/clang/Sema/Scope.h
@@ -43,6 +43,9 @@ class Scope {
   /// ScopeFlags - These are bitfields that are or'd together when creating a
   /// scope, which defines the sorts of things the scope contains.
   enum ScopeFlags {
+    // A bitfield value representing no scopes.
+    NoScope = 0,
+
     /// This indicates that the scope corresponds to a function, which
     /// means that labels are set here.
     FnScope = 0x01,
@@ -521,10 +524,17 @@ class Scope {
     return getFlags() & Scope::OpenACCComputeConstructScope;
   }
 
-  bool isInOpenACCComputeConstructScope() const {
+  /// Determine if this scope (or its parents) are a compute construct. If the
+  /// argument is provided, the search will stop at any of the specified scopes.
+  /// Otherwise, it will stop only at the normal 'no longer search' scopes.
+  bool isInOpenACCComputeConstructScope(ScopeFlags Flags = NoScope) const {
     for (const Scope *S = this; S; S = S->getParent()) {
-      if (S->getFlags() & Scope::OpenACCComputeConstructScope)
+      if (S->isOpenACCComputeConstructScope())
         return true;
+
+      if (S->getFlags() & Flags)
+        return false;
+
       else if (S->getFlags() &
                (Scope::FnScope | Scope::ClassScope | Scope::BlockScope |
                 Scope::TemplateParamScope | Scope::FunctionPrototypeScope |
@@ -534,25 +544,6 @@ class Scope {
     return false;
   }
 
-  /// Determine if this scope (or its parents) are a compute construct inside of
-  /// the nearest 'switch' scope.  This is needed to check whether we are inside
-  /// of a 'duffs' device, which is an illegal branch into a compute construct.
-  bool isInOpenACCComputeConstructBeforeSwitch() const {
-    for (const Scope *S = this; S; S = S->getParent()) {
-      if (S->getFlags() & Scope::OpenACCComputeConstructScope)
-        return true;
-      if (S->getFlags() & Scope::SwitchScope)
-        return false;
-
-      if (S->getFlags() &
-          (Scope::FnScope | Scope::ClassScope | Scope::BlockScope |
-           Scope::TemplateParamScope | Scope::FunctionPrototypeScope |
-           Scope::AtCatchScope | Scope::ObjCMethodScope))
-        return false;
-    }
-    return false;
-  }
-
   /// Determine whether this scope is a while/do/for statement, which can have
   /// continue statements embedded into it.
   bool isContinueScope() const {
diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp
index c4750ce78fa9c..c34a40fa7c81a 100644
--- a/clang/lib/Sema/SemaExprCXX.cpp
+++ b/clang/lib/Sema/SemaExprCXX.cpp
@@ -890,6 +890,12 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex,
   if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
     Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw";
 
+  // Exceptions that escape a compute construct are ill-formed.
+  if (getLangOpts().OpenACC && getCurScope() &&
+      getCurScope()->isInOpenACCComputeConstructScope(Scope::TryScope))
+    Diag(OpLoc, diag::err_acc_branch_in_out_compute_construct)
+        << /*throw*/ 2 << /*out of*/ 0;
+
   if (Ex && !Ex->isTypeDependent()) {
     // Initialize the exception result.  This implicitly weeds out
     // abstract types or types with inaccessible copy constructors.
diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp
index 4a15a8f6effd3..e72397adec24f 100644
--- a/clang/lib/Sema/SemaStmt.cpp
+++ b/clang/lib/Sema/SemaStmt.cpp
@@ -528,7 +528,7 @@ Sema::ActOnCaseStmt(SourceLocation CaseLoc, ExprResult LHSVal,
   }
 
   if (LangOpts.OpenACC &&
-      getCurScope()->isInOpenACCComputeConstructBeforeSwitch()) {
+      getCurScope()->isInOpenACCComputeConstructScope(Scope::SwitchScope)) {
     Diag(CaseLoc, diag::err_acc_branch_in_out_compute_construct)
         << /*branch*/ 0 << /*into*/ 1;
     return StmtError();
@@ -554,7 +554,7 @@ Sema::ActOnDefaultStmt(SourceLocation DefaultLoc, SourceLocation ColonLoc,
   }
 
   if (LangOpts.OpenACC &&
-      getCurScope()->isInOpenACCComputeConstructBeforeSwitch()) {
+      getCurScope()->isInOpenACCComputeConstructScope(Scope::SwitchScope)) {
     Diag(DefaultLoc, diag::err_acc_branch_in_out_compute_construct)
         << /*branch*/ 0 << /*into*/ 1;
     return StmtError();
diff --git a/clang/test/SemaOpenACC/no-branch-in-out.cpp b/clang/test/SemaOpenACC/no-branch-in-out.cpp
index e7d5683f9bc78..6ee4553cd3035 100644
--- a/clang/test/SemaOpenACC/no-branch-in-out.cpp
+++ b/clang/test/SemaOpenACC/no-branch-in-out.cpp
@@ -136,6 +136,66 @@ void DuffsDevice() {
   }
 }
 
+void Exceptions() {
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    throw 5; // expected-error{{invalid throw out of OpenACC Compute Construct}}
+  }
+
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    throw; // expected-error{{invalid throw out of OpenACC Compute Construct}}
+  }
+
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    try {
+    throw 5;
+    } catch(float f) {
+    }
+  }
+
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    try {
+    throw 5;
+    } catch(int f) {
+    }
+  }
+
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    try {
+    throw 5;
+    } catch(...) {
+    }
+  }
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    try {
+    throw;
+    } catch(...) {
+    }
+  }
+
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    try {
+    throw;
+    } catch(...) {
+      throw; // expected-error{{invalid throw out of OpenACC Compute Construct}}
+    }
+  }
+#pragma acc parallel
+  for(int i = 0; i < 5; ++i) {
+    try {
+    throw;
+    } catch(int f) {
+      throw; // expected-error{{invalid throw out of OpenACC Compute Construct}}
+    }
+  }
+}
+
 void Instantiate() {
   BreakContinue<int>();
   DuffsDevice<int>();