[clang] [OpenMP] Do not GEP 'inbounds' for multi-dimensional OpenMP device maps (PR #218436)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 13:50:23 PDT 2026
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/218436
>From 3309f46ff171692856e281d4b2bf57ed9a89db08 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Sun, 23 Aug 2026 20:14:54 -0500
Subject: [PATCH 1/3] [OpenMP] Do not GEP 'inbounds' for multi-dimensional
OpenMP device maps
The C standard treats array subscripts as inbounds of the object, but
OpenMP section maps intentionally index from a host-relative base that
is not the mapped allocation. Drop inbounds on those GEPs so later
passes do not assume the slice lives inside the object.
---
clang/lib/CodeGen/CGExprScalar.cpp | 30 ++++++++++++++-----
.../OpenMP/target_map_section_inbounds.cpp | 18 +++++++++++
2 files changed, 41 insertions(+), 7 deletions(-)
create mode 100644 clang/test/OpenMP/target_map_section_inbounds.cpp
diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp
index 173931213ce68..7d848feed4eda 100644
--- a/clang/lib/CodeGen/CGExprScalar.cpp
+++ b/clang/lib/CodeGen/CGExprScalar.cpp
@@ -35,6 +35,7 @@
#include "clang/Basic/TargetInfo.h"
#include "llvm/ADT/APFixedPoint.h"
#include "llvm/ADT/ScopeExit.h"
+#include "llvm/Analysis/ValueTracking.h"
#include "llvm/IR/Argument.h"
#include "llvm/IR/CFG.h"
#include "llvm/IR/Constants.h"
@@ -6434,6 +6435,23 @@ static GEPOffsetAndOverflow EmitGEPOffsetInBytes(Value *BasePtr, Value *GEPVal,
return {TotalOffset, OffsetOverflows};
}
+// OpenMP section maps pass `section - host_offset`. A GEP of an array-of-arrays
+// from that pointer is not inbounds of the allocation.
+static llvm::GEPNoWrapFlags
+inBoundsGEPFlags(const CodeGenFunction &CGF, const llvm::Value *Ptr,
+ llvm::Type *SrcTy, bool SignedIndices, bool IsSubtraction) {
+ llvm::GEPNoWrapFlags NW;
+ auto *AT = dyn_cast<llvm::ArrayType>(SrcTy);
+ bool SectionGEP = CGF.getLangOpts().OpenMPIsTargetDevice && AT &&
+ AT->getElementType()->isArrayTy() &&
+ !isa<llvm::AllocaInst>(llvm::getUnderlyingObject(Ptr));
+ if (!SectionGEP)
+ NW = llvm::GEPNoWrapFlags::inBounds();
+ if (!SignedIndices && !IsSubtraction)
+ NW |= llvm::GEPNoWrapFlags::noUnsignedWrap();
+ return NW;
+}
+
Value *
CodeGenFunction::EmitCheckedInBoundsGEP(llvm::Type *ElemTy, Value *Ptr,
ArrayRef<Value *> IdxList,
@@ -6441,9 +6459,8 @@ CodeGenFunction::EmitCheckedInBoundsGEP(llvm::Type *ElemTy, Value *Ptr,
SourceLocation Loc, const Twine &Name) {
llvm::Type *PtrTy = Ptr->getType();
- llvm::GEPNoWrapFlags NWFlags = llvm::GEPNoWrapFlags::inBounds();
- if (!SignedIndices && !IsSubtraction)
- NWFlags |= llvm::GEPNoWrapFlags::noUnsignedWrap();
+ llvm::GEPNoWrapFlags NWFlags =
+ inBoundsGEPFlags(*this, Ptr, ElemTy, SignedIndices, IsSubtraction);
Value *GEPVal = Builder.CreateGEP(ElemTy, Ptr, IdxList, Name, NWFlags);
@@ -6554,10 +6571,9 @@ Address CodeGenFunction::EmitCheckedInBoundsGEP(
bool SignedIndices, bool IsSubtraction, SourceLocation Loc, CharUnits Align,
const Twine &Name) {
if (!SanOpts.has(SanitizerKind::PointerOverflow)) {
- llvm::GEPNoWrapFlags NWFlags = llvm::GEPNoWrapFlags::inBounds();
- if (!SignedIndices && !IsSubtraction)
- NWFlags |= llvm::GEPNoWrapFlags::noUnsignedWrap();
-
+ llvm::GEPNoWrapFlags NWFlags =
+ inBoundsGEPFlags(*this, Addr.getBasePointer(), Addr.getElementType(),
+ SignedIndices, IsSubtraction);
return Builder.CreateGEP(Addr, IdxList, elementType, Align, Name, NWFlags);
}
diff --git a/clang/test/OpenMP/target_map_section_inbounds.cpp b/clang/test/OpenMP/target_map_section_inbounds.cpp
new file mode 100644
index 0000000000000..f3f380ef597bc
--- /dev/null
+++ b/clang/test/OpenMP/target_map_section_inbounds.cpp
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - \
+// RUN: | FileCheck %s
+
+int Data[8][4];
+
+void section_2d(int P, int I) {
+#pragma omp target map(tofrom : Data[P : 1][0 : 4]) firstprivate(P, I)
+ {
+ Data[P][I] = 1;
+ }
+}
+
+// CHECK-LABEL: define {{.*}}@{{__omp_offloading_.+}}section_2d
+// CHECK-NOT: getelementptr inbounds{{.*}}[8 x [4 x i32]]
+// CHECK: getelementptr{{.*}}[8 x [4 x i32]]
+// CHECK: getelementptr inbounds{{.*}}[4 x i32]
>From c7ef2f7fb264666e4bdef5ff36cae04e13ca05e6 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Tue, 25 Aug 2026 17:10:37 -0500
Subject: [PATCH 2/3] Update clang/lib/CodeGen/CGExprScalar.cpp
Co-authored-by: Alexey Bataev <a.bataev at gmx.com>
---
clang/lib/CodeGen/CGExprScalar.cpp | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp
index 7d848feed4eda..538efafc4a44f 100644
--- a/clang/lib/CodeGen/CGExprScalar.cpp
+++ b/clang/lib/CodeGen/CGExprScalar.cpp
@@ -6441,9 +6441,8 @@ static llvm::GEPNoWrapFlags
inBoundsGEPFlags(const CodeGenFunction &CGF, const llvm::Value *Ptr,
llvm::Type *SrcTy, bool SignedIndices, bool IsSubtraction) {
llvm::GEPNoWrapFlags NW;
- auto *AT = dyn_cast<llvm::ArrayType>(SrcTy);
- bool SectionGEP = CGF.getLangOpts().OpenMPIsTargetDevice && AT &&
- AT->getElementType()->isArrayTy() &&
+ bool SectionGEP = CGF.getLangOpts().OpenMPIsTargetDevice &&
+ isa<llvm::ArrayType>(SrcTy) &&
!isa<llvm::AllocaInst>(llvm::getUnderlyingObject(Ptr));
if (!SectionGEP)
NW = llvm::GEPNoWrapFlags::inBounds();
>From 97064bf77f7a9722244b5a981f85bc05e560978d Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 2 Oct 2026 15:36:00 -0500
Subject: [PATCH 3/3] Only drop inbounds for arrays of arrays, update tests
---
clang/lib/CodeGen/CGExprScalar.cpp | 3 +-
clang/test/OpenMP/nvptx_target_codegen.cpp | 4 +-
..._teams_distribute_parallel_for_codegen.cpp | 12 ++---
...s_distribute_parallel_for_simd_codegen.cpp | 4 +-
...x_target_teams_distribute_simd_codegen.cpp | 12 ++---
...vptx_target_teams_generic_loop_codegen.cpp | 12 ++---
.../OpenMP/target_map_section_inbounds.cpp | 4 ++
clang/test/OpenMP/target_parallel_codegen.cpp | 4 +-
.../OpenMP/target_parallel_debug_codegen.cpp | 50 +++++++++----------
.../OpenMP/target_parallel_for_codegen.cpp | 4 +-
.../target_parallel_for_debug_codegen.cpp | 50 +++++++++----------
.../target_parallel_for_simd_codegen.cpp | 8 +--
...target_parallel_generic_loop_codegen-3.cpp | 50 +++++++++----------
clang/test/OpenMP/target_teams_codegen.cpp | 4 +-
.../target_teams_distribute_codegen.cpp | 4 +-
.../target_teams_distribute_simd_codegen.cpp | 8 +--
16 files changed, 119 insertions(+), 114 deletions(-)
diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp
index fe436081265ee..0c101e40defdd 100644
--- a/clang/lib/CodeGen/CGExprScalar.cpp
+++ b/clang/lib/CodeGen/CGExprScalar.cpp
@@ -6468,7 +6468,8 @@ inBoundsGEPFlags(const CodeGenFunction &CGF, const llvm::Value *Ptr,
llvm::Type *SrcTy, bool SignedIndices, bool IsSubtraction) {
llvm::GEPNoWrapFlags NW;
bool SectionGEP = CGF.getLangOpts().OpenMPIsTargetDevice &&
- isa<llvm::ArrayType>(SrcTy) &&
+ SrcTy->isArrayTy() &&
+ SrcTy->getArrayElementType()->isArrayTy() &&
!isa<llvm::AllocaInst>(llvm::getUnderlyingObject(Ptr));
if (!SectionGEP)
NW = llvm::GEPNoWrapFlags::inBounds();
diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp
index 39ad99e7cfe2a..7b0cf4cef2141 100644
--- a/clang/test/OpenMP/nvptx_target_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -283,7 +283,7 @@ void unreachable_call() {
// CHECK1-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00
// CHECK1-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float
// CHECK1-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4
-// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i64 0, i64 2
// CHECK1-NEXT: [[TMP12:%.*]] = load double, ptr [[ARRAYIDX12]], align 8
// CHECK1-NEXT: [[ADD13:%.*]] = fadd double [[TMP12]], 1.000000e+00
@@ -669,7 +669,7 @@ void unreachable_call() {
// CHECK2-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00
// CHECK2-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float
// CHECK2-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4
-// CHECK2-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK2-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK2-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i32 0, i32 2
// CHECK2-NEXT: [[TMP12:%.*]] = load double, ptr [[ARRAYIDX12]], align 8
// CHECK2-NEXT: [[ADD13:%.*]] = fadd double [[TMP12]], 1.000000e+00
diff --git a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
index 2ec30f5ce8ae4..8ebe1fcad87ab 100644
--- a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
@@ -1033,7 +1033,7 @@ int bar(int n){
// CHECK1-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP14]]
// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP15]] to i64
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4
// CHECK1-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP16]] to i64
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -1341,7 +1341,7 @@ int bar(int n){
// CHECK1-NEXT: [[ADD33:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[I9]], align 4
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP23]] to i64
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[J10]], align 4
// CHECK1-NEXT: [[IDXPROM34:%.*]] = sext i32 [[TMP24]] to i64
// CHECK1-NEXT: [[ARRAYIDX35:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM34]]
@@ -2593,7 +2593,7 @@ int bar(int n){
// CHECK2-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP14]]
// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4
// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP15]] to i64
-// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4
// CHECK2-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP16]] to i64
// CHECK2-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -2896,7 +2896,7 @@ int bar(int n){
// CHECK2-NEXT: [[ADD29:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK2-NEXT: [[TMP23:%.*]] = load i32, ptr [[I9]], align 4
// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP23]] to i64
-// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[J10]], align 4
// CHECK2-NEXT: [[IDXPROM30:%.*]] = sext i32 [[TMP24]] to i64
// CHECK2-NEXT: [[ARRAYIDX31:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM30]]
@@ -4123,7 +4123,7 @@ int bar(int n){
// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[K]], align 4
// CHECK3-NEXT: [[ADD8:%.*]] = add nsw i32 [[ADD7]], [[TMP14]]
// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4
-// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP15]]
+// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP15]]
// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4
// CHECK3-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP16]]
// CHECK3-NEXT: store i32 [[ADD8]], ptr [[ARRAYIDX9]], align 4
@@ -4434,7 +4434,7 @@ int bar(int n){
// CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[J12]], align 4
// CHECK3-NEXT: [[ADD36:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[I11]], align 4
-// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP23]]
+// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP23]]
// CHECK3-NEXT: [[TMP24:%.*]] = load i32, ptr [[J12]], align 4
// CHECK3-NEXT: [[ARRAYIDX37:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP24]]
// CHECK3-NEXT: store i32 [[ADD36]], ptr [[ARRAYIDX37]], align 4
diff --git a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp
index 7a1ad6bab8d29..1dfb22fac6f3b 100644
--- a/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp
@@ -1091,7 +1091,7 @@ int bar(int n){
// CHECK1-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP14]]
// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP32]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP15]] to i64
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP32]]
// CHECK1-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP16]] to i64
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -2127,7 +2127,7 @@ int bar(int n){
// CHECK2-NEXT: [[TMP14:%.*]] = load i32, ptr [[K]], align 4, !llvm.access.group [[ACC_GRP32]]
// CHECK2-NEXT: [[ADD8:%.*]] = add nsw i32 [[ADD7]], [[TMP14]]
// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP32]]
-// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP15]]
+// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP15]]
// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP32]]
// CHECK2-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP16]]
// CHECK2-NEXT: store i32 [[ADD8]], ptr [[ARRAYIDX9]], align 4, !llvm.access.group [[ACC_GRP32]]
diff --git a/clang/test/OpenMP/nvptx_target_teams_distribute_simd_codegen.cpp b/clang/test/OpenMP/nvptx_target_teams_distribute_simd_codegen.cpp
index cacfc06b0e434..f306bf93acafb 100644
--- a/clang/test/OpenMP/nvptx_target_teams_distribute_simd_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_distribute_simd_codegen.cpp
@@ -618,7 +618,7 @@ int bar(int n){
// CHECK45-64-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP16]]
// CHECK45-64-NEXT: [[TMP17:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK45-64-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i64
-// CHECK45-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK45-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK45-64-NEXT: [[TMP18:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK45-64-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP18]] to i64
// CHECK45-64-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -1203,7 +1203,7 @@ int bar(int n){
// CHECK45-32-NEXT: [[TMP16:%.*]] = load i32, ptr [[K]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK45-32-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP16]]
// CHECK45-32-NEXT: [[TMP17:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP20]]
-// CHECK45-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
+// CHECK45-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
// CHECK45-32-NEXT: [[TMP18:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK45-32-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP18]]
// CHECK45-32-NEXT: store i32 [[ADD10]], ptr [[ARRAYIDX11]], align 4, !llvm.access.group [[ACC_GRP20]]
@@ -1787,7 +1787,7 @@ int bar(int n){
// CHECK45-32-EX-NEXT: [[TMP16:%.*]] = load i32, ptr [[K]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK45-32-EX-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP16]]
// CHECK45-32-EX-NEXT: [[TMP17:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP20]]
-// CHECK45-32-EX-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
+// CHECK45-32-EX-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
// CHECK45-32-EX-NEXT: [[TMP18:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK45-32-EX-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP18]]
// CHECK45-32-EX-NEXT: store i32 [[ADD10]], ptr [[ARRAYIDX11]], align 4, !llvm.access.group [[ACC_GRP20]]
@@ -2375,7 +2375,7 @@ int bar(int n){
// CHECK-64-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP16]]
// CHECK-64-NEXT: [[TMP17:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK-64-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP17]] to i64
-// CHECK-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK-64-NEXT: [[TMP18:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK-64-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP18]] to i64
// CHECK-64-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -2960,7 +2960,7 @@ int bar(int n){
// CHECK-32-NEXT: [[TMP16:%.*]] = load i32, ptr [[K]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK-32-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP16]]
// CHECK-32-NEXT: [[TMP17:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP20]]
-// CHECK-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
+// CHECK-32-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
// CHECK-32-NEXT: [[TMP18:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK-32-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP18]]
// CHECK-32-NEXT: store i32 [[ADD10]], ptr [[ARRAYIDX11]], align 4, !llvm.access.group [[ACC_GRP20]]
@@ -3544,7 +3544,7 @@ int bar(int n){
// CHECK-32-EX-NEXT: [[TMP16:%.*]] = load i32, ptr [[K]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK-32-EX-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP16]]
// CHECK-32-EX-NEXT: [[TMP17:%.*]] = load i32, ptr [[I]], align 4, !llvm.access.group [[ACC_GRP20]]
-// CHECK-32-EX-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
+// CHECK-32-EX-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP17]]
// CHECK-32-EX-NEXT: [[TMP18:%.*]] = load i32, ptr [[J]], align 4, !llvm.access.group [[ACC_GRP20]]
// CHECK-32-EX-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP18]]
// CHECK-32-EX-NEXT: store i32 [[ADD10]], ptr [[ARRAYIDX11]], align 4, !llvm.access.group [[ACC_GRP20]]
diff --git a/clang/test/OpenMP/nvptx_target_teams_generic_loop_codegen.cpp b/clang/test/OpenMP/nvptx_target_teams_generic_loop_codegen.cpp
index 5b508c61f4319..c97fa6b299688 100644
--- a/clang/test/OpenMP/nvptx_target_teams_generic_loop_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_teams_generic_loop_codegen.cpp
@@ -723,7 +723,7 @@ int bar(int n){
// CHECK1-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP14]]
// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP15]] to i64
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4
// CHECK1-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP16]] to i64
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -1031,7 +1031,7 @@ int bar(int n){
// CHECK1-NEXT: [[ADD33:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[I9]], align 4
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP23]] to i64
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[J10]], align 4
// CHECK1-NEXT: [[IDXPROM34:%.*]] = sext i32 [[TMP24]] to i64
// CHECK1-NEXT: [[ARRAYIDX35:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM34]]
@@ -1979,7 +1979,7 @@ int bar(int n){
// CHECK2-NEXT: [[ADD10:%.*]] = add nsw i32 [[ADD9]], [[TMP14]]
// CHECK2-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4
// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP15]] to i64
-// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK2-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4
// CHECK2-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP16]] to i64
// CHECK2-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM11]]
@@ -2282,7 +2282,7 @@ int bar(int n){
// CHECK2-NEXT: [[ADD29:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK2-NEXT: [[TMP23:%.*]] = load i32, ptr [[I9]], align 4
// CHECK2-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP23]] to i64
-// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
+// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]
// CHECK2-NEXT: [[TMP24:%.*]] = load i32, ptr [[J10]], align 4
// CHECK2-NEXT: [[IDXPROM30:%.*]] = sext i32 [[TMP24]] to i64
// CHECK2-NEXT: [[ARRAYIDX31:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i64 0, i64 [[IDXPROM30]]
@@ -3212,7 +3212,7 @@ int bar(int n){
// CHECK3-NEXT: [[TMP14:%.*]] = load i32, ptr [[K]], align 4
// CHECK3-NEXT: [[ADD8:%.*]] = add nsw i32 [[ADD7]], [[TMP14]]
// CHECK3-NEXT: [[TMP15:%.*]] = load i32, ptr [[I]], align 4
-// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP15]]
+// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP15]]
// CHECK3-NEXT: [[TMP16:%.*]] = load i32, ptr [[J]], align 4
// CHECK3-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP16]]
// CHECK3-NEXT: store i32 [[ADD8]], ptr [[ARRAYIDX9]], align 4
@@ -3523,7 +3523,7 @@ int bar(int n){
// CHECK3-NEXT: [[TMP22:%.*]] = load i32, ptr [[J12]], align 4
// CHECK3-NEXT: [[ADD36:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK3-NEXT: [[TMP23:%.*]] = load i32, ptr [[I11]], align 4
-// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP23]]
+// CHECK3-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP0]], i32 0, i32 [[TMP23]]
// CHECK3-NEXT: [[TMP24:%.*]] = load i32, ptr [[J12]], align 4
// CHECK3-NEXT: [[ARRAYIDX37:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX]], i32 0, i32 [[TMP24]]
// CHECK3-NEXT: store i32 [[ADD36]], ptr [[ARRAYIDX37]], align 4
diff --git a/clang/test/OpenMP/target_map_section_inbounds.cpp b/clang/test/OpenMP/target_map_section_inbounds.cpp
index f3f380ef597bc..2ac5862937f47 100644
--- a/clang/test/OpenMP/target_map_section_inbounds.cpp
+++ b/clang/test/OpenMP/target_map_section_inbounds.cpp
@@ -1,8 +1,12 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple amdgcn-amd-amdhsa \
// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
// RUN: -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - \
// RUN: | FileCheck %s
+// expected-no-diagnostics
+
int Data[8][4];
void section_2d(int P, int I) {
diff --git a/clang/test/OpenMP/target_parallel_codegen.cpp b/clang/test/OpenMP/target_parallel_codegen.cpp
index 8969b74ed3e24..10388c475e14d 100644
--- a/clang/test/OpenMP/target_parallel_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_codegen.cpp
@@ -2917,7 +2917,7 @@ int bar(int n){
// CHECK9-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00
// CHECK9-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float
// CHECK9-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4
-// CHECK9-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK9-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK9-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i64 0, i64 2
// CHECK9-NEXT: [[TMP11:%.*]] = load double, ptr [[ARRAYIDX12]], align 8
// CHECK9-NEXT: [[ADD13:%.*]] = fadd double [[TMP11]], 1.000000e+00
@@ -3387,7 +3387,7 @@ int bar(int n){
// CHECK11-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00
// CHECK11-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float
// CHECK11-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4
-// CHECK11-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK11-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK11-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i32 0, i32 2
// CHECK11-NEXT: [[TMP11:%.*]] = load double, ptr [[ARRAYIDX12]], align 8
// CHECK11-NEXT: [[ADD13:%.*]] = fadd double [[TMP11]], 1.000000e+00
diff --git a/clang/test/OpenMP/target_parallel_debug_codegen.cpp b/clang/test/OpenMP/target_parallel_debug_codegen.cpp
index cbd7e1346a580..78793222ed7a1 100644
--- a/clang/test/OpenMP/target_parallel_debug_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_debug_codegen.cpp
@@ -197,8 +197,8 @@ int main() {
// CHECK1-NEXT: #dbg_declare(ptr [[B3]], [[META88:![0-9]+]], !DIExpression(), [[META77]])
// CHECK1-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[B3]], ptr align 4 [[TMP4]], i64 400, i1 false), !dbg [[DBG87]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META89:![0-9]+]], !DIExpression(), [[META92:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG93:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG93]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG93:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX4:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG93]]
// CHECK1-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX4]], i64 0, i64 1, !dbg [[DBG93]]
// CHECK1-NEXT: store ptr [[ARRAYIDX5]], ptr [[F]], align 8, !dbg [[META92]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META94:![0-9]+]], !DIExpression(), [[META95:![0-9]+]])
@@ -215,14 +215,14 @@ int main() {
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP8]] to i64, !dbg [[DBG102]]
// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX8]], i64 0, i64 [[IDXPROM]], !dbg [[DBG102]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX9]], align 4, !dbg [[DBG104:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG105:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX10]], i64 0, i64 0, !dbg [[DBG105]]
+// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG105:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX10]], i64 0, i64 0, !dbg [[DBG105]]
// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG106:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM12:%.*]] = sext i32 [[TMP9]] to i64, !dbg [[DBG105]]
// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX11]], i64 0, i64 [[IDXPROM12]], !dbg [[DBG105]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX13]], align 4, !dbg [[DBG107:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG108:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX14]], i64 0, i64 0, !dbg [[DBG108]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG108:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX14]], i64 0, i64 0, !dbg [[DBG108]]
// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG109:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM16:%.*]] = sext i32 [[TMP10]] to i64, !dbg [[DBG108]]
// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX15]], i64 0, i64 [[IDXPROM16]], !dbg [[DBG108]]
@@ -416,37 +416,37 @@ int main() {
// CHECK1-NEXT: store ptr [[TMP7]], ptr [[_TMP2]], align 8, !dbg [[DBG171]]
// CHECK1-NEXT: [[TMP8:%.*]] = load ptr, ptr [[_TMP2]], align 8, !dbg [[DBG171]], !nonnull [[META37]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META172:![0-9]+]], !DIExpression(), [[META174:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG175:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG175]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG175:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX3:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG175]]
// CHECK1-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX3]], i64 0, i64 1, !dbg [[DBG175]]
// CHECK1-NEXT: store ptr [[ARRAYIDX4]], ptr [[F]], align 8, !dbg [[META174]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META176:![0-9]+]], !DIExpression(), [[META177:![0-9]+]])
// CHECK1-NEXT: store ptr [[A_ADDR]], ptr [[G]], align 8, !dbg [[META177]]
// CHECK1-NEXT: #dbg_declare(ptr [[H]], [[META178:![0-9]+]], !DIExpression(), [[META179:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 1, !dbg [[DBG180:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX5:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 1, !dbg [[DBG180:![0-9]+]]
// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX5]], i64 0, i64 1, !dbg [[DBG180]]
// CHECK1-NEXT: store ptr [[ARRAYIDX6]], ptr [[H]], align 8, !dbg [[META179]]
// CHECK1-NEXT: #dbg_declare(ptr [[D]], [[META181:![0-9]+]], !DIExpression(), [[META182:![0-9]+]])
// CHECK1-NEXT: store i32 15, ptr [[D]], align 4, !dbg [[META182]]
// CHECK1-NEXT: store i32 5, ptr [[A_ADDR]], align 4, !dbg [[DBG183:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG184:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG184:![0-9]+]]
// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG185:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP9]] to i64, !dbg [[DBG184]]
// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX7]], i64 0, i64 [[IDXPROM]], !dbg [[DBG184]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX8]], align 4, !dbg [[DBG186:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG187:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX9]], i64 0, i64 0, !dbg [[DBG187]]
+// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG187:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX9]], i64 0, i64 0, !dbg [[DBG187]]
// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG188:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM11:%.*]] = sext i32 [[TMP10]] to i64, !dbg [[DBG187]]
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX10]], i64 0, i64 [[IDXPROM11]], !dbg [[DBG187]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX12]], align 4, !dbg [[DBG189:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG190:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG190]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG190:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG190]]
// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG191:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM15:%.*]] = sext i32 [[TMP11]] to i64, !dbg [[DBG190]]
// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX14]], i64 0, i64 [[IDXPROM15]], !dbg [[DBG190]]
// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[ARRAYIDX16]], align 4, !dbg [[DBG190]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG192:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG192:![0-9]+]]
// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG193:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM18:%.*]] = sext i32 [[TMP13]] to i64, !dbg [[DBG192]]
// CHECK1-NEXT: [[ARRAYIDX19:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX17]], i64 0, i64 [[IDXPROM18]], !dbg [[DBG192]]
@@ -635,42 +635,42 @@ int main() {
// CHECK1-NEXT: store ptr [[TMP10]], ptr [[_TMP3]], align 8, !dbg [[DBG252]]
// CHECK1-NEXT: [[TMP11:%.*]] = load ptr, ptr [[_TMP3]], align 8, !dbg [[DBG252]], !nonnull [[META37]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META253:![0-9]+]], !DIExpression(), [[META255:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG256:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG256]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG256:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX4:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG256]]
// CHECK1-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX4]], i64 0, i64 1, !dbg [[DBG256]]
// CHECK1-NEXT: store ptr [[ARRAYIDX5]], ptr [[F]], align 8, !dbg [[META255]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META257:![0-9]+]], !DIExpression(), [[META258:![0-9]+]])
// CHECK1-NEXT: store ptr [[TMP5]], ptr [[G]], align 8, !dbg [[META258]]
// CHECK1-NEXT: #dbg_declare(ptr [[H]], [[META259:![0-9]+]], !DIExpression(), [[META260:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 1, !dbg [[DBG261:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 1, !dbg [[DBG261:![0-9]+]]
// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX6]], i64 0, i64 1, !dbg [[DBG261]]
// CHECK1-NEXT: store ptr [[ARRAYIDX7]], ptr [[H]], align 8, !dbg [[META260]]
// CHECK1-NEXT: #dbg_declare(ptr [[D]], [[META262:![0-9]+]], !DIExpression(), [[META263:![0-9]+]])
// CHECK1-NEXT: store i32 15, ptr [[D]], align 4, !dbg [[META263]]
// CHECK1-NEXT: store i32 5, ptr [[TMP5]], align 4, !dbg [[DBG264:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG265:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG265:![0-9]+]]
// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG266:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP12]] to i64, !dbg [[DBG265]]
// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX8]], i64 0, i64 [[IDXPROM]], !dbg [[DBG265]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX9]], align 4, !dbg [[DBG267:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG268:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX10]], i64 0, i64 0, !dbg [[DBG268]]
+// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG268:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX10]], i64 0, i64 0, !dbg [[DBG268]]
// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG269:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM12:%.*]] = sext i32 [[TMP13]] to i64, !dbg [[DBG268]]
// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX11]], i64 0, i64 [[IDXPROM12]], !dbg [[DBG268]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX13]], align 4, !dbg [[DBG270:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG271:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX14]], i64 0, i64 0, !dbg [[DBG271]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG271:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX14]], i64 0, i64 0, !dbg [[DBG271]]
// CHECK1-NEXT: [[TMP14:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG272:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM16:%.*]] = sext i32 [[TMP14]] to i64, !dbg [[DBG271]]
// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX15]], i64 0, i64 [[IDXPROM16]], !dbg [[DBG271]]
// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[ARRAYIDX17]], align 4, !dbg [[DBG271]]
-// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG273:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG273:![0-9]+]]
// CHECK1-NEXT: [[TMP16:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG274:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM19:%.*]] = sext i32 [[TMP16]] to i64, !dbg [[DBG273]]
// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX18]], i64 0, i64 [[IDXPROM19]], !dbg [[DBG273]]
// CHECK1-NEXT: store i32 [[TMP15]], ptr [[ARRAYIDX20]], align 4, !dbg [[DBG275:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX21:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG276:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX21:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG276:![0-9]+]]
// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG277:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM22:%.*]] = sext i32 [[TMP17]] to i64, !dbg [[DBG276]]
// CHECK1-NEXT: [[ARRAYIDX23:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX21]], i64 0, i64 [[IDXPROM22]], !dbg [[DBG276]]
diff --git a/clang/test/OpenMP/target_parallel_for_codegen.cpp b/clang/test/OpenMP/target_parallel_for_codegen.cpp
index 6c337e4dad399..2d94ee297d564 100644
--- a/clang/test/OpenMP/target_parallel_for_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_codegen.cpp
@@ -4275,7 +4275,7 @@ int bar(int n){
// CHECK9-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK9-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK9-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4
-// CHECK9-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK9-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK9-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i64 0, i64 2
// CHECK9-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8
// CHECK9-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
@@ -5135,7 +5135,7 @@ int bar(int n){
// CHECK11-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK11-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK11-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4
-// CHECK11-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK11-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK11-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i32 0, i32 2
// CHECK11-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8
// CHECK11-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
diff --git a/clang/test/OpenMP/target_parallel_for_debug_codegen.cpp b/clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
index 95c54d9ac3c17..dafb37bec13e4 100644
--- a/clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
@@ -250,8 +250,8 @@ int main() {
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]], !dbg [[DBG93]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !dbg [[DBG93]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META94:![0-9]+]], !DIExpression(), [[META97:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG98:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG98]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG98:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG98]]
// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX7]], i64 0, i64 1, !dbg [[DBG98]]
// CHECK1-NEXT: store ptr [[ARRAYIDX8]], ptr [[F]], align 8, !dbg [[META97]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META99:![0-9]+]], !DIExpression(), [[META100:![0-9]+]])
@@ -268,14 +268,14 @@ int main() {
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP18]] to i64, !dbg [[DBG107]]
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX11]], i64 0, i64 [[IDXPROM]], !dbg [[DBG107]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX12]], align 4, !dbg [[DBG109:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG110:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG110]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG110:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG110]]
// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG111:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM15:%.*]] = sext i32 [[TMP19]] to i64, !dbg [[DBG110]]
// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX14]], i64 0, i64 [[IDXPROM15]], !dbg [[DBG110]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX16]], align 4, !dbg [[DBG112:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG113:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG113]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG113:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG113]]
// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG114:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM19:%.*]] = sext i32 [[TMP20]] to i64, !dbg [[DBG113]]
// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX18]], i64 0, i64 [[IDXPROM19]], !dbg [[DBG113]]
@@ -542,37 +542,37 @@ int main() {
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]], !dbg [[DBG191]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !dbg [[DBG191]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META192:![0-9]+]], !DIExpression(), [[META194:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG195:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG195]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG195:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG195]]
// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX6]], i64 0, i64 1, !dbg [[DBG195]]
// CHECK1-NEXT: store ptr [[ARRAYIDX7]], ptr [[F]], align 8, !dbg [[META194]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META196:![0-9]+]], !DIExpression(), [[META197:![0-9]+]])
// CHECK1-NEXT: store ptr [[A_ADDR]], ptr [[G]], align 8, !dbg [[META197]]
// CHECK1-NEXT: #dbg_declare(ptr [[H]], [[META198:![0-9]+]], !DIExpression(), [[META199:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 1, !dbg [[DBG200:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 1, !dbg [[DBG200:![0-9]+]]
// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX8]], i64 0, i64 1, !dbg [[DBG200]]
// CHECK1-NEXT: store ptr [[ARRAYIDX9]], ptr [[H]], align 8, !dbg [[META199]]
// CHECK1-NEXT: #dbg_declare(ptr [[D]], [[META201:![0-9]+]], !DIExpression(), [[META202:![0-9]+]])
// CHECK1-NEXT: store i32 15, ptr [[D]], align 4, !dbg [[META202]]
// CHECK1-NEXT: store i32 5, ptr [[A_ADDR]], align 4, !dbg [[DBG203:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG204:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG204:![0-9]+]]
// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG205:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i64, !dbg [[DBG204]]
// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX10]], i64 0, i64 [[IDXPROM]], !dbg [[DBG204]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX11]], align 4, !dbg [[DBG206:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG207:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX12]], i64 0, i64 0, !dbg [[DBG207]]
+// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG207:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX12]], i64 0, i64 0, !dbg [[DBG207]]
// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG208:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM14:%.*]] = sext i32 [[TMP20]] to i64, !dbg [[DBG207]]
// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX13]], i64 0, i64 [[IDXPROM14]], !dbg [[DBG207]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX15]], align 4, !dbg [[DBG209:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG210:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX16]], i64 0, i64 0, !dbg [[DBG210]]
+// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG210:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX16]], i64 0, i64 0, !dbg [[DBG210]]
// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG211:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM18:%.*]] = sext i32 [[TMP21]] to i64, !dbg [[DBG210]]
// CHECK1-NEXT: [[ARRAYIDX19:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX17]], i64 0, i64 [[IDXPROM18]], !dbg [[DBG210]]
// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[ARRAYIDX19]], align 4, !dbg [[DBG210]]
-// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG212:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG212:![0-9]+]]
// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG213:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM21:%.*]] = sext i32 [[TMP23]] to i64, !dbg [[DBG212]]
// CHECK1-NEXT: [[ARRAYIDX22:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX20]], i64 0, i64 [[IDXPROM21]], !dbg [[DBG212]]
@@ -834,42 +834,42 @@ int main() {
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]], !dbg [[DBG285]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !dbg [[DBG285]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META286:![0-9]+]], !DIExpression(), [[META288:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG289:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG289]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG289:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG289]]
// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX7]], i64 0, i64 1, !dbg [[DBG289]]
// CHECK1-NEXT: store ptr [[ARRAYIDX8]], ptr [[F]], align 8, !dbg [[META288]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META290:![0-9]+]], !DIExpression(), [[META291:![0-9]+]])
// CHECK1-NEXT: store ptr [[TMP5]], ptr [[G]], align 8, !dbg [[META291]]
// CHECK1-NEXT: #dbg_declare(ptr [[H]], [[META292:![0-9]+]], !DIExpression(), [[META293:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 1, !dbg [[DBG294:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 1, !dbg [[DBG294:![0-9]+]]
// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX9]], i64 0, i64 1, !dbg [[DBG294]]
// CHECK1-NEXT: store ptr [[ARRAYIDX10]], ptr [[H]], align 8, !dbg [[META293]]
// CHECK1-NEXT: #dbg_declare(ptr [[D]], [[META295:![0-9]+]], !DIExpression(), [[META296:![0-9]+]])
// CHECK1-NEXT: store i32 15, ptr [[D]], align 4, !dbg [[META296]]
// CHECK1-NEXT: store i32 5, ptr [[TMP5]], align 4, !dbg [[DBG297:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG298:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG298:![0-9]+]]
// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG299:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP22]] to i64, !dbg [[DBG298]]
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX11]], i64 0, i64 [[IDXPROM]], !dbg [[DBG298]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX12]], align 4, !dbg [[DBG300:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG301:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG301]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG301:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG301]]
// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG302:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM15:%.*]] = sext i32 [[TMP23]] to i64, !dbg [[DBG301]]
// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX14]], i64 0, i64 [[IDXPROM15]], !dbg [[DBG301]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX16]], align 4, !dbg [[DBG303:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG304:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG304]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG304:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG304]]
// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG305:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM19:%.*]] = sext i32 [[TMP24]] to i64, !dbg [[DBG304]]
// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX18]], i64 0, i64 [[IDXPROM19]], !dbg [[DBG304]]
// CHECK1-NEXT: [[TMP25:%.*]] = load i32, ptr [[ARRAYIDX20]], align 4, !dbg [[DBG304]]
-// CHECK1-NEXT: [[ARRAYIDX21:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG306:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX21:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG306:![0-9]+]]
// CHECK1-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG307:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM22:%.*]] = sext i32 [[TMP26]] to i64, !dbg [[DBG306]]
// CHECK1-NEXT: [[ARRAYIDX23:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX21]], i64 0, i64 [[IDXPROM22]], !dbg [[DBG306]]
// CHECK1-NEXT: store i32 [[TMP25]], ptr [[ARRAYIDX23]], align 4, !dbg [[DBG308:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX24:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG309:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX24:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG309:![0-9]+]]
// CHECK1-NEXT: [[TMP27:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG310:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM25:%.*]] = sext i32 [[TMP27]] to i64, !dbg [[DBG309]]
// CHECK1-NEXT: [[ARRAYIDX26:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX24]], i64 0, i64 [[IDXPROM25]], !dbg [[DBG309]]
diff --git a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp
index 6663071e23cd5..1ff8e1ce2e8c3 100644
--- a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp
@@ -9928,7 +9928,7 @@ int bar(int n){
// CHECK17-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK17-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK17-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP32]]
-// CHECK17-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK17-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK17-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i64 0, i64 2
// CHECK17-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8, !llvm.access.group [[ACC_GRP32]]
// CHECK17-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
@@ -10813,7 +10813,7 @@ int bar(int n){
// CHECK19-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK19-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK19-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP32]]
-// CHECK19-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK19-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK19-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i32 0, i32 2
// CHECK19-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8, !llvm.access.group [[ACC_GRP32]]
// CHECK19-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
@@ -11698,7 +11698,7 @@ int bar(int n){
// CHECK21-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK21-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK21-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP32]]
-// CHECK21-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK21-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK21-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i64 0, i64 2
// CHECK21-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8, !llvm.access.group [[ACC_GRP32]]
// CHECK21-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
@@ -12674,7 +12674,7 @@ int bar(int n){
// CHECK23-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK23-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK23-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP32]]
-// CHECK23-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK23-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK23-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i32 0, i32 2
// CHECK23-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8, !llvm.access.group [[ACC_GRP32]]
// CHECK23-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
diff --git a/clang/test/OpenMP/target_parallel_generic_loop_codegen-3.cpp b/clang/test/OpenMP/target_parallel_generic_loop_codegen-3.cpp
index 18206912f92a0..cf5bd7df38a5b 100644
--- a/clang/test/OpenMP/target_parallel_generic_loop_codegen-3.cpp
+++ b/clang/test/OpenMP/target_parallel_generic_loop_codegen-3.cpp
@@ -250,8 +250,8 @@ int main() {
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]], !dbg [[DBG93]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !dbg [[DBG93]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META94:![0-9]+]], !DIExpression(), [[META97:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG98:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG98]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG98:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG98]]
// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX7]], i64 0, i64 1, !dbg [[DBG98]]
// CHECK1-NEXT: store ptr [[ARRAYIDX8]], ptr [[F]], align 8, !dbg [[META97]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META99:![0-9]+]], !DIExpression(), [[META100:![0-9]+]])
@@ -268,14 +268,14 @@ int main() {
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP18]] to i64, !dbg [[DBG107]]
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX11]], i64 0, i64 [[IDXPROM]], !dbg [[DBG107]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX12]], align 4, !dbg [[DBG109:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG110:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG110]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG110:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG110]]
// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG111:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM15:%.*]] = sext i32 [[TMP19]] to i64, !dbg [[DBG110]]
// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX14]], i64 0, i64 [[IDXPROM15]], !dbg [[DBG110]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX16]], align 4, !dbg [[DBG112:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG113:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG113]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG113:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG113]]
// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG114:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM19:%.*]] = sext i32 [[TMP20]] to i64, !dbg [[DBG113]]
// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX18]], i64 0, i64 [[IDXPROM19]], !dbg [[DBG113]]
@@ -542,37 +542,37 @@ int main() {
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]], !dbg [[DBG191]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !dbg [[DBG191]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META192:![0-9]+]], !DIExpression(), [[META194:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG195:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG195]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG195:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX6:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG195]]
// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX6]], i64 0, i64 1, !dbg [[DBG195]]
// CHECK1-NEXT: store ptr [[ARRAYIDX7]], ptr [[F]], align 8, !dbg [[META194]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META196:![0-9]+]], !DIExpression(), [[META197:![0-9]+]])
// CHECK1-NEXT: store ptr [[A_ADDR]], ptr [[G]], align 8, !dbg [[META197]]
// CHECK1-NEXT: #dbg_declare(ptr [[H]], [[META198:![0-9]+]], !DIExpression(), [[META199:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 1, !dbg [[DBG200:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 1, !dbg [[DBG200:![0-9]+]]
// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX8]], i64 0, i64 1, !dbg [[DBG200]]
// CHECK1-NEXT: store ptr [[ARRAYIDX9]], ptr [[H]], align 8, !dbg [[META199]]
// CHECK1-NEXT: #dbg_declare(ptr [[D]], [[META201:![0-9]+]], !DIExpression(), [[META202:![0-9]+]])
// CHECK1-NEXT: store i32 15, ptr [[D]], align 4, !dbg [[META202]]
// CHECK1-NEXT: store i32 5, ptr [[A_ADDR]], align 4, !dbg [[DBG203:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG204:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG204:![0-9]+]]
// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG205:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP19]] to i64, !dbg [[DBG204]]
// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX10]], i64 0, i64 [[IDXPROM]], !dbg [[DBG204]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX11]], align 4, !dbg [[DBG206:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG207:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX12]], i64 0, i64 0, !dbg [[DBG207]]
+// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG207:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX12]], i64 0, i64 0, !dbg [[DBG207]]
// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG208:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM14:%.*]] = sext i32 [[TMP20]] to i64, !dbg [[DBG207]]
// CHECK1-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX13]], i64 0, i64 [[IDXPROM14]], !dbg [[DBG207]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX15]], align 4, !dbg [[DBG209:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG210:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX16]], i64 0, i64 0, !dbg [[DBG210]]
+// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG210:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX16]], i64 0, i64 0, !dbg [[DBG210]]
// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG211:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM18:%.*]] = sext i32 [[TMP21]] to i64, !dbg [[DBG210]]
// CHECK1-NEXT: [[ARRAYIDX19:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX17]], i64 0, i64 [[IDXPROM18]], !dbg [[DBG210]]
// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[ARRAYIDX19]], align 4, !dbg [[DBG210]]
-// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG212:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP5]], i64 0, i64 0, !dbg [[DBG212:![0-9]+]]
// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[A_ADDR]], align 4, !dbg [[DBG213:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM21:%.*]] = sext i32 [[TMP23]] to i64, !dbg [[DBG212]]
// CHECK1-NEXT: [[ARRAYIDX22:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX20]], i64 0, i64 [[IDXPROM21]], !dbg [[DBG212]]
@@ -834,42 +834,42 @@ int main() {
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]], !dbg [[DBG285]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[I]], align 4, !dbg [[DBG285]]
// CHECK1-NEXT: #dbg_declare(ptr [[F]], [[META286:![0-9]+]], !DIExpression(), [[META288:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG289:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG289]]
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 1, !dbg [[DBG289:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX7:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX]], i64 0, i64 1, !dbg [[DBG289]]
// CHECK1-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX7]], i64 0, i64 1, !dbg [[DBG289]]
// CHECK1-NEXT: store ptr [[ARRAYIDX8]], ptr [[F]], align 8, !dbg [[META288]]
// CHECK1-NEXT: #dbg_declare(ptr [[G]], [[META290:![0-9]+]], !DIExpression(), [[META291:![0-9]+]])
// CHECK1-NEXT: store ptr [[TMP5]], ptr [[G]], align 8, !dbg [[META291]]
// CHECK1-NEXT: #dbg_declare(ptr [[H]], [[META292:![0-9]+]], !DIExpression(), [[META293:![0-9]+]])
-// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 1, !dbg [[DBG294:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX9:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 1, !dbg [[DBG294:![0-9]+]]
// CHECK1-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX9]], i64 0, i64 1, !dbg [[DBG294]]
// CHECK1-NEXT: store ptr [[ARRAYIDX10]], ptr [[H]], align 8, !dbg [[META293]]
// CHECK1-NEXT: #dbg_declare(ptr [[D]], [[META295:![0-9]+]], !DIExpression(), [[META296:![0-9]+]])
// CHECK1-NEXT: store i32 15, ptr [[D]], align 4, !dbg [[META296]]
// CHECK1-NEXT: store i32 5, ptr [[TMP5]], align 4, !dbg [[DBG297:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG298:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG298:![0-9]+]]
// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG299:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP22]] to i64, !dbg [[DBG298]]
// CHECK1-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX11]], i64 0, i64 [[IDXPROM]], !dbg [[DBG298]]
// CHECK1-NEXT: store i32 10, ptr [[ARRAYIDX12]], align 4, !dbg [[DBG300:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG301:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG301]]
+// CHECK1-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG301:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX13]], i64 0, i64 0, !dbg [[DBG301]]
// CHECK1-NEXT: [[TMP23:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG302:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM15:%.*]] = sext i32 [[TMP23]] to i64, !dbg [[DBG301]]
// CHECK1-NEXT: [[ARRAYIDX16:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX14]], i64 0, i64 [[IDXPROM15]], !dbg [[DBG301]]
// CHECK1-NEXT: store i32 11, ptr [[ARRAYIDX16]], align 4, !dbg [[DBG303:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG304:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG304]]
+// CHECK1-NEXT: [[ARRAYIDX17:%.*]] = getelementptr [10 x [10 x [10 x i32]]], ptr [[TMP2]], i64 0, i64 0, !dbg [[DBG304:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX18:%.*]] = getelementptr [10 x [10 x i32]], ptr [[ARRAYIDX17]], i64 0, i64 0, !dbg [[DBG304]]
// CHECK1-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG305:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM19:%.*]] = sext i32 [[TMP24]] to i64, !dbg [[DBG304]]
// CHECK1-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX18]], i64 0, i64 [[IDXPROM19]], !dbg [[DBG304]]
// CHECK1-NEXT: [[TMP25:%.*]] = load i32, ptr [[ARRAYIDX20]], align 4, !dbg [[DBG304]]
-// CHECK1-NEXT: [[ARRAYIDX21:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG306:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX21:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG306:![0-9]+]]
// CHECK1-NEXT: [[TMP26:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG307:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM22:%.*]] = sext i32 [[TMP26]] to i64, !dbg [[DBG306]]
// CHECK1-NEXT: [[ARRAYIDX23:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX21]], i64 0, i64 [[IDXPROM22]], !dbg [[DBG306]]
// CHECK1-NEXT: store i32 [[TMP25]], ptr [[ARRAYIDX23]], align 4, !dbg [[DBG308:![0-9]+]]
-// CHECK1-NEXT: [[ARRAYIDX24:%.*]] = getelementptr inbounds [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG309:![0-9]+]]
+// CHECK1-NEXT: [[ARRAYIDX24:%.*]] = getelementptr [10 x [10 x i32]], ptr [[TMP8]], i64 0, i64 0, !dbg [[DBG309:![0-9]+]]
// CHECK1-NEXT: [[TMP27:%.*]] = load i32, ptr [[TMP5]], align 4, !dbg [[DBG310:![0-9]+]]
// CHECK1-NEXT: [[IDXPROM25:%.*]] = sext i32 [[TMP27]] to i64, !dbg [[DBG309]]
// CHECK1-NEXT: [[ARRAYIDX26:%.*]] = getelementptr inbounds [10 x i32], ptr [[ARRAYIDX24]], i64 0, i64 [[IDXPROM25]], !dbg [[DBG309]]
diff --git a/clang/test/OpenMP/target_teams_codegen.cpp b/clang/test/OpenMP/target_teams_codegen.cpp
index a63c5b432d6a2..68a61f7ec5f24 100644
--- a/clang/test/OpenMP/target_teams_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_codegen.cpp
@@ -4530,7 +4530,7 @@ int bar(int n){
// CHECK9-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00
// CHECK9-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float
// CHECK9-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4
-// CHECK9-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK9-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK9-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i64 0, i64 2
// CHECK9-NEXT: [[TMP11:%.*]] = load double, ptr [[ARRAYIDX12]], align 8
// CHECK9-NEXT: [[ADD13:%.*]] = fadd double [[TMP11]], 1.000000e+00
@@ -5322,7 +5322,7 @@ int bar(int n){
// CHECK11-NEXT: [[ADD9:%.*]] = fadd double [[CONV8]], 1.000000e+00
// CHECK11-NEXT: [[CONV10:%.*]] = fptrunc double [[ADD9]] to float
// CHECK11-NEXT: store float [[CONV10]], ptr [[ARRAYIDX7]], align 4
-// CHECK11-NEXT: [[ARRAYIDX11:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK11-NEXT: [[ARRAYIDX11:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK11-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX11]], i32 0, i32 2
// CHECK11-NEXT: [[TMP11:%.*]] = load double, ptr [[ARRAYIDX12]], align 8
// CHECK11-NEXT: [[ADD13:%.*]] = fadd double [[TMP11]], 1.000000e+00
diff --git a/clang/test/OpenMP/target_teams_distribute_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_codegen.cpp
index 233ddd53f489c..aeab804626b88 100644
--- a/clang/test/OpenMP/target_teams_distribute_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_codegen.cpp
@@ -4335,7 +4335,7 @@ int bar(int n){
// CHECK9-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK9-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK9-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP22]]
-// CHECK9-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK9-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK9-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i64 0, i64 2
// CHECK9-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8, !llvm.access.group [[ACC_GRP22]]
// CHECK9-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
@@ -5243,7 +5243,7 @@ int bar(int n){
// CHECK11-NEXT: [[ADD12:%.*]] = fadd double [[CONV11]], 1.000000e+00
// CHECK11-NEXT: [[CONV13:%.*]] = fptrunc double [[ADD12]] to float
// CHECK11-NEXT: store float [[CONV13]], ptr [[ARRAYIDX10]], align 4, !llvm.access.group [[ACC_GRP22]]
-// CHECK11-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK11-NEXT: [[ARRAYIDX14:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK11-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX14]], i32 0, i32 2
// CHECK11-NEXT: [[TMP22:%.*]] = load double, ptr [[ARRAYIDX15]], align 8, !llvm.access.group [[ACC_GRP22]]
// CHECK11-NEXT: [[ADD16:%.*]] = fadd double [[TMP22]], 1.000000e+00
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp
index 5c7c798512c55..75a9289d6f3cd 100644
--- a/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp
@@ -10304,7 +10304,7 @@ int bar(int n){
// CHECK17-NEXT: [[ADD11:%.*]] = fadd double [[CONV10]], 1.000000e+00
// CHECK17-NEXT: [[CONV12:%.*]] = fptrunc double [[ADD11]] to float
// CHECK17-NEXT: store float [[CONV12]], ptr [[ARRAYIDX9]], align 4, !llvm.access.group [[ACC_GRP33]]
-// CHECK17-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK17-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK17-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX13]], i64 0, i64 2
// CHECK17-NEXT: [[TMP19:%.*]] = load double, ptr [[ARRAYIDX14]], align 8, !llvm.access.group [[ACC_GRP33]]
// CHECK17-NEXT: [[ADD15:%.*]] = fadd double [[TMP19]], 1.000000e+00
@@ -11243,7 +11243,7 @@ int bar(int n){
// CHECK19-NEXT: [[ADD11:%.*]] = fadd double [[CONV10]], 1.000000e+00
// CHECK19-NEXT: [[CONV12:%.*]] = fptrunc double [[ADD11]] to float
// CHECK19-NEXT: store float [[CONV12]], ptr [[ARRAYIDX9]], align 4, !llvm.access.group [[ACC_GRP33]]
-// CHECK19-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK19-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK19-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX13]], i32 0, i32 2
// CHECK19-NEXT: [[TMP19:%.*]] = load double, ptr [[ARRAYIDX14]], align 8, !llvm.access.group [[ACC_GRP33]]
// CHECK19-NEXT: [[ADD15:%.*]] = fadd double [[TMP19]], 1.000000e+00
@@ -12182,7 +12182,7 @@ int bar(int n){
// CHECK21-NEXT: [[ADD11:%.*]] = fadd double [[CONV10]], 1.000000e+00
// CHECK21-NEXT: [[CONV12:%.*]] = fptrunc double [[ADD11]] to float
// CHECK21-NEXT: store float [[CONV12]], ptr [[ARRAYIDX9]], align 4, !llvm.access.group [[ACC_GRP33]]
-// CHECK21-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
+// CHECK21-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i64 0, i64 1
// CHECK21-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX13]], i64 0, i64 2
// CHECK21-NEXT: [[TMP19:%.*]] = load double, ptr [[ARRAYIDX14]], align 8, !llvm.access.group [[ACC_GRP33]]
// CHECK21-NEXT: [[ADD15:%.*]] = fadd double [[TMP19]], 1.000000e+00
@@ -13178,7 +13178,7 @@ int bar(int n){
// CHECK23-NEXT: [[ADD11:%.*]] = fadd double [[CONV10]], 1.000000e+00
// CHECK23-NEXT: [[CONV12:%.*]] = fptrunc double [[ADD11]] to float
// CHECK23-NEXT: store float [[CONV12]], ptr [[ARRAYIDX9]], align 4, !llvm.access.group [[ACC_GRP33]]
-// CHECK23-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
+// CHECK23-NEXT: [[ARRAYIDX13:%.*]] = getelementptr [5 x [10 x double]], ptr [[TMP3]], i32 0, i32 1
// CHECK23-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds [10 x double], ptr [[ARRAYIDX13]], i32 0, i32 2
// CHECK23-NEXT: [[TMP19:%.*]] = load double, ptr [[ARRAYIDX14]], align 8, !llvm.access.group [[ACC_GRP33]]
// CHECK23-NEXT: [[ADD15:%.*]] = fadd double [[TMP19]], 1.000000e+00
More information about the cfe-commits
mailing list