[OpenACC] implement 'detach' clause sema

This is another new clause specific to 'exit data' that takes a pointer
argument. This patch implements this the same way we do a few other
clauses (like attach) that have the same restrictions.

clang/include/clang/AST/OpenACCClause.h

@@ -744,6 +744,28 @@ public:
          ArrayRef<Expr *> VarList, SourceLocation EndLoc);
 };
 
+class OpenACCDetachClause final
+    : public OpenACCClauseWithVarList,
+      public llvm::TrailingObjects<OpenACCDetachClause, Expr *> {
+
+  OpenACCDetachClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                      ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+      : OpenACCClauseWithVarList(OpenACCClauseKind::Detach, BeginLoc, LParenLoc,
+                                 EndLoc) {
+    std::uninitialized_copy(VarList.begin(), VarList.end(),
+                            getTrailingObjects<Expr *>());
+    setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+  }
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::Detach;
+  }
+  static OpenACCDetachClause *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+         ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
 class OpenACCNoCreateClause final
     : public OpenACCClauseWithVarList,
       public llvm::TrailingObjects<OpenACCNoCreateClause, Expr *> {

clang/include/clang/Basic/OpenACCClauses.def

@@ -38,6 +38,7 @@ VISIT_CLAUSE(Create)
 CLAUSE_ALIAS(PCreate, Create, true)
 CLAUSE_ALIAS(PresentOrCreate, Create, true)
 VISIT_CLAUSE(Default)
+VISIT_CLAUSE(Detach)
 VISIT_CLAUSE(DevicePtr)
 VISIT_CLAUSE(DeviceType)
 CLAUSE_ALIAS(DType, DeviceType, false)

clang/include/clang/Sema/SemaOpenACC.h

@@ -399,6 +399,7 @@ public:
               ClauseKind == OpenACCClauseKind::PCreate ||
               ClauseKind == OpenACCClauseKind::PresentOrCreate ||
               ClauseKind == OpenACCClauseKind::Attach ||
+              ClauseKind == OpenACCClauseKind::Detach ||
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::Reduction ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -535,6 +536,7 @@ public:
               ClauseKind == OpenACCClauseKind::PCreate ||
               ClauseKind == OpenACCClauseKind::PresentOrCreate ||
               ClauseKind == OpenACCClauseKind::Attach ||
+              ClauseKind == OpenACCClauseKind::Detach ||
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
              "Parsed clause kind does not have a var-list");
@@ -571,6 +573,7 @@ public:
               ClauseKind == OpenACCClauseKind::PCreate ||
               ClauseKind == OpenACCClauseKind::PresentOrCreate ||
               ClauseKind == OpenACCClauseKind::Attach ||
+              ClauseKind == OpenACCClauseKind::Detach ||
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
              "Parsed clause kind does not have a var-list");

clang/lib/AST/OpenACCClause.cpp

@@ -33,7 +33,8 @@ bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
          OpenACCFirstPrivateClause::classof(C) ||
          OpenACCDevicePtrClause::classof(C) ||
          OpenACCDevicePtrClause::classof(C) ||
-         OpenACCAttachClause::classof(C) || OpenACCNoCreateClause::classof(C) ||
+         OpenACCDetachClause::classof(C) || OpenACCAttachClause::classof(C) ||
+         OpenACCNoCreateClause::classof(C) ||
          OpenACCPresentClause::classof(C) || OpenACCCopyClause::classof(C) ||
          OpenACCCopyInClause::classof(C) || OpenACCCopyOutClause::classof(C) ||
          OpenACCReductionClause::classof(C) || OpenACCCreateClause::classof(C);
@@ -277,6 +278,16 @@ OpenACCAttachClause *OpenACCAttachClause::Create(const ASTContext &C,
   return new (Mem) OpenACCAttachClause(BeginLoc, LParenLoc, VarList, EndLoc);
 }
 
+OpenACCDetachClause *OpenACCDetachClause::Create(const ASTContext &C,
+                                                 SourceLocation BeginLoc,
+                                                 SourceLocation LParenLoc,
+                                                 ArrayRef<Expr *> VarList,
+                                                 SourceLocation EndLoc) {
+  void *Mem =
+      C.Allocate(OpenACCDetachClause::totalSizeToAlloc<Expr *>(VarList.size()));
+  return new (Mem) OpenACCDetachClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
 OpenACCDevicePtrClause *OpenACCDevicePtrClause::Create(const ASTContext &C,
                                                        SourceLocation BeginLoc,
                                                        SourceLocation LParenLoc,
@@ -546,6 +557,13 @@ void OpenACCClausePrinter::VisitAttachClause(const OpenACCAttachClause &C) {
   OS << ")";
 }
 
+void OpenACCClausePrinter::VisitDetachClause(const OpenACCDetachClause &C) {
+  OS << "detach(";
+  llvm::interleaveComma(C.getVarList(), OS,
+                        [&](const Expr *E) { printExpr(E); });
+  OS << ")";
+}
+
 void OpenACCClausePrinter::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &C) {
   OS << "deviceptr(";

clang/lib/AST/StmtProfile.cpp

@@ -2605,6 +2605,12 @@ void OpenACCClauseProfiler::VisitAttachClause(
     Profiler.VisitStmt(E);
 }
 
+void OpenACCClauseProfiler::VisitDetachClause(
+    const OpenACCDetachClause &Clause) {
+  for (auto *E : Clause.getVarList())
+    Profiler.VisitStmt(E);
+}
+
 void OpenACCClauseProfiler::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &Clause) {
   for (auto *E : Clause.getVarList())

clang/lib/AST/TextNodeDumper.cpp

@@ -411,6 +411,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::If:
     case OpenACCClauseKind::IfPresent:
     case OpenACCClauseKind::Independent:
+    case OpenACCClauseKind::Detach:
     case OpenACCClauseKind::DevicePtr:
     case OpenACCClauseKind::Finalize:
     case OpenACCClauseKind::FirstPrivate:

clang/lib/Parse/ParseOpenACC.cpp

@@ -999,7 +999,6 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       assert(DirKind == OpenACCDirectiveKind::Update);
       [[fallthrough]];
     case OpenACCClauseKind::Delete:
-    case OpenACCClauseKind::Detach:
     case OpenACCClauseKind::Device:
     case OpenACCClauseKind::DeviceResident:
     case OpenACCClauseKind::Host:
@@ -1008,6 +1007,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       ParseOpenACCVarList(ClauseKind);
       break;
     case OpenACCClauseKind::Attach:
+    case OpenACCClauseKind::Detach:
     case OpenACCClauseKind::DevicePtr:
       ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
                                      /*IsReadOnly=*/false, /*IsZero=*/false);

clang/lib/Sema/SemaOpenACC.cpp

@@ -425,6 +425,14 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::Detach: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::ExitData:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -1043,6 +1051,21 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
                                      Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitDetachClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  // ActOnVar ensured that everything is a valid variable reference, but we
+  // still have to make sure it is a pointer type.
+  llvm::SmallVector<Expr *> VarList{Clause.getVarList()};
+  llvm::erase_if(VarList, [&](Expr *E) {
+    return SemaRef.CheckVarIsPointerType(OpenACCClauseKind::Detach, E);
+  });
+  Clause.setVarListDetails(VarList,
+                           /*IsReadOnly=*/false, /*IsZero=*/false);
+  return OpenACCDetachClause::Create(Ctx, Clause.getBeginLoc(),
+                                     Clause.getLParenLoc(), Clause.getVarList(),
+                                     Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // Restrictions only properly implemented on 'compute'/'combined'/'data'

clang/lib/Sema/TreeTransform.h

@@ -11755,6 +11755,28 @@ void OpenACCClauseTransform<Derived>::VisitAttachClause(
       ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDetachClause(
+    const OpenACCDetachClause &C) {
+  llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList());
+
+  // Ensure each var is a pointer type.
+  VarList.erase(
+      std::remove_if(VarList.begin(), VarList.end(),
+                     [&](Expr *E) {
+                       return Self.getSema().OpenACC().CheckVarIsPointerType(
+                           OpenACCClauseKind::Detach, E);
+                     }),
+      VarList.end());
+
+  ParsedClause.setVarListDetails(VarList,
+                                 /*IsReadOnly=*/false, /*IsZero=*/false);
+  NewClause = OpenACCDetachClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &C) {

clang/lib/Serialization/ASTReader.cpp

@@ -12442,6 +12442,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCAttachClause::Create(getContext(), BeginLoc, LParenLoc,
                                        VarList, EndLoc);
   }
+  case OpenACCClauseKind::Detach: {
+    SourceLocation LParenLoc = readSourceLocation();
+    llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+    return OpenACCDetachClause::Create(getContext(), BeginLoc, LParenLoc,
+                                       VarList, EndLoc);
+  }
   case OpenACCClauseKind::DevicePtr: {
     SourceLocation LParenLoc = readSourceLocation();
     llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
@@ -12586,7 +12592,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Delete:
-  case OpenACCClauseKind::Detach:
   case OpenACCClauseKind::Device:
   case OpenACCClauseKind::DeviceResident:
   case OpenACCClauseKind::Host:

clang/lib/Serialization/ASTWriter.cpp

@@ -8356,6 +8356,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     writeOpenACCVarList(AC);
     return;
   }
+  case OpenACCClauseKind::Detach: {
+    const auto *DC = cast<OpenACCDetachClause>(C);
+    writeSourceLocation(DC->getLParenLoc());
+    writeOpenACCVarList(DC);
+    return;
+  }
   case OpenACCClauseKind::DevicePtr: {
     const auto *DPC = cast<OpenACCDevicePtrClause>(C);
     writeSourceLocation(DPC->getLParenLoc());
@@ -8501,7 +8507,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Delete:
-  case OpenACCClauseKind::Detach:
   case OpenACCClauseKind::Device:
   case OpenACCClauseKind::DeviceResident:
   case OpenACCClauseKind::Host:

clang/test/AST/ast-print-openacc-data-construct.cpp

@@ -117,4 +117,7 @@ void foo() {
 // CHECK: #pragma acc host_data if_present
 #pragma acc host_data use_device(i) if_present
   ;
+// CHECK: #pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
+#pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
+
 }

clang/test/ParserOpenACC/parse-clauses.c

@@ -488,14 +488,13 @@ void VarListClauses() {
 #pragma acc serial attach(IsPointer), self
   for(int i = 0; i < 5;++i) {}
 
-  // expected-error@+2{{expected ','}}
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented, clause ignored}}
-#pragma acc serial detach(s.array[s.value] s.array[s.value :5] ), self
-  for(int i = 0; i < 5;++i) {}
+  // expected-error@+4{{expected ','}}
+  // expected-error@+3{{expected pointer in 'detach' clause, type is 'char'}}
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc exit data copyout(s) detach(s.array[s.value] s.array[s.value :5])
 
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented, clause ignored}}
-#pragma acc serial detach(s.array[s.value : 5], s.value), self
-  for(int i = 0; i < 5;++i) {}
+#pragma acc exit data copyout(s) detach(IsPointer)
 
   // expected-error@+1{{expected ','}}
 #pragma acc serial private(s.array[s.value] s.array[s.value :5] ), self

clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c

@@ -72,7 +72,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc parallel loop auto delete(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto detach(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -189,7 +189,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc parallel loop delete(Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop detach(Var) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -307,7 +307,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc parallel loop independent delete(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent detach(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -424,7 +424,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc parallel loop delete(Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop detach(Var) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -550,7 +550,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc parallel loop seq delete(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq detach(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -673,7 +673,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc parallel loop delete(Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop detach(Var) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}

clang/test/SemaOpenACC/combined-construct-device_type-clause.c

@@ -99,8 +99,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc serial loop device_type(*) delete(Var)
   for(int i = 0; i < 5; ++i);
-  // expected-error@+2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'kernels loop' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'kernels loop' directive}}
 #pragma acc kernels loop device_type(*) detach(Var)
   for(int i = 0; i < 5; ++i);
   // expected-error@+2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'parallel loop' construct}}

clang/test/SemaOpenACC/compute-construct-device_type-clause.c

@@ -103,8 +103,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc kernels device_type(*) delete(Var)
   while(1);
-  // expected-error@+2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'kernels' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) detach(Var)
   while(1);
   // expected-error@+2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'kernels' construct}}

clang/test/SemaOpenACC/data-construct-detach-ast.cpp

@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+int Global;
+short GlobalArray[5];
+
+void NormalUses(float *PointerParam) {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK: ParmVarDecl
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc exit data copyout(Global) detach(PointerParam)
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+  // CHECK-NEXT: copyout clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
+  // CHECK-NEXT: detach clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+}
+
+template<typename T>
+void TemplUses(T *t) {
+  // CHECK-NEXT: FunctionTemplateDecl
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T *)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T *'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc exit data copyout(Global) detach(t)
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+  // CHECK-NEXT: copyout clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
+  // CHECK-NEXT: detach clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 't' 'T *'
+
+
+  // Check the instantiated versions of the above.
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (int *)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used t 'int *'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+  // CHECK-NEXT: copyout clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
+  // CHECK-NEXT: detach clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 't' 'int *'
+
+}
+
+void Inst() {
+  int i;
+  TemplUses(&i);
+}
+#endif

clang/test/SemaOpenACC/data-construct-detach-clause.c

@@ -0,0 +1,68 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+  int IntMem;
+  int *PtrMem;
+};
+
+void uses() {
+  int LocalInt;
+  int *LocalPtr;
+  int Array[5];
+  int *PtrArray[5];
+  struct S s;
+
+  // expected-error@+1{{expected pointer in 'detach' clause, type is 'int'}}
+#pragma acc exit data copyout(LocalInt) detach(LocalInt)
+  ;
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc exit data copyout(LocalInt) detach(&LocalInt)
+  ;
+
+
+  // expected-error@+1{{expected pointer in 'detach' clause, type is 'int[5]'}}
+#pragma acc exit data copyout(LocalInt) detach(Array)
+
+  // expected-error@+1{{expected pointer in 'detach' clause, type is 'int'}}
+#pragma acc exit data copyout(LocalInt) detach(Array[0])
+  ;
+
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc exit data copyout(LocalInt) detach(Array[0:1])
+  ;
+
+  // expected-error@+1{{expected pointer in 'detach' clause, type is 'int *[5]'}}
+#pragma acc exit data copyout(LocalInt) detach(PtrArray)
+  ;
+
+#pragma acc exit data copyout(LocalInt) detach(PtrArray[0])
+  ;
+
+  // expected-error@+2{{OpenACC sub-array is not allowed here}}
+  // expected-note@+1{{expected variable of pointer type}}
+#pragma acc exit data copyout(LocalInt) detach(PtrArray[0:1])
+  ;
+
+  // expected-error@+1{{expected pointer in 'detach' clause, type is 'struct S'}}
+#pragma acc exit data copyout(LocalInt) detach(s)
+  ;
+
+  // expected-error@+1{{expected pointer in 'detach' clause, type is 'int'}}
+#pragma acc exit data copyout(LocalInt) detach(s.IntMem)
+  ;
+
+#pragma acc exit data copyout(LocalInt) detach(s.PtrMem)
+  ;
+
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'data' directive}}
+#pragma acc data copyin(LocalInt) detach(PtrArray[0])
+  ;
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'enter data' directive}}
+#pragma acc enter data copyin(LocalInt) detach(PtrArray[0])
+  // expected-warning@+2{{OpenACC clause 'use_device' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'host_data' directive}}
+#pragma acc host_data use_device(LocalInt) detach(PtrArray[0])
+  ;
+}

clang/test/SemaOpenACC/data-construct.cpp

@@ -75,8 +75,7 @@ void AtLeastOneOf() {
 #pragma acc exit data copyout(Var)
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc exit data delete(Var)
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
-#pragma acc exit data detach(Var)
+#pragma acc exit data detach(VarPtr)
 
   // OpenACC TODO: The following 'exit data' directives should diagnose, since
   // they don't have at least one of the above clauses.

clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c

@@ -77,7 +77,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc loop auto delete(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop auto detach(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -211,7 +211,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc loop delete(Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop detach(Var) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -346,7 +346,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc loop independent delete(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop independent detach(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -480,7 +480,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc loop delete(Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop detach(Var) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -623,7 +623,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc loop seq delete(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop seq detach(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}
@@ -763,7 +763,7 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
 #pragma acc loop delete(Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'detach' not yet implemented}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop detach(Var) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device' not yet implemented}}

clang/test/SemaOpenACC/loop-construct-device_type-clause.c

@@ -91,8 +91,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc loop device_type(*) delete(Var)
   for(int i = 0; i < 5; ++i);
-  // expected-error@+2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'loop' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
 #pragma acc loop device_type(*) detach(Var)
   for(int i = 0; i < 5; ++i);
   // expected-error@+2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'loop' construct}}

clang/tools/libclang/CIndex.cpp

@@ -2886,6 +2886,11 @@ void OpenACCClauseEnqueue::VisitCreateClause(const OpenACCCreateClause &C) {
 void OpenACCClauseEnqueue::VisitAttachClause(const OpenACCAttachClause &C) {
   VisitVarList(C);
 }
+
+void OpenACCClauseEnqueue::VisitDetachClause(const OpenACCDetachClause &C) {
+  VisitVarList(C);
+}
+
 void OpenACCClauseEnqueue::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &C) {
   VisitVarList(C);