[OpenACC] Start enforcing 'device_type' clause values

Researching in prep of doing the implementation for lowering, I found that the source of the valid identifiers list from flang is in the frontend. This patch adds the same list to the frontend, but does it as a sema diagnostic, so we still parse it as an identifier/identifier-like thing, but then diagnose it as invalid later.
2025-04-15 22:16:30 +00:00 · 2025-04-09 14:20:15 -07:00 · 2025-04-09 14:20:15 -07:00 · d47401e376
commit d47401e376
parent af9c04fa68
34 changed files with 388 additions and 281 deletions
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@ -13061,6 +13061,8 @@ def err_acc_decl_for_routine
    : Error<"expected function or lambda declaration for 'routine' construct">;
 def err_acc_invalid_modifier
    : Error<"OpenACC '%0' modifier not valid on '%1' clause">;
+def err_acc_invalid_default_type
+    : Error<"invalid value %0 in '%1' clause; valid values are %2">;

 // AMDGCN builtins diagnostics
 def err_amdgcn_load_lds_size_invalid_value : Error<"invalid size value">;
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@ -1338,13 +1338,50 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause(
      checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
    return nullptr;

-  // TODO OpenACC: Once we get enough of the CodeGen implemented that we have
-  // a source for the list of valid architectures, we need to warn on unknown
-  // identifiers here.
+  // The list of valid device_type values. Flang also has these hardcoded in
+  // openacc_parsers.cpp, as there does not seem to be a reliable backend
+  // source. The list below is sourced from Flang, though NVC++ supports only
+  // 'nvidia', 'host', 'multicore', and 'default'.
+  const std::array<llvm::StringLiteral, 6> ValidValues{
+      "default", "nvidia", "acc_device_nvidia", "radeon", "host", "multicore"};
+  // As an optimization, we have a manually maintained list of valid values
+  // below, rather than trying to calculate from above. These should be kept in
+  // sync if/when the above list ever changes.
+  std::string ValidValuesString =
+      "'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'";
+
+  llvm::SmallVector<DeviceTypeArgument> Architectures{
+      Clause.getDeviceTypeArchitectures()};
+
+  // The parser has ensured that we either have a single entry of just '*'
+  // (represented by a nullptr IdentifierInfo), or a list.
+
+  bool Diagnosed = false;
+  auto FilterPred = [&](const DeviceTypeArgument &Arch) {
+    // The '*' case.
+    if (!Arch.first)
+      return false;
+    return llvm::find_if(ValidValues, [&](StringRef RHS) {
+             return Arch.first->getName().equals_insensitive(RHS);
+           }) == ValidValues.end();
+  };
+
+  auto Diagnose = [&](const DeviceTypeArgument &Arch) {
+    Diagnosed = SemaRef.Diag(Arch.second, diag::err_acc_invalid_default_type)
+                << Arch.first << Clause.getClauseKind() << ValidValuesString;
+  };
+
+  // There aren't stable enumertor versions of 'for-each-then-erase', so do it
+  // here.  We DO keep track of whether we diagnosed something to make sure we
+  // don't do the 'erase_if' in the event that the first list didn't find
+  // anything.
+  llvm::for_each(llvm::make_filter_range(Architectures, FilterPred), Diagnose);
+  if (Diagnosed)
+    llvm::erase_if(Architectures, FilterPred);

  return OpenACCDeviceTypeClause::Create(
      Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
-      Clause.getDeviceTypeArchitectures(), Clause.getEndLoc());
+      Architectures, Clause.getEndLoc());
 }

 OpenACCClause *SemaOpenACCClauseVisitor::VisitAutoClause(
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@ -37,40 +37,40 @@ void foo() {
  bool SomeB;
  struct SomeStruct{} SomeStructImpl;

-//CHECK: #pragma acc parallel loop dtype(SomeB)
+//CHECK: #pragma acc parallel loop dtype(default)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc parallel loop dtype(SomeB)
+#pragma acc parallel loop dtype(default)
  for(int i = 0;i<5;++i);

-//CHECK: #pragma acc serial loop device_type(SomeStruct)
+//CHECK: #pragma acc serial loop device_type(radeon, host)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc serial loop device_type(SomeStruct)
+#pragma acc serial loop device_type(radeon, host)
  for(int i = 0;i<5;++i);

-//CHECK: #pragma acc kernels loop device_type(int)
+//CHECK: #pragma acc kernels loop device_type(nvidia)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc kernels loop device_type(int)
+#pragma acc kernels loop device_type(nvidia)
  for(int i = 0;i<5;++i);

-//CHECK: #pragma acc parallel loop dtype(bool)
+//CHECK: #pragma acc parallel loop dtype(multicore)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc parallel loop dtype(bool)
+#pragma acc parallel loop dtype(multicore)
  for(int i = 0;i<5;++i);

-//CHECK: #pragma acc serial loop device_type(SomeStructImpl)
+//CHECK: #pragma acc serial loop device_type(default)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc serial loop device_type (SomeStructImpl)
+#pragma acc serial loop device_type (default)
  for(int i = 0;i<5;++i);

-// CHECK: #pragma acc kernels loop dtype(AnotherIdent)
+// CHECK: #pragma acc kernels loop dtype(host)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc kernels loop dtype(AnotherIdent)
+#pragma acc kernels loop dtype(host)
  for(int i = 0;i<5;++i);

  int i;
--- a/clang/test/AST/ast-print-openacc-compute-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-compute-construct.cpp
@ -115,24 +115,24 @@ void foo() {
  bool SomeB;
  struct SomeStruct{} SomeStructImpl;

-//CHECK: #pragma acc parallel dtype(SomeB)
-#pragma acc parallel dtype(SomeB)
+//CHECK: #pragma acc parallel dtype(default)
+#pragma acc parallel dtype(default)
  while(true);

-//CHECK: #pragma acc parallel device_type(SomeStruct)
-#pragma acc parallel device_type(SomeStruct)
+//CHECK: #pragma acc parallel device_type(radeon)
+#pragma acc parallel device_type(radeon)
  while(true);

-//CHECK: #pragma acc parallel device_type(int)
-#pragma acc parallel device_type(int)
+//CHECK: #pragma acc parallel device_type(nvidia)
+#pragma acc parallel device_type(nvidia)
  while(true);

-//CHECK: #pragma acc parallel dtype(bool)
-#pragma acc parallel dtype(bool)
+//CHECK: #pragma acc parallel dtype(multicore)
+#pragma acc parallel dtype(multicore)
  while(true);

-//CHECK: #pragma acc parallel device_type(SomeStructImpl)
-#pragma acc parallel device_type (SomeStructImpl)
+//CHECK: #pragma acc parallel device_type(host)
+#pragma acc parallel device_type (host)
  while(true);

 //CHECK: #pragma acc parallel reduction(+: iPtr)
--- a/clang/test/AST/ast-print-openacc-data-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-data-construct.cpp
@ -2,16 +2,12 @@

 void foo() {
  int Var;
-  // TODO OpenACC: These are only legal if they have one of a list of clauses on
-  // them, so the 'check' lines should start to include those once we implement
-  // them.  For now, they don't emit those because they are 'not implemented'.
-
 // CHECK: #pragma acc data default(none)
 #pragma acc data default(none)
  ;

-// CHECK: #pragma acc data default(none) device_type(int)
-#pragma acc data default(none) device_type(int)
+// CHECK: #pragma acc data default(none) device_type(radeon)
+#pragma acc data default(none) device_type(radeon)
  ;

 // CHECK: #pragma acc enter data copyin(Var)
--- a/clang/test/AST/ast-print-openacc-init-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-init-construct.cpp
@ -9,6 +9,6 @@ void uses() {
 #pragma acc init device_type(*) device_num(Int)
 // CHECK: #pragma acc init device_type(*) if(Int == 5)
 #pragma acc init device_type(*) if (Int == 5)
-// CHECK: #pragma acc init device_type(SomeName)
-#pragma acc init device_type(SomeName)
+// CHECK: #pragma acc init device_type(radeon)
+#pragma acc init device_type(radeon)
 }
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@ -10,28 +10,28 @@ void foo() {
 #pragma acc loop
  for(int i = 0;i<5;++i);

-// CHECK: #pragma acc loop device_type(SomeStruct)
+// CHECK: #pragma acc loop device_type(default)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc loop device_type(SomeStruct)
+#pragma acc loop device_type(default)
  for(int i = 0;i<5;++i);

-// CHECK: #pragma acc loop device_type(int)
+// CHECK: #pragma acc loop device_type(nvidia)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc loop device_type(int)
+#pragma acc loop device_type(nvidia)
  for(int i = 0;i<5;++i);

-// CHECK: #pragma acc loop dtype(bool)
+// CHECK: #pragma acc loop dtype(radeon)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc loop dtype(bool)
+#pragma acc loop dtype(radeon)
  for(int i = 0;i<5;++i);

-// CHECK: #pragma acc loop dtype(AnotherIdent)
+// CHECK: #pragma acc loop dtype(host)
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
-#pragma acc loop dtype(AnotherIdent)
+#pragma acc loop dtype(host)
  for(int i = 0;i<5;++i);

 // CHECK: #pragma acc loop independent
--- a/clang/test/AST/ast-print-openacc-routine-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp
@ -8,10 +8,10 @@ auto Lambda = [](){};
 int function();
 #pragma acc routine (function) vector nohost bind("string")
 // CHECK: #pragma acc routine(function) vector nohost bind("string")
-#pragma acc routine(function) device_type(Something) seq
-// CHECK-NEXT: #pragma acc routine(function) device_type(Something) seq
-#pragma acc routine(function) dtype(Something) seq
-// CHECK-NEXT: #pragma acc routine(function) dtype(Something) seq
+#pragma acc routine(function) device_type(multicore) seq
+// CHECK-NEXT: #pragma acc routine(function) device_type(multicore) seq
+#pragma acc routine(function) dtype(radeon) seq
+// CHECK-NEXT: #pragma acc routine(function) dtype(radeon) seq

 #pragma acc routine nohost vector
 int function2();
@ -136,8 +136,8 @@ struct DepS {

 #pragma acc routine (MemFunc) worker dtype(*)
 // CHECK-NEXT: #pragma acc routine(MemFunc) worker dtype(*)
-#pragma acc routine (MemFunc) device_type(Lambda) vector
-// CHECK-NEXT: #pragma acc routine(MemFunc) device_type(Lambda) vector
+#pragma acc routine (MemFunc) device_type(nvidia) vector
+// CHECK-NEXT: #pragma acc routine(MemFunc) device_type(nvidia) vector
 };

 // CHECK: #pragma acc routine(DepS<int>::Lambda) gang bind("string")
--- a/clang/test/AST/ast-print-openacc-set-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-set-construct.cpp
@ -3,22 +3,22 @@
 unsigned Int;

 void uses() {
-// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_type(I) device_num(Int)
-#pragma acc set default_async(Int) if (Int == 5) device_type(I) device_num(Int)
-// CHECK: #pragma acc set default_async(Int) device_type(I) device_num(Int)
-#pragma acc set default_async(Int) device_type(I) device_num(Int)
+// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_type(default) device_num(Int)
+#pragma acc set default_async(Int) if (Int == 5) device_type(default) device_num(Int)
+// CHECK: #pragma acc set default_async(Int) device_type(nvidia) device_num(Int)
+#pragma acc set default_async(Int) device_type(nvidia) device_num(Int)
 // CHECK: #pragma acc set default_async(Int) if(Int == 5) device_num(Int)
 #pragma acc set default_async(Int) if (Int == 5) device_num(Int)
-// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_type(I)
-#pragma acc set default_async(Int) if (Int == 5) device_type(I)
-// CHECK: #pragma acc set if(Int == 5) device_type(I) device_num(Int)
-#pragma acc set if (Int == 5) device_type(I) device_num(Int)
+// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_type(host)
+#pragma acc set default_async(Int) if (Int == 5) device_type(host)
+// CHECK: #pragma acc set if(Int == 5) device_type(multicore) device_num(Int)
+#pragma acc set if (Int == 5) device_type(multicore) device_num(Int)
 // CHECK: #pragma acc set default_async(Int)
 #pragma acc set default_async(Int)
 // CHECK: #pragma acc set if(Int == 5)
 #pragma acc set if (Int == 5)
-// CHECK: #pragma acc set device_type(I)
-#pragma acc set device_type(I)
+// CHECK: #pragma acc set device_type(radeon)
+#pragma acc set device_type(radeon)
 // CHECK: #pragma acc set device_num(Int)
 #pragma acc set device_num(Int)
 }
--- a/clang/test/AST/ast-print-openacc-shutdown-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-shutdown-construct.cpp
@ -9,6 +9,6 @@ void uses() {
 #pragma acc shutdown device_type(*) device_num(Int)
 // CHECK: #pragma acc shutdown device_type(*) if(Int == 5)
 #pragma acc shutdown device_type(*) if (Int == 5)
-// CHECK: #pragma acc shutdown device_type(SomeName)
-#pragma acc shutdown device_type(SomeName)
+// CHECK: #pragma acc shutdown device_type(radeon)
+#pragma acc shutdown device_type(radeon)
 }
--- a/clang/test/AST/ast-print-openacc-update-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-update-construct.cpp
@ -30,11 +30,11 @@ void uses(bool cond) {
 // CHECK: #pragma acc update self(I) wait(devnum: I : queues: *iPtr, I) if(I == array[I]) async(I)
 #pragma acc update self(I) wait(devnum:I:queues:*iPtr, I) if(I == array[I]) async(I)

-// CHECK: #pragma acc update self(I) device_type(I) dtype(H)
-#pragma acc update self(I) device_type(I) dtype(H)
+// CHECK: #pragma acc update self(I) device_type(nvidia) dtype(host)
+#pragma acc update self(I) device_type(nvidia) dtype(host)

-// CHECK: #pragma acc update self(I) device_type(J) dtype(K)
-#pragma acc update self(I) device_type(J) dtype(K)
+// CHECK: #pragma acc update self(I) device_type(radeon) dtype(multicore)
+#pragma acc update self(I) device_type(radeon) dtype(multicore)

 // CHECK: #pragma acc update self(I, iPtr, array, array[1], array[1:2])
 #pragma acc update self(I, iPtr, array, array[1], array[1:2])
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@ -1080,11 +1080,27 @@ void device_type() {
 #pragma acc parallel dtype(31, "bar")
  {}

+  // expected-error@+4{{invalid value 'ident' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+3{{invalid value 'auto' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+2{{invalid value 'int' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+1{{invalid value 'float' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
 #pragma acc parallel device_type(ident, auto, int, float)
  {}
+  // expected-error@+4{{invalid value 'ident' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+3{{invalid value 'auto' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+2{{invalid value 'int' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+1{{invalid value 'float' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
 #pragma acc parallel dtype(ident, auto, int, float)
  {}

+  // expected-error@+8{{invalid value 'ident' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+7{{invalid value 'auto' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+6{{invalid value 'int' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+5{{invalid value 'float' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+4{{invalid value 'ident' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+3{{invalid value 'auto' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+2{{invalid value 'int' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+1{{invalid value 'float' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
 #pragma acc parallel device_type(ident, auto, int, float) dtype(ident, auto, int, float)
  {}
 }
--- a/clang/test/SemaOpenACC/combined-construct-device_type-ast.cpp
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-ast.cpp
@ -19,11 +19,11 @@ void NormalUses() {
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} SomeVar 'bool'

-#pragma acc parallel loop device_type(SomeS) dtype(SomeImpl)
+#pragma acc parallel loop device_type(*) dtype(default)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
-  // CHECK-NEXT: device_type(SomeS)
-  // CHECK-NEXT: dtype(SomeImpl)
+  // CHECK-NEXT: device_type(*)
+  // CHECK-NEXT: dtype(default)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -36,11 +36,11 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc serial loop device_type(SomeVar) dtype(int)
+#pragma acc serial loop device_type(nvidia) dtype(acc_device_nvidia)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
-  // CHECK-NEXT: device_type(SomeVar)
-  // CHECK-NEXT: dtype(int)
+  // CHECK-NEXT: device_type(nvidia)
+  // CHECK-NEXT: dtype(acc_device_nvidia)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -53,11 +53,11 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc kernels loop device_type(private) dtype(struct)
+#pragma acc kernels loop device_type(radeon) dtype(host)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
-  // CHECK-NEXT: device_type(private)
-  // CHECK-NEXT: dtype(struct)
+  // CHECK-NEXT: device_type(radeon)
+  // CHECK-NEXT: dtype(host)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -70,27 +70,10 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel loop device_type(private) dtype(class)
+#pragma acc parallel loop device_type(multicore) dtype(*)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
-  // CHECK-NEXT: device_type(private)
-  // CHECK-NEXT: dtype(class)
-  // CHECK-NEXT: ForStmt
-  // CHECK-NEXT: DeclStmt
-  // CHECK-NEXT: VarDecl{{.*}} i 'int'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
-  // CHECK-NEXT: <<<NULL>>>
-  // CHECK-NEXT: BinaryOperator{{.*}}'<'
-  // CHECK-NEXT: ImplicitCastExpr
-  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
-  // CHECK-NEXT: UnaryOperator{{.*}}++
-  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
-  // CHECK-NEXT: CompoundStmt
-#pragma acc serial loop device_type(float) dtype(*)
-  for(int i = 0; i < 5; ++i){}
-  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
-  // CHECK-NEXT: device_type(float)
+  // CHECK-NEXT: device_type(multicore)
  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
@ -104,10 +87,27 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc kernels loop device_type(float, int) dtype(*)
+#pragma acc serial loop device_type(default, nvidia) dtype(*)
+  for(int i = 0; i < 5; ++i){}
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: device_type(default, nvidia)
+  // CHECK-NEXT: dtype(*)
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: CompoundStmt
+#pragma acc kernels loop device_type(acc_device_nvidia, radeon) dtype(*)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
-  // CHECK-NEXT: device_type(float, int)
+  // CHECK-NEXT: device_type(acc_device_nvidia, radeon)
  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
@ -129,11 +129,11 @@ void TemplUses() {
  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel loop device_type(T) dtype(T)
+#pragma acc parallel loop device_type(host, multicore) dtype(*)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(host, multicore)
+  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -155,8 +155,8 @@ void TemplUses() {
  // CHECK-NEXT: CompoundStmt

  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(host, multicore)
+  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@ -6,22 +6,22 @@ void uses() {
  typedef struct S{} STy;
  STy SImpl;

-#pragma acc parallel loop device_type(I)
+#pragma acc parallel loop device_type(radeon)
  for(int i = 0; i < 5; ++i);
-#pragma acc serial loop device_type(S) dtype(STy)
+#pragma acc serial loop device_type(multicore) dtype(host)
  for(int i = 0; i < 5; ++i);
-#pragma acc kernels loop dtype(SImpl)
+#pragma acc kernels loop dtype(acc_device_nvidia)
  for(int i = 0; i < 5; ++i);
-#pragma acc kernels loop dtype(int) device_type(*)
+#pragma acc kernels loop dtype(default) device_type(*)
  for(int i = 0; i < 5; ++i);
-#pragma acc kernels loop dtype(true) device_type(false)
+#pragma acc kernels loop dtype(nvidia) device_type(radeon)
  for(int i = 0; i < 5; ++i);

  // expected-error@+1{{expected identifier}}
-#pragma acc kernels loop dtype(int, *)
+#pragma acc kernels loop dtype(radeon, *)
  for(int i = 0; i < 5; ++i);

-#pragma acc parallel loop device_type(I, int)
+#pragma acc parallel loop device_type(host, multicore)
  for(int i = 0; i < 5; ++i);
  // expected-error@+2{{expected ','}}
  // expected-error@+1{{expected identifier}}
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.cpp
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.cpp
@ -2,17 +2,17 @@

 template<typename T>
 void TemplUses() {
-#pragma acc parallel loop device_type(I)
+#pragma acc parallel loop device_type(default)
  for(int i = 0; i < 5; ++i);
 #pragma acc serial loop dtype(*)
  for(int i = 0; i < 5; ++i);
-#pragma acc kernels loop device_type(class)
+#pragma acc kernels loop device_type(nvidia)
  for(int i = 0; i < 5; ++i);
-#pragma acc parallel loop device_type(private)
+#pragma acc parallel loop device_type(radeon)
  for(int i = 0; i < 5; ++i);
-#pragma acc serial loop device_type(bool)
+#pragma acc serial loop device_type(host)
  for(int i = 0; i < 5; ++i);
-#pragma acc kernels loop dtype(true) device_type(false)
+#pragma acc kernels loop dtype(multicore) device_type(host)
  for(int i = 0; i < 5; ++i);
  // expected-error@+2{{expected ','}}
  // expected-error@+1{{expected identifier}}
--- a/clang/test/SemaOpenACC/compute-construct-device_type-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-ast.cpp
@ -19,50 +19,50 @@ void NormalUses() {
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} SomeVar 'bool'

-#pragma acc parallel device_type(SomeS) dtype(SomeImpl)
+#pragma acc parallel device_type(default) dtype(nvidia)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(SomeS)
-  // CHECK-NEXT: dtype(SomeImpl)
+  // CHECK-NEXT: device_type(default)
+  // CHECK-NEXT: dtype(nvidia)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel device_type(SomeVar) dtype(int)
+#pragma acc parallel device_type(radeon) dtype(host)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(SomeVar)
-  // CHECK-NEXT: dtype(int)
+  // CHECK-NEXT: device_type(radeon)
+  // CHECK-NEXT: dtype(host)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel device_type(private) dtype(struct)
+#pragma acc parallel device_type(multicore) dtype(default)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(private)
-  // CHECK-NEXT: dtype(struct)
+  // CHECK-NEXT: device_type(multicore)
+  // CHECK-NEXT: dtype(default)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel device_type(private) dtype(class)
+#pragma acc parallel device_type(nvidia) dtype(acc_device_nvidia)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(private)
-  // CHECK-NEXT: dtype(class)
+  // CHECK-NEXT: device_type(nvidia)
+  // CHECK-NEXT: dtype(acc_device_nvidia)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel device_type(float) dtype(*)
+#pragma acc parallel device_type(radeon) dtype(*)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(float)
+  // CHECK-NEXT: device_type(radeon)
  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel device_type(float, int) dtype(*)
+#pragma acc parallel device_type(host, multicore) dtype(*)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(float, int)
+  // CHECK-NEXT: device_type(host, multicore)
  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
@ -75,11 +75,11 @@ void TemplUses() {
  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
  // CHECK-NEXT: CompoundStmt
-#pragma acc parallel device_type(T) dtype(T)
+#pragma acc parallel device_type(host) dtype(multicore)
  while(true){}
  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(host)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
@ -92,8 +92,8 @@ void TemplUses() {
  // CHECK-NEXT: CompoundStmt

  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(host)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: WhileStmt
  // CHECK-NEXT: CXXBoolLiteralExpr
  // CHECK-NEXT: CompoundStmt
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@ -6,22 +6,22 @@ void uses() {
  typedef struct S{} STy;
  STy SImpl;

-#pragma acc parallel device_type(I)
+#pragma acc parallel device_type(radeon)
  while(1);
-#pragma acc serial device_type(S) dtype(STy)
+#pragma acc serial device_type(nvidia) dtype(multicore)
  while(1);
-#pragma acc kernels dtype(SImpl)
+#pragma acc kernels dtype(host)
  while(1);
-#pragma acc kernels dtype(int) device_type(*)
+#pragma acc kernels dtype(default) device_type(*)
  while(1);
-#pragma acc kernels dtype(true) device_type(false)
+#pragma acc kernels dtype(nvidia) device_type(multicore)
  while(1);

  // expected-error@+1{{expected identifier}}
-#pragma acc kernels dtype(int, *)
+#pragma acc kernels dtype(default, *)
  while(1);

-#pragma acc parallel device_type(I, int)
+#pragma acc parallel device_type(host, radeon)
  while(1);
  // expected-error@+2{{expected ','}}
  // expected-error@+1{{expected identifier}}
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.cpp
@ -2,17 +2,17 @@

 template<typename T>
 void TemplUses() {
-#pragma acc parallel device_type(I)
+#pragma acc parallel device_type(default)
  while(true);
 #pragma acc parallel dtype(*)
  while(true);
-#pragma acc parallel device_type(class)
+#pragma acc parallel device_type(nvidia)
  while(true);
-#pragma acc parallel device_type(private)
+#pragma acc parallel device_type(radeon)
  while(true);
-#pragma acc parallel device_type(bool)
+#pragma acc parallel device_type(host)
  while(true);
-#pragma acc kernels dtype(true) device_type(false)
+#pragma acc kernels dtype(multicore) device_type(host)
  while(true);
  // expected-error@+2{{expected ','}}
  // expected-error@+1{{expected identifier}}
--- a/clang/test/SemaOpenACC/data-construct-device_type-ast.cpp
+++ b/clang/test/SemaOpenACC/data-construct-device_type-ast.cpp
@ -13,12 +13,12 @@ void TemplUses() {
  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
  // CHECK-NEXT: CompoundStmt

-#pragma acc data default(none) device_type(T) dtype(T)
+#pragma acc data default(none) device_type(host) dtype(multicore)
  ;
  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
  // CHECK-NEXT: default(none)
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(host)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: NullStmt

  // Instantiations
@ -30,8 +30,8 @@ void TemplUses() {
  // Argument to 'device-type' is just an identifier, so we don't transform it.
  // CHECK-NEXT: OpenACCDataConstruct{{.*}} data
  // CHECK-NEXT: default(none)
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(host)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: NullStmt
 }
 void Inst() {
--- a/clang/test/SemaOpenACC/data-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-device_type-clause.c
@ -2,54 +2,54 @@

 void uses() {
  int Var;
-#pragma acc data default(none) device_type(foo) async
+#pragma acc data default(none) device_type(radeon) async
  ;
-#pragma acc data default(none) device_type(foo) wait
+#pragma acc data default(none) device_type(radeon) wait
  ;
-#pragma acc data default(none) device_type(foo) dtype(false)
+#pragma acc data default(none) device_type(radeon) dtype(nvidia)
  ;
-#pragma acc data default(none) dtype(foo) device_type(false)
+#pragma acc data default(none) dtype(radeon) device_type(nvidia)
  ;

  // expected-error@+2{{OpenACC clause 'if' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) if(1)
+#pragma acc data default(none) device_type(radeon) if(1)
  ;
  // expected-error@+2{{OpenACC clause 'copy' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) copy(Var)
+#pragma acc data default(none) device_type(radeon) copy(Var)
  ;
  // expected-error@+2{{OpenACC clause 'copyin' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) copyin(Var)
+#pragma acc data default(none) device_type(radeon) copyin(Var)
  ;
  // expected-error@+2{{OpenACC clause 'copyout' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) copyout(Var)
+#pragma acc data default(none) device_type(radeon) copyout(Var)
  ;
  // expected-error@+2{{OpenACC clause 'create' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) create(Var)
+#pragma acc data default(none) device_type(radeon) create(Var)
  ;
  // expected-error@+2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) no_create(Var)
+#pragma acc data default(none) device_type(radeon) no_create(Var)
  ;
  // expected-error@+2{{OpenACC clause 'present' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) present(Var)
+#pragma acc data default(none) device_type(radeon) present(Var)
  ;
  // expected-error@+2{{OpenACC clause 'deviceptr' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) deviceptr(Var)
+#pragma acc data default(none) device_type(radeon) deviceptr(Var)
  ;
  // expected-error@+2{{OpenACC clause 'attach' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data default(none) device_type(foo) attach(Var)
+#pragma acc data default(none) device_type(radeon) attach(Var)
  ;
  // expected-error@+3{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}}
  // expected-error@+2{{OpenACC clause 'default' may not follow a 'device_type' clause in a 'data' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc data device_type(foo) default(none)
+#pragma acc data device_type(radeon) default(none)
  ;
 }
--- a/clang/test/SemaOpenACC/device_type_valid_values.c
+++ b/clang/test/SemaOpenACC/device_type_valid_values.c
@ -0,0 +1,56 @@
+// RUN: %clang_cc1 %s -fopenacc -verify 
+
+void foo(int i) {
+#pragma acc parallel device_type(*)
+  ;
+#pragma acc serial device_type(default)
+  ;
+#pragma acc kernels device_type(Default)
+  ;
+#pragma acc data copy(i) device_type(nvidia)
+  ;
+#pragma acc init device_type(acc_device_nvidia)
+  ;
+#pragma acc shutdown device_type(radeon)
+  ;
+#pragma acc set device_type(host)
+  ;
+#pragma acc update self(i) device_type(multicore)
+  ;
+#pragma acc loop device_type(DEFAULT)
+  for(unsigned I = 0; I < 5; ++I);
+#pragma acc parallel loop device_type(HOST)
+  for(unsigned I = 0; I < 5; ++I);
+#pragma acc serial loop device_type(MuLtIcOrE)
+  for(unsigned I = 0; I < 5; ++I);
+#pragma acc kernels loop device_type(radEon)
+  for(unsigned I = 0; I < 5; ++I);
+}
+
+#pragma acc routine(foo) seq device_type(host)
+
+#pragma acc routine seq device_type(acc_device_nvidia)
+void bar(){}
+
+// Anything outside of the above list is invalid.
+void invalid() {
+  // expected-error@+1{{invalid value 'other' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+#pragma acc parallel device_type(other)
+  ;
+  // expected-error@+1{{invalid value 'invalid' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+#pragma acc kernels device_type(invalid)
+  ;
+  // expected-error@+1{{invalid value 'invalid' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+#pragma acc kernels device_type(invalid, nvidia)
+  ;
+  // expected-error@+1{{invalid value 'invalid' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+#pragma acc kernels device_type(nvidia, invalid)
+  ;
+  // expected-error@+1{{invalid value 'invalid' in 'dtype' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+#pragma acc kernels dtype(nvidia, invalid, radeon)
+  ;
+  // expected-error@+2{{invalid value 'invalid' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+  // expected-error@+1{{invalid value 'invalid2' in 'device_type' clause; valid values are 'default', 'nvidia', 'acc_device_nvidia', 'radeon', 'host', 'multicore'}}
+#pragma acc kernels device_type(nvidia, invalid, radeon, invalid2)
+  ;
+}
--- a/clang/test/SemaOpenACC/init-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/init-construct-ast.cpp
@ -34,10 +34,10 @@ _Pragma("acc init")
  // CHECK-NEXT: CallExpr
  // CHECK-NEXT: ImplicitCastExpr
  // CHECK-NEXT: DeclRefExpr{{.*}}'some_int' 'int ()'
-#pragma acc init device_type(T)
+#pragma acc init device_type(default)
  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
-  // CHECK-NEXT: device_type(T)
-#pragma acc init if (some_int() < some_long()) device_type(T) device_num(some_int())
+  // CHECK-NEXT: device_type(default)
+#pragma acc init if (some_int() < some_long()) device_type(nvidia) device_num(some_int())
  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
@ -48,7 +48,7 @@ _Pragma("acc init")
  // CHECK-NEXT: CallExpr
  // CHECK-NEXT: ImplicitCastExpr
  // CHECK-NEXT: DeclRefExpr{{.*}}'some_long' 'long ()'
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(nvidia)
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: CallExpr
  // CHECK-NEXT: ImplicitCastExpr
@ -76,17 +76,17 @@ void TemplFunc(T t) {
  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
-#pragma acc init device_type(T)
+#pragma acc init device_type(radeon)
  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
-  // CHECK-NEXT: device_type(T)
-#pragma acc init if (T::value > t) device_type(T) device_num(t)
+  // CHECK-NEXT: device_type(radeon)
+#pragma acc init if (T::value > t) device_type(host) device_num(t)
  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '>'
  // CHECK-NEXT: DependentScopeDeclRefExpr
  // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(host)
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'

@ -119,7 +119,7 @@ void TemplFunc(T t) {
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'

  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(radeon)

  // CHECK-NEXT: OpenACCInitConstruct{{.*}}init
  // CHECK-NEXT: if clause
@ -131,7 +131,7 @@ void TemplFunc(T t) {
  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'unsigned int'
  // CHECK-NEXT: MemberExpr{{.*}} .operator unsigned int
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(host)
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: ImplicitCastExpr{{.*}}'unsigned int'
  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'unsigned int'
--- a/clang/test/SemaOpenACC/init-construct.cpp
+++ b/clang/test/SemaOpenACC/init-construct.cpp
@ -18,9 +18,9 @@ void uses() {
 #pragma acc init
 #pragma acc init if (getI() < getS())
 #pragma acc init device_num(getI())
-#pragma acc init device_type(SOMETHING) device_num(getI())
-#pragma acc init device_type(SOMETHING) if (getI() < getS())
-#pragma acc init device_type(SOMETHING) device_num(getI()) if (getI() < getS())
+#pragma acc init device_type(host) device_num(getI())
+#pragma acc init device_type(nvidia) if (getI() < getS())
+#pragma acc init device_type(multicore) device_num(getI()) if (getI() < getS())

  // expected-error@+1{{value of type 'struct NotConvertible' is not contextually convertible to 'bool'}}
 #pragma acc init if (NC)
@ -41,8 +41,8 @@ void TestInst() {
  T t;
 #pragma acc init
 #pragma acc init if (T::value < T{})
-#pragma acc init device_type(SOMETHING) device_num(getI()) if (getI() < getS())
-#pragma acc init device_type(SOMETHING) device_type(T) device_num(t) if (t < T::value) device_num(getI()) if (getI() < getS())
+#pragma acc init device_type(radeon) device_num(getI()) if (getI() < getS())
+#pragma acc init device_type(multicore) device_type(host) device_num(t) if (t < T::value) device_num(getI()) if (getI() < getS())

  // expected-error@+1{{value of type 'const NotConvertible' is not contextually convertible to 'bool'}}
 #pragma acc init if (T::NCValue)
--- a/clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-ast.cpp
@ -19,11 +19,11 @@ void NormalUses() {
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} SomeVar 'bool'

-#pragma acc loop device_type(SomeS) dtype(SomeImpl)
+#pragma acc loop device_type(default) dtype(nvidia)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(SomeS)
-  // CHECK-NEXT: dtype(SomeImpl)
+  // CHECK-NEXT: device_type(default)
+  // CHECK-NEXT: dtype(nvidia)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -36,11 +36,11 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc loop device_type(SomeVar) dtype(int)
+#pragma acc loop device_type(acc_device_nvidia) dtype(radeon)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(SomeVar)
-  // CHECK-NEXT: dtype(int)
+  // CHECK-NEXT: device_type(acc_device_nvidia)
+  // CHECK-NEXT: dtype(radeon)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -53,11 +53,11 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc loop device_type(private) dtype(struct)
+#pragma acc loop device_type(host) dtype(multicore)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(private)
-  // CHECK-NEXT: dtype(struct)
+  // CHECK-NEXT: device_type(host)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -70,27 +70,10 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc loop device_type(private) dtype(class)
+#pragma acc loop device_type(default) dtype(*)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(private)
-  // CHECK-NEXT: dtype(class)
-  // CHECK-NEXT: ForStmt
-  // CHECK-NEXT: DeclStmt
-  // CHECK-NEXT: VarDecl{{.*}} i 'int'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
-  // CHECK-NEXT: <<<NULL>>>
-  // CHECK-NEXT: BinaryOperator{{.*}}'<'
-  // CHECK-NEXT: ImplicitCastExpr
-  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
-  // CHECK-NEXT: UnaryOperator{{.*}}++
-  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
-  // CHECK-NEXT: CompoundStmt
-#pragma acc loop device_type(float) dtype(*)
-  for(int i = 0; i < 5; ++i){}
-  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(float)
+  // CHECK-NEXT: device_type(default)
  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
@ -104,10 +87,27 @@ void NormalUses() {
  // CHECK-NEXT: UnaryOperator{{.*}}++
  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
  // CHECK-NEXT: CompoundStmt
-#pragma acc loop device_type(float, int) dtype(*)
+#pragma acc loop device_type(nvidia) dtype(*)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(float, int)
+  // CHECK-NEXT: device_type(nvidia)
+  // CHECK-NEXT: dtype(*)
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: CompoundStmt
+#pragma acc loop device_type(radeon, host) dtype(*)
+  for(int i = 0; i < 5; ++i){}
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
+  // CHECK-NEXT: device_type(radeon, host)
  // CHECK-NEXT: dtype(*)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
@ -129,11 +129,11 @@ void TemplUses() {
  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
  // CHECK-NEXT: CompoundStmt
-#pragma acc loop device_type(T) dtype(T)
+#pragma acc loop device_type(default) dtype(multicore)
  for(int i = 0; i < 5; ++i){}
  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(default)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
@ -155,8 +155,8 @@ void TemplUses() {
  // CHECK-NEXT: CompoundStmt

  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(T)
+  // CHECK-NEXT: device_type(default)
+  // CHECK-NEXT: dtype(multicore)
  // CHECK-NEXT: ForStmt
  // CHECK-NEXT: DeclStmt
  // CHECK-NEXT: VarDecl{{.*}} i 'int'
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@ -6,22 +6,22 @@ void uses() {
  typedef struct S{} STy;
  STy SImpl;

-#pragma acc loop device_type(I)
+#pragma acc loop device_type(radeon)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop device_type(S) dtype(STy)
+#pragma acc loop device_type(host) dtype(multicore)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop dtype(SImpl)
+#pragma acc loop dtype(nvidia)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop dtype(int) device_type(*)
+#pragma acc loop dtype(default) device_type(*)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop dtype(true) device_type(false)
+#pragma acc loop dtype(radeon) device_type(nvidia)
  for(int i = 0; i < 5; ++i);

  // expected-error@+1{{expected identifier}}
 #pragma acc loop dtype(int, *)
  for(int i = 0; i < 5; ++i);

-#pragma acc loop device_type(I, int)
+#pragma acc loop device_type(default, nvidia)
  for(int i = 0; i < 5; ++i);
  // expected-error@+2{{expected ','}}
  // expected-error@+1{{expected identifier}}
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.cpp
@ -2,17 +2,17 @@

 template<typename T>
 void TemplUses() {
-#pragma acc loop device_type(I)
+#pragma acc loop device_type(host)
  for(int i = 0; i < 5; ++i);
 #pragma acc loop dtype(*)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop device_type(class)
+#pragma acc loop device_type(default)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop device_type(private)
+#pragma acc loop device_type(multicore)
  for(int i = 0; i < 5; ++i);
-#pragma acc loop device_type(bool)
+#pragma acc loop device_type(host)
  for(int i = 0; i < 5; ++i);
-#pragma acc kernels dtype(true) device_type(false)
+#pragma acc kernels dtype(radeon) device_type(nvidia)
  for(int i = 0; i < 5; ++i);
  // expected-error@+2{{expected ','}}
  // expected-error@+1{{expected identifier}}
--- a/clang/test/SemaOpenACC/routine-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp
@ -24,16 +24,16 @@ int function();
 // CHECK-NEXT: vector clause
 // CHECK-NEXT: bind clause identifier 'identifier'

-#pragma acc routine(function) device_type(Something) seq
+#pragma acc routine(function) device_type(acc_device_nvidia) seq
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
 // CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
-// CHECK-NEXT: device_type(Something)
+// CHECK-NEXT: device_type(acc_device_nvidia)
 // CHECK-NEXT: seq clause
-#pragma acc routine(function) nohost dtype(Something) vector
+#pragma acc routine(function) nohost dtype(nvidia) vector
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
 // CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
 // CHECK-NEXT: nohost clause
-// CHECK-NEXT: dtype(Something)
+// CHECK-NEXT: dtype(nvidia)
 // CHECK-NEXT: vector clause

 #pragma acc routine nohost vector
@ -172,15 +172,15 @@ struct S {
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
 // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
 // CHECK-NEXT: worker clause
-#pragma acc routine(Lambda) worker device_type(Lambda)
+#pragma acc routine(Lambda) worker device_type(host)
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
 // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
 // CHECK-NEXT: worker clause
-// CHECK-NEXT: device_type(Lambda)
-#pragma acc routine(Lambda) dtype(Lambda) vector
+// CHECK-NEXT: device_type(host)
+#pragma acc routine(Lambda) dtype(radeon) vector
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
 // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
-// CHECK-NEXT: dtype(Lambda)
+// CHECK-NEXT: dtype(radeon)
 // CHECK-NEXT: vector clause
 };

@ -294,12 +294,12 @@ struct DepS {
 // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
 // CHECK-NEXT: worker clause
-#pragma acc routine(DepS<T>::StaticMemFunc) worker device_type(T)
+#pragma acc routine(DepS<T>::StaticMemFunc) worker device_type(multicore)
 // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine 
 // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
 // CHECK-NEXT: worker clause
-// CHECK-NEXT: device_type(T)
+// CHECK-NEXT: device_type(multicore)

 // Instantiation:
 // CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS
@ -398,7 +398,7 @@ struct DepS {
 // CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
 // CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
 // CHECK-NEXT: worker clause
-// CHECK-NEXT: device_type(T)
+// CHECK-NEXT: device_type(multicore)
 };

 #pragma acc routine(DepS<int>::Lambda) gang(dim:1)
--- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
@ -139,8 +139,8 @@ void Inst() {
  DependentT<HasFuncs> T;// expected-note{{in instantiation of}}
 }

-#pragma acc routine(Func) gang device_type(Inst)
-#pragma acc routine(Func) gang dtype(Inst)
+#pragma acc routine(Func) gang device_type(host)
+#pragma acc routine(Func) gang dtype(multicore)
 #pragma acc routine(Func) device_type(*) worker
 #pragma acc routine(Func) dtype(*) worker

--- a/clang/test/SemaOpenACC/set-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/set-construct-ast.cpp
@ -13,7 +13,7 @@ void NormalFunc() {
  // CHECK-LABEL: NormalFunc
  // CHECK-NEXT: CompoundStmt

-#pragma acc set default_async(some_int()) device_num(some_long()) device_type(DT) if (some_int() < some_long())
+#pragma acc set default_async(some_int()) device_num(some_long()) device_type(acc_device_nvidia) if (some_int() < some_long())
  // CHECK-NEXT: OpenACCSetConstruct{{.*}}set
  // CHECK-NEXT: default_async clause
  // CHECK-NEXT: CallExpr{{.*}}'int'
@ -23,7 +23,7 @@ void NormalFunc() {
  // CHECK-NEXT: CallExpr{{.*}} 'long'
  // CHECK-NEXT: ImplicitCastExpr
  // CHECK-NEXT: DeclRefExpr{{.*}}'some_long' 'long ()'
-  // CHECK-NEXT: device_type(DT)
+  // CHECK-NEXT: device_type(acc_device_nvidia)
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
  // CHECK-NEXT: ImplicitCastExpr{{.*}}'long'
@ -43,14 +43,14 @@ void TemplFunc(T t) {
  // CHECK-NEXT: ParmVarDecl{{.*}} t 'T'
  // CHECK-NEXT: CompoundStmt

-#pragma acc set default_async(T::value) device_num(t) device_type(DT) if (T::value < t)
+#pragma acc set default_async(T::value) device_num(t) device_type(multicore) if (T::value < t)
  // CHECK-NEXT: OpenACCSetConstruct{{.*}}set
  // CHECK-NEXT: default_async clause
  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>'
  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: DeclRefExpr{{.*}}'t' 'T'
-  // CHECK-NEXT: device_type(DT)
+  // CHECK-NEXT: device_type(multicore)
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>'
@ -75,7 +75,7 @@ void TemplFunc(T t) {
  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'unsigned int'
  // CHECK-NEXT: MemberExpr{{.*}}.operator unsigned int
  // CHECK-NEXT: DeclRefExpr{{.*}}'t' 'SomeStruct'
-  // CHECK-NEXT: device_type(DT)
+  // CHECK-NEXT: device_type(multicore)
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
  // CHECK-NEXT: ImplicitCastExpr {{.*}}'unsigned int'
--- a/clang/test/SemaOpenACC/set-construct.cpp
+++ b/clang/test/SemaOpenACC/set-construct.cpp
@ -17,11 +17,11 @@ struct ExplicitConvertOnly {
 void uses() {
 #pragma acc set default_async(getI())
 #pragma acc set device_num(getI())
-#pragma acc set device_type(getI)
-#pragma acc set device_type(getI) if (getI() < getS())
+#pragma acc set device_type(multicore)
+#pragma acc set device_type(host) if (getI() < getS())

  // expected-error@+1{{value of type 'struct NotConvertible' is not contextually convertible to 'bool'}}
-#pragma acc set if (NC) device_type(I)
+#pragma acc set if (NC) device_type(radeon)

  // expected-error@+2{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}}
  // expected-error@+1{{OpenACC clause 'device_num' requires expression of integer type ('struct NotConvertible' invalid)}}
@ -62,8 +62,8 @@ void uses() {

  // expected-error@+2{{'device_type' clause cannot appear more than once on a 'set' directive}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc set device_type(I) device_type(I)
+#pragma acc set device_type(nvidia) device_type(default)
  // expected-error@+2{{'if' clause cannot appear more than once on a 'set' directive}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc set device_type(I) if(true) if (true)
+#pragma acc set device_type(acc_device_nvidia) if(true) if (true)
 }
--- a/clang/test/SemaOpenACC/shutdown-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/shutdown-construct-ast.cpp
@ -34,10 +34,10 @@ _Pragma("acc shutdown")
  // CHECK-NEXT: CallExpr
  // CHECK-NEXT: ImplicitCastExpr
  // CHECK-NEXT: DeclRefExpr{{.*}}'some_int' 'int ()'
-#pragma acc shutdown device_type(T)
+#pragma acc shutdown device_type(multicore)
  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
-  // CHECK-NEXT: device_type(T)
-#pragma acc shutdown if (some_int() < some_long()) device_type(T) device_num(some_int())
+  // CHECK-NEXT: device_type(multicore)
+#pragma acc shutdown if (some_int() < some_long()) device_type(host) device_num(some_int())
  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
@ -48,7 +48,7 @@ _Pragma("acc shutdown")
  // CHECK-NEXT: CallExpr
  // CHECK-NEXT: ImplicitCastExpr
  // CHECK-NEXT: DeclRefExpr{{.*}}'some_long' 'long ()'
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(host)
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: CallExpr
  // CHECK-NEXT: ImplicitCastExpr
@ -76,17 +76,17 @@ void TemplFunc(T t) {
  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
-#pragma acc shutdown device_type(T)
+#pragma acc shutdown device_type(radeon)
  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
-  // CHECK-NEXT: device_type(T)
-#pragma acc shutdown if (T::value > t) device_type(T) device_num(t)
+  // CHECK-NEXT: device_type(radeon)
+#pragma acc shutdown if (T::value > t) device_type(acc_device_nvidia) device_num(t)
  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
  // CHECK-NEXT: if clause
  // CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '>'
  // CHECK-NEXT: DependentScopeDeclRefExpr
  // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(acc_device_nvidia)
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'

@ -119,7 +119,7 @@ void TemplFunc(T t) {
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'

  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(radeon)

  // CHECK-NEXT: OpenACCShutdownConstruct{{.*}}shutdown
  // CHECK-NEXT: if clause
@ -131,7 +131,7 @@ void TemplFunc(T t) {
  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'unsigned int'
  // CHECK-NEXT: MemberExpr{{.*}} .operator unsigned int
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
-  // CHECK-NEXT: device_type(T)
+  // CHECK-NEXT: device_type(acc_device_nvidia)
  // CHECK-NEXT: device_num clause
  // CHECK-NEXT: ImplicitCastExpr{{.*}}'unsigned int'
  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'unsigned int'
--- a/clang/test/SemaOpenACC/shutdown-construct.cpp
+++ b/clang/test/SemaOpenACC/shutdown-construct.cpp
@ -18,9 +18,9 @@ void uses() {
 #pragma acc shutdown
 #pragma acc shutdown if (getI() < getS())
 #pragma acc shutdown device_num(getI())
-#pragma acc shutdown device_type(SOMETHING) device_num(getI())
-#pragma acc shutdown device_type(SOMETHING) if (getI() < getS())
-#pragma acc shutdown device_type(SOMETHING) device_num(getI()) if (getI() < getS())
+#pragma acc shutdown device_type(host) device_num(getI())
+#pragma acc shutdown device_type(nvidia) if (getI() < getS())
+#pragma acc shutdown device_type(radeon) device_num(getI()) if (getI() < getS())

  // expected-error@+1{{value of type 'struct NotConvertible' is not contextually convertible to 'bool'}}
 #pragma acc shutdown if (NC)
@ -41,8 +41,8 @@ void TestInst() {
  T t;
 #pragma acc shutdown
 #pragma acc shutdown if (T::value < T{})
-#pragma acc shutdown device_type(SOMETHING) device_num(getI()) if (getI() < getS())
-#pragma acc shutdown device_type(SOMETHING) device_type(T) device_num(t) if (t < T::value) device_num(getI()) if (getI() < getS())
+#pragma acc shutdown device_type(multicore) device_num(getI()) if (getI() < getS())
+#pragma acc shutdown device_type(default) device_type(radeon) device_num(t) if (t < T::value) device_num(getI()) if (getI() < getS())

  // expected-error@+1{{value of type 'const NotConvertible' is not contextually convertible to 'bool'}}
 #pragma acc shutdown if (T::NCValue)
--- a/clang/test/SemaOpenACC/update-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/update-construct-ast.cpp
@ -33,15 +33,15 @@ void NormalFunc() {
  // CHECK-NEXT: ImplicitCastExpr
  // CHECK-NEXT: DeclRefExpr{{.*}}'some_long' 'long ()'

-#pragma acc update self(Global) wait async device_type(A) dtype(B)
+#pragma acc update self(Global) wait async device_type(nvidia) dtype(radeon)
  // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
  // CHECK-NEXT: self clause
  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
  // CHECK-NEXT: wait clause
  // CHECK-NEXT: <<<NULL>>>
  // CHECK-NEXT: async clause
-  // CHECK-NEXT: device_type(A)
-  // CHECK-NEXT: dtype(B)
+  // CHECK-NEXT: device_type(nvidia)
+  // CHECK-NEXT: dtype(radeon)
 #pragma acc update self(Global) wait(some_int(), some_long()) async(some_int())
  // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
  // CHECK-NEXT: self clause
@ -147,15 +147,15 @@ void TemplFunc(T t) {
  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
  // CHECK-NEXT: DeclRefExpr{{.*}}'t' 'T'

-#pragma acc update self(t) wait async device_type(T) dtype(U)
+#pragma acc update self(t) wait async device_type(default) dtype(host)
  // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
  // CHECK-NEXT: self clause
  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
  // CHECK-NEXT: wait clause
  // CHECK-NEXT: <<<NULL>>>
  // CHECK-NEXT: async clause
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(U)
+  // CHECK-NEXT: device_type(default)
+  // CHECK-NEXT: dtype(host)
 #pragma acc update self(t) wait(T::value, t) async(T::value)
  // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
  // CHECK-NEXT: self clause
@ -261,8 +261,8 @@ void TemplFunc(T t) {
  // CHECK-NEXT: wait clause
  // CHECK-NEXT: <<<NULL>>>
  // CHECK-NEXT: async clause
-  // CHECK-NEXT: device_type(T)
-  // CHECK-NEXT: dtype(U)
+  // CHECK-NEXT: device_type(default)
+  // CHECK-NEXT: dtype(host)

  // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
  // CHECK-NEXT: self clause
--- a/clang/test/SemaOpenACC/update-construct.cpp
+++ b/clang/test/SemaOpenACC/update-construct.cpp
@ -6,7 +6,7 @@ void uses() {
  int Var;
 #pragma acc update async self(Var)
 #pragma acc update wait self(Var)
-#pragma acc update self(Var) device_type(I)
+#pragma acc update self(Var) device_type(host)
 #pragma acc update if(true) self(Var)
 #pragma acc update if_present self(Var)
 #pragma acc update self(Var)
@ -15,24 +15,24 @@ void uses() {

  // expected-error@+2{{OpenACC clause 'if' may not follow a 'device_type' clause in a 'update' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc update self(Var) device_type(I) if(true)
+#pragma acc update self(Var) device_type(default) if(true)
  // expected-error@+2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a 'update' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc update self(Var) device_type(I) if_present
+#pragma acc update self(Var) device_type(multicore) if_present
  // expected-error@+2{{OpenACC clause 'self' may not follow a 'device_type' clause in a 'update' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc update self(Var) device_type(I) self(Var)
+#pragma acc update self(Var) device_type(nvidia) self(Var)
  // expected-error@+2{{OpenACC clause 'host' may not follow a 'device_type' clause in a 'update' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc update self(Var) device_type(I) host(Var)
+#pragma acc update self(Var) device_type(radeon) host(Var)
  // expected-error@+2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'update' construct}}
  // expected-note@+1{{previous clause is here}}
-#pragma acc update self(Var) device_type(I) device(Var)
+#pragma acc update self(Var) device_type(multicore) device(Var)
  // These 2 are OK.
-#pragma acc update self(Var) device_type(I) async
-#pragma acc update self(Var) device_type(I) wait
+#pragma acc update self(Var) device_type(default) async
+#pragma acc update self(Var) device_type(nvidia) wait
  // Unless otherwise specified, we assume 'device_type' can happen after itself.
-#pragma acc update self(Var) device_type(I) device_type(I)
+#pragma acc update self(Var) device_type(acc_device_nvidia) device_type(multicore)

  // These diagnose because there isn't at least 1 of 'self', 'host', or
  // 'device'.
@ -41,14 +41,14 @@ void uses() {
  // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}}
 #pragma acc update wait
  // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}}
-#pragma acc update device_type(I)
+#pragma acc update device_type(host)
  // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}}
 #pragma acc update if(true)
  // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}}
 #pragma acc update if_present

  // expected-error@+1{{value of type 'struct NotConvertible' is not contextually convertible to 'bool'}}
-#pragma acc update self(Var) if (NC) device_type(I)
+#pragma acc update self(Var) if (NC) device_type(radeon)

  // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'update' directive}}
  // expected-note@+1{{previous clause is here}}