[OpenACC] Implement no throw out of Compute construct

Exception handling SHOULD be possible depending on codegen (and if not,
we can make it trap and add a warning when we make that decision), but
throwing out of a compute construct is ill formed.

This patch adds an error for a 'throw' that isn't in a 'try' scope. This
error is imperfect as it won't diagnose a 'throw' that escapes its
'try', or one in a separate function, but it catches the obvious
mistakes. The other cases will need to be handled as runtime failures.

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

clang/include/clang/Sema/Scope.h

@@ -43,6 +43,9 @@ public:
   /// 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 @@ public:
     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 @@ public:
     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 {

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.

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();

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>();