[OpenACC] Implement 'at least one of' restriction for 'update'

This completes the implementation of 'update' by implementing its last restriction. This restriction requires at least 1 of the 'self', 'host', or 'device' clauses.
arm · Jan 9, 2025 · 553fa20 · 553fa20
1 parent be32621
commit 553fa20
Show file tree

Hide file tree

Showing 6 changed files with 129 additions and 87 deletions.
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
@@ -3732,8 +3732,17 @@ bool SemaOpenACC::ActOnStartStmtDirective(
                                 OpenACCClauseKind::DeviceType,
                                 OpenACCClauseKind::If});
 
-  // TODO: OpenACC: 'Update' construct needs to have one of 'self', 'host', or
-  // 'device'.  Implement here.
+  // OpenACC3.3 2.14.4: At least one self, host, or device clause must appear on
+  // an update directive.
+  if (K == OpenACCDirectiveKind::Update &&
+      llvm::find_if(Clauses, llvm::IsaPred<OpenACCSelfClause, OpenACCHostClause,
+                                           OpenACCDeviceClause>) ==
+          Clauses.end())
+    return Diag(StartLoc, diag::err_acc_construct_one_clause_of)
+           << K
+           << GetListOfClauses({OpenACCClauseKind::Self,
+                                OpenACCClauseKind::Host,
+                                OpenACCClauseKind::Device});
 
   return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/true);
 }

diff --git a/clang/test/AST/ast-print-openacc-update-construct.cpp b/clang/test/AST/ast-print-openacc-update-construct.cpp
@@ -3,38 +3,38 @@ void uses(bool cond) {
   int I;
   int *iPtr;
   int array[5];
-  // CHECK: #pragma acc update
-#pragma acc update
+  // CHECK: #pragma acc update self(I)
+#pragma acc update self(I)
 
-// CHECK: #pragma acc update if_present
-#pragma acc update if_present
-// CHECK: #pragma acc update if(cond)
-#pragma acc update if(cond)
+// CHECK: #pragma acc update self(I) if_present
+#pragma acc update self(I) if_present
+// CHECK: #pragma acc update self(I) if(cond)
+#pragma acc update self(I) if(cond)
 
-// CHECK: #pragma acc update async
-#pragma acc update async
-// CHECK: #pragma acc update async(*iPtr)
-#pragma acc update async(*iPtr)
-// CHECK: #pragma acc update async(I)
-#pragma acc update async(I)
+// CHECK: #pragma acc update self(I) async
+#pragma acc update self(I) async
+// CHECK: #pragma acc update self(I) async(*iPtr)
+#pragma acc update self(I) async(*iPtr)
+// CHECK: #pragma acc update self(I) async(I)
+#pragma acc update self(I) async(I)
 
-// CHECK: #pragma acc update wait(*iPtr, I) async
-#pragma acc update wait(*iPtr, I) async
+// CHECK: #pragma acc update self(I) wait(*iPtr, I) async
+#pragma acc update self(I) wait(*iPtr, I) async
 
-// CHECK: #pragma acc update wait(queues: *iPtr, I) async(*iPtr)
-#pragma acc update wait(queues:*iPtr, I) async(*iPtr)
+// CHECK: #pragma acc update self(I) wait(queues: *iPtr, I) async(*iPtr)
+#pragma acc update self(I) wait(queues:*iPtr, I) async(*iPtr)
 
-// CHECK: #pragma acc update wait(devnum: I : *iPtr, I) async(I)
-#pragma acc update wait(devnum:I:*iPtr, I) async(I)
+// CHECK: #pragma acc update self(I) wait(devnum: I : *iPtr, I) async(I)
+#pragma acc update self(I) wait(devnum:I:*iPtr, I) async(I)
 
-// CHECK: #pragma acc update wait(devnum: I : queues: *iPtr, I) if(I == array[I]) async(I)
-#pragma acc update wait(devnum:I:queues:*iPtr, I) if(I == array[I]) async(I)
+// 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 device_type(I) dtype(H)
-#pragma acc update device_type(I) dtype(H)
+// 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 device_type(J) dtype(K)
-#pragma acc update device_type(J) dtype(K)
+// 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, iPtr, array, array[1], array[1:2])
 #pragma acc update self(I, iPtr, array, array[1], array[1:2])

diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
@@ -344,7 +344,7 @@ void SelfUpdate() {
   struct Members s;
 
   // expected-error@+1{{expected '('}}
-#pragma acc update self
+#pragma acc update host(s) self
   for(int i = 0; i < 5;++i) {}
 
   // expected-error@+3{{use of undeclared identifier 'zero'}}

diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
@@ -151,6 +151,7 @@ void func() {
   // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}}
 #pragma acc set clause list
   for(;;){}
+  // expected-error@+2{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}}
   // expected-error@+1{{invalid OpenACC clause 'clause'}}
 #pragma acc update clause list
   for(;;){}

diff --git a/clang/test/SemaOpenACC/update-construct-ast.cpp b/clang/test/SemaOpenACC/update-construct-ast.cpp
@@ -18,8 +18,10 @@ void NormalFunc() {
   // CHECK-LABEL: NormalFunc
   // CHECK-NEXT: CompoundStmt
 
-#pragma acc update if_present if (some_int() < some_long())
+#pragma acc update self(Global) if_present if (some_int() < some_long())
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
   // CHECK-NEXT: if_present clause
   // CHECK-NEXT: if clause
   // CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
@@ -31,15 +33,19 @@ void NormalFunc() {
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'some_long' 'long ()'
 
-#pragma acc update wait async device_type(A) dtype(B)
+#pragma acc update self(Global) wait async device_type(A) dtype(B)
   // 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)
-#pragma acc update wait(some_int(), some_long()) async(some_int())
+#pragma acc update self(Global) wait(some_int(), some_long()) async(some_int())
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: CallExpr{{.*}}'int'
@@ -52,8 +58,10 @@ void NormalFunc() {
   // CHECK-NEXT: CallExpr{{.*}}'int'
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'some_int' 'int ()'
-#pragma acc update wait(queues:some_int(), some_long())
+#pragma acc update self(Global) wait(queues:some_int(), some_long())
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: CallExpr{{.*}}'int'
@@ -62,8 +70,10 @@ void NormalFunc() {
   // CHECK-NEXT: CallExpr{{.*}}'long'
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'some_long' 'long ()'
-#pragma acc update wait(devnum: some_int() :some_int(), some_long())
+#pragma acc update self(Global) wait(devnum: some_int() :some_int(), some_long())
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: CallExpr{{.*}}'int'
   // CHECK-NEXT: ImplicitCastExpr
@@ -83,12 +93,12 @@ void NormalFunc() {
   // CHECK-NEXT: ArraySubscriptExpr{{.*}} 'short' lvalue
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 0
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 0
-  // CHECK-NEXT: IntegerLiteral{{.*}} 1
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
 
 #pragma acc update host(Global, GlobalArray, GlobalArray[0], GlobalArray[0:1]) device(Global, GlobalArray, GlobalArray[0], GlobalArray[0:1])
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
@@ -98,24 +108,24 @@ void NormalFunc() {
   // CHECK-NEXT: ArraySubscriptExpr{{.*}} 'short' lvalue
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 0
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 0
-  // CHECK-NEXT: IntegerLiteral{{.*}} 1
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
   // CHECK-NEXT: device clause
   // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
   // CHECK-NEXT: ArraySubscriptExpr{{.*}} 'short' lvalue
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 0
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'short[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}} 0
-  // CHECK-NEXT: IntegerLiteral{{.*}} 1
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
 }
 
 template<typename T>
@@ -126,24 +136,30 @@ void TemplFunc(T t) {
   // CHECK-NEXT: ParmVarDecl{{.*}} t 'T'
   // CHECK-NEXT: CompoundStmt
 
-#pragma acc update if_present if (T::value < t)
+#pragma acc update self(t) if_present if (T::value < t)
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
   // CHECK-NEXT: if_present clause
   // CHECK-NEXT: if clause
   // CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '<'
   // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>'
   // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
   // CHECK-NEXT: DeclRefExpr{{.*}}'t' 'T'
 
-#pragma acc update wait async device_type(T) dtype(U)
+#pragma acc update self(t) wait async device_type(T) dtype(U)
   // 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)
-#pragma acc update wait(T::value, t) async(T::value)
+#pragma acc update self(t) wait(T::value, t) async(T::value)
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
@@ -152,17 +168,21 @@ void TemplFunc(T t) {
   // CHECK-NEXT: async clause
   // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
   // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
-#pragma acc update wait(queues:T::value, t) async(t)
+#pragma acc update self(t) wait(queues:T::value, t) async(t)
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
   // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
   // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
   // CHECK-NEXT: async clause
   // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
-#pragma acc update wait(devnum: T::value:t, T::value)
+#pragma acc update self(t) wait(devnum: T::value:t, T::value)
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'T'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
   // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
@@ -184,11 +204,11 @@ void TemplFunc(T t) {
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
   // CHECK-NEXT: ArraySubscriptExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
-  // CHECK-NEXT: IntegerLiteral{{.*}}1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
 
 #pragma acc update host(Local, LocalArray, LocalArray[0], LocalArray[0:1]) device(Local, LocalArray, LocalArray[0], LocalArray[0:1])
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
@@ -197,21 +217,21 @@ void TemplFunc(T t) {
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
   // CHECK-NEXT: ArraySubscriptExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
-  // CHECK-NEXT: IntegerLiteral{{.*}}1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
   // CHECK-NEXT: device clause
   // CHECK-NEXT: DeclRefExpr{{.*}} 'Local' 'decltype(T::value)'
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
   // CHECK-NEXT: ArraySubscriptExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(T::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
-  // CHECK-NEXT: IntegerLiteral{{.*}}1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
 
   // Instantiation:
   // CHECK-NEXT: FunctionDecl{{.*}} TemplFunc 'void (SomeStruct)' implicit_instantiation
@@ -222,6 +242,8 @@ void TemplFunc(T t) {
   // CHECK-NEXT: CompoundStmt
 
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
   // CHECK-NEXT: if_present clause
   // CHECK-NEXT: if clause
   // CHECK-NEXT: BinaryOperator{{.*}}'bool' '<'
@@ -234,13 +256,17 @@ void TemplFunc(T t) {
   // CHECK-NEXT: DeclRefExpr{{.*}}'t' 'SomeStruct'
 
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: async clause
   // CHECK-NEXT: device_type(T)
   // CHECK-NEXT: dtype(U)
 
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: ImplicitCastExpr{{.*}}'unsigned int'
@@ -256,6 +282,8 @@ void TemplFunc(T t) {
   // CHECK-NEXT: NestedNameSpecifier TypeSpec 'SomeStruct'
 
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: <<<NULL>>>
   // CHECK-NEXT: ImplicitCastExpr{{.*}}'unsigned int'
@@ -272,6 +300,8 @@ void TemplFunc(T t) {
   // CHECK-NEXT: DeclRefExpr{{.*}}'t' 'SomeStruct'
 
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
+  // CHECK-NEXT: self clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 't' 'SomeStruct'
   // CHECK-NEXT: wait clause
   // CHECK-NEXT: ImplicitCastExpr{{.*}}'unsigned int'
   // CHECK-NEXT: DeclRefExpr{{.*}}'value' 'const unsigned int'
@@ -299,12 +329,12 @@ void TemplFunc(T t) {
   // CHECK-NEXT: ArraySubscriptExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
-  // CHECK-NEXT: IntegerLiteral{{.*}}1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
 
   // CHECK-NEXT: OpenACCUpdateConstruct{{.*}}update
   // CHECK-NEXT: host clause
@@ -313,24 +343,24 @@ void TemplFunc(T t) {
   // CHECK-NEXT: ArraySubscriptExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
-  // CHECK-NEXT: IntegerLiteral{{.*}}1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
   // CHECK-NEXT: device clause
   // CHECK-NEXT: DeclRefExpr{{.*}} 'Local' 'decltype(SomeStruct::value)':'const unsigned int'
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
   // CHECK-NEXT: ArraySubscriptExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
   // CHECK-NEXT: ArraySectionExpr
   // CHECK-NEXT: ImplicitCastExpr
   // CHECK-NEXT: DeclRefExpr{{.*}} 'LocalArray' 'decltype(SomeStruct::value)[5]'
-  // CHECK-NEXT: IntegerLiteral{{.*}}0
-  // CHECK-NEXT: IntegerLiteral{{.*}}1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 0
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
 }
 
 struct SomeStruct{