[clang] [CIR] Add inbounds and nuw flags to cir.get_element (PR #234774)
Adam Smith via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 22 14:50:20 PDT 2026
https://github.com/adams381 created https://github.com/llvm/llvm-project/pull/234774
CIR lowered array subscripts to a getelementptr with no no-wrap flags, where classic codegen marks them inbounds, and also nuw when the index is unsigned.
cir.get_element now has optional inbounds and nuw flags, and CIRGen sets them where classic does: inbounds unless -fwrapv-pointer is given, nuw for an unsigned index expression, and no inbounds for an OpenMP device subscript of an array of arrays whose base is not an alloca. The flags live on the op because both depend on language options and on the C index type, which the lowering cannot see. LowerToLLVM maps them to the GEP's no-wrap flags.
Assisted-by: Cursor / Claude Opus 5.5
>From 7398f6cb58c889edb5ea458f069f4ae1c84df7ef Mon Sep 17 00:00:00 2001
From: Adam Smith <adams at nvidia.com>
Date: Thu, 22 Oct 2026 14:42:26 -0700
Subject: [PATCH] [CIR] Add inbounds and nuw flags to cir.get_element
CIR lowered array subscripts to a getelementptr with no no-wrap flags,
where classic codegen marks them inbounds, and also nuw when the index
is unsigned.
cir.get_element now has optional inbounds and nuw flags, and CIRGen sets
them where classic does: inbounds unless -fwrapv-pointer is given, nuw
for an unsigned index expression, and no inbounds for an OpenMP device
subscript of an array of arrays whose base is not an alloca. The flags
live on the op because both depend on language options and on the C
index type, which the lowering cannot see. LowerToLLVM maps them to the
GEP's no-wrap flags.
Assisted-by: Cursor / Claude Opus 5.5
---
.../CIR/Dialect/Builder/CIRBaseBuilder.h | 13 +
clang/include/clang/CIR/Dialect/IR/CIROps.td | 36 ++-
clang/lib/CIR/CodeGen/CIRGenBuilder.cpp | 9 +-
clang/lib/CIR/CodeGen/CIRGenBuilder.h | 18 +-
clang/lib/CIR/CodeGen/CIRGenBuiltinX86.cpp | 5 +-
clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 82 +++++--
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 9 +-
.../CIR/CodeGen/amdgpu-array-addrspace.cpp | 4 +-
.../CIR/CodeGen/array-init-loop-exprs.cpp | 16 +-
.../CIR/CodeGen/array-subscript-inbounds.c | 232 ++++++++++++++++++
.../CIR/CodeGen/array-subscript-order.cpp | 4 +-
clang/test/CIR/CodeGen/array.cpp | 34 +--
clang/test/CIR/CodeGen/binassign.c | 4 +-
clang/test/CIR/CodeGen/bitint-record-layout.c | 6 +-
clang/test/CIR/CodeGen/bitint-wide.c | 12 +-
clang/test/CIR/CodeGen/builtin-cpu-supports.c | 12 +-
clang/test/CIR/CodeGen/complex.cpp | 4 +-
clang/test/CIR/CodeGen/const-label-addr.c | 4 +-
clang/test/CIR/CodeGen/empty-union.cpp | 2 +-
clang/test/CIR/CodeGen/no-odr-use.cpp | 20 +-
clang/test/CIR/CodeGen/union.c | 8 +-
.../CIR/CodeGenBuiltins/builtin-constant-p.c | 2 +-
.../CodeGenHIP/string-literal-addrspace.hip | 12 +-
clang/test/CIR/CodeGenOpenACC/combined-copy.c | 2 +-
.../CodeGenOpenCL/string-literal-addrspace.cl | 12 +-
.../CodeGenOpenMP/device-array-subscript.c | 104 ++++++++
clang/test/CIR/IR/array.cir | 18 ++
clang/test/CIR/Lowering/array.cpp | 8 +-
clang/test/CIR/Lowering/get-element.cir | 40 +++
29 files changed, 616 insertions(+), 116 deletions(-)
create mode 100644 clang/test/CIR/CodeGen/array-subscript-inbounds.c
create mode 100644 clang/test/CIR/CodeGenOpenMP/device-array-subscript.c
create mode 100644 clang/test/CIR/Lowering/get-element.cir
diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
index 9560eae0dc1153..87ec07963e26d8 100644
--- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
+++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
@@ -16,6 +16,7 @@
#include "clang/CIR/Dialect/IR/CIRDialect.h"
#include "clang/CIR/Dialect/IR/CIRTypes.h"
#include "clang/CIR/MissingFeatures.h"
+#include "llvm/ADT/BitmaskEnum.h"
#include "llvm/ADT/STLForwardCompat.h"
#include "llvm/IR/FPEnv.h"
#include "llvm/Support/ErrorHandling.h"
@@ -62,6 +63,18 @@ constexpr bool testFlag(OverflowBehavior ob, OverflowBehavior flag) {
return (ob & flag) != OverflowBehavior::None;
}
+LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
+
+/// No-wrap flags of a pointer offset computation: the subset of
+/// llvm::GEPNoWrapFlags that `cir.get_element` carries as its `in_bounds` and
+/// `no_unsigned_wrap` properties.
+enum class PtrOffsetFlags {
+ None = 0,
+ InBounds = 1 << 0,
+ NoUnsignedWrap = 1 << 1,
+ LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/NoUnsignedWrap)
+};
+
class CIRBaseBuilderTy : public mlir::OpBuilder {
public:
diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td
index 4dde0a667ac5a1..d3f552a2dd157a 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIROps.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td
@@ -4225,18 +4225,42 @@ def CIR_GetElementOp : CIR_Op<"get_element", [
It expects a pointer to the `base` array and the `index` of the element.
The result pointer preserves the address space of the base pointer.
+ The optional `inbounds` and `nuw` (no unsigned wrap) flags have the
+ semantics of the same flags on LLVM's `getelementptr`. With `inbounds`,
+ unless `index` is zero, `base` and the result both point into the same
+ allocated object, or to its end. `inbounds` also guarantees that scaling
+ `index` by the element size does not overflow as a signed value, and that
+ adding that signed offset to the unsigned address of `base` does not wrap.
+ With `nuw`, the scaling and the addition do not wrap when `index` and the
+ address are both treated as unsigned. The result is poison if the
+ guarantee of a present flag does not hold.
+
Example:
```
- // Suppose we have a array.
+ // Suppose we have an array.
!s32i = !cir.int<s, 32>
+ !u32i = !cir.int<u, 32>
!arr_ty = !cir.array<!s32i x 4>
// Get the address of the element at index 1.
- %elem_1 = cir.get_element %0[1 : !s32i] : !cir.ptr<!array_ty> -> !cir.ptr<!s32i>
+ %elem_1 = cir.get_element %0[1 : !s32i] :
+ !cir.ptr<!arr_ty> -> !cir.ptr<!s32i>
// Get the address of the element at index %i.
%i = ...
- %elem_i = cir.get_element %0[%i : !s32i] : !cir.ptr<!array_ty> -> !cir.ptr<!s32i>
+ %elem_i = cir.get_element %0[%i : !s32i] :
+ !cir.ptr<!arr_ty> -> !cir.ptr<!s32i>
+
+ // The same, when %0 and the result are known to be in bounds of the same
+ // allocated object.
+ %elem_ib = cir.get_element inbounds %0[%i : !s32i] :
+ !cir.ptr<!arr_ty> -> !cir.ptr<!s32i>
+
+ // The same, when scaling %u and adding it to %0 are also known not to
+ // wrap as unsigned arithmetic.
+ %u = ...
+ %elem_ibnuw = cir.get_element inbounds nuw %0[%u : !u32i] :
+ !cir.ptr<!arr_ty> -> !cir.ptr<!s32i>
// With address space (e.g., GPU private memory):
%elem_gpu = cir.get_element %gpu_arr[%i : !s32i] :
@@ -4248,11 +4272,15 @@ def CIR_GetElementOp : CIR_Op<"get_element", [
let arguments = (ins
Arg<CIR_PtrToArray, "the base address of the array ">:$base,
- Arg<CIR_AnyFundamentalIntType, "the index of the element">:$index
+ Arg<CIR_AnyFundamentalIntType, "the index of the element">:$index,
+ UnitProp:$in_bounds,
+ UnitProp:$no_unsigned_wrap
);
let results = (outs CIR_PointerType : $result);
let assemblyFormat = [{
+ (`inbounds` $in_bounds^)?
+ (`nuw` $no_unsigned_wrap^)?
$base`[` $index `:` type($index) `]` attr-dict
`:` qualified(type($base)) `->` qualified(type($result))
}];
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.cpp b/clang/lib/CIR/CodeGen/CIRGenBuilder.cpp
index 4c39cf2d872528..f74c392c7cbe8b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuilder.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.cpp
@@ -36,7 +36,8 @@ mlir::Value CIRGenBuilderTy::getArrayElement(mlir::Location arrayLocBegin,
mlir::Location arrayLocEnd,
mlir::Value arrayPtr,
mlir::Type eltTy, mlir::Value idx,
- bool shouldDecay) {
+ bool shouldDecay,
+ cir::PtrOffsetFlags flags) {
auto arrayPtrTy = mlir::dyn_cast<cir::PointerType>(arrayPtr.getType());
assert(arrayPtrTy && "expected pointer type");
// If the array pointer is not decayed, emit a GetElementOp.
@@ -49,8 +50,10 @@ mlir::Value CIRGenBuilderTy::getArrayElement(mlir::Location arrayLocBegin,
if (shouldDecay && arrayTy && arrayTy == eltTy) {
auto eltPtrTy =
getPointerTo(arrayTy.getElementType(), arrayPtrTy.getAddrSpace());
- return cir::GetElementOp::create(*this, arrayLocEnd, eltPtrTy, arrayPtr,
- idx);
+ return cir::GetElementOp::create(
+ *this, arrayLocEnd, eltPtrTy, arrayPtr, idx,
+ any(flags & cir::PtrOffsetFlags::InBounds),
+ any(flags & cir::PtrOffsetFlags::NoUnsignedWrap));
}
// If we don't have sufficient type information, emit a PtrStrideOp.
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h
index b9cc67e394accf..af8cbfacd993a1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuilder.h
+++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.h
@@ -653,13 +653,17 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy {
memberPtr);
}
- /// Create a cir.ptr_stride operation to get access to an array element.
- /// \p idx is the index of the element to access, \p shouldDecay is true if
- /// the result should decay to a pointer to the element type.
- mlir::Value getArrayElement(mlir::Location arrayLocBegin,
- mlir::Location arrayLocEnd, mlir::Value arrayPtr,
- mlir::Type eltTy, mlir::Value idx,
- bool shouldDecay);
+ /// Create a cir.get_element or cir.ptr_stride operation to get access to an
+ /// array element. \p idx is the index of the element to access,
+ /// \p shouldDecay is true if the result should decay to a pointer to the
+ /// element type. \p flags are set on the cir.get_element emitted when
+ /// \p shouldDecay is true and \p arrayPtr points to the array type
+ /// \p eltTy, and are dropped when a cir.ptr_stride is emitted.
+ mlir::Value
+ getArrayElement(mlir::Location arrayLocBegin, mlir::Location arrayLocEnd,
+ mlir::Value arrayPtr, mlir::Type eltTy, mlir::Value idx,
+ bool shouldDecay,
+ cir::PtrOffsetFlags flags = cir::PtrOffsetFlags::None);
/// Returns a decayed pointer to the first element of the array
/// pointed to by \p arrayPtr.
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinX86.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinX86.cpp
index c4bd6b687c8f57..3ba228d4d02d95 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinX86.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinX86.cpp
@@ -1131,7 +1131,8 @@ CIRGenFunction::emitX86CpuSupports(mlir::Location loc,
mlir::Value cpuFeatVal = builder.getArrayElement(
loc, loc, cpuFeatArrPtr, arrTy,
- /*index=*/builder.getUInt32(0, loc), /*shouldDecay=*/true);
+ /*idx=*/builder.getUInt32(0, loc), /*shouldDecay=*/true,
+ cir::PtrOffsetFlags::InBounds);
cir::LoadOp features = builder.createAlignedLoad(
loc, u32, cpuFeatVal, CharUnits::fromQuantity(4));
@@ -1151,7 +1152,7 @@ CIRGenFunction::emitX86CpuSupports(mlir::Location loc,
loc, loc, getCpuFeatures2,
cast<cir::PointerType>(getCpuFeatures2.getType()).getPointee(),
builder.getUInt32(/*index=*/i - 1, loc),
- /*shouldDecay=*/true);
+ /*shouldDecay=*/true, cir::PtrOffsetFlags::InBounds);
cir::LoadOp features = builder.createAlignedLoad(
loc, u32, getArrayElt, CharUnits::fromQuantity(4));
diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
index dfd6fa725e77fe..3d3a318f84d40a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
@@ -1384,24 +1384,73 @@ static CharUnits getArrayElementAlign(CharUnits arrayAlign, mlir::Value idx,
return arrayAlign.alignmentOfArrayElement(eltSize);
}
+/// Return true if \p ptr is derived from a cir.alloca through cir.get_element,
+/// cir.ptr_stride, or a cir.cast that is array_to_ptrdecay or for which
+/// isAllocaPreservingCast() holds, so the alloca is its underlying object.
+static bool isBasedOnAlloca(mlir::Value ptr) {
+ while (mlir::Operation *op = ptr.getDefiningOp()) {
+ if (mlir::isa<cir::AllocaOp>(op))
+ return true;
+ if (auto elementOp = mlir::dyn_cast<cir::GetElementOp>(op))
+ ptr = elementOp.getBase();
+ else if (auto strideOp = mlir::dyn_cast<cir::PtrStrideOp>(op))
+ ptr = strideOp.getBase();
+ else if (auto castOp = mlir::dyn_cast<cir::CastOp>(op);
+ castOp && (castOp.isAllocaPreservingCast() ||
+ castOp.getKind() == cir::CastKind::array_to_ptrdecay))
+ ptr = castOp.getSrc();
+ else
+ return false;
+ }
+ return false;
+}
+
+/// Return the no-wrap flags for an array subscript of \p ptr: none when
+/// pointer overflow is defined, otherwise inbounds unless \p ptr may come from
+/// an OpenMP section map, plus nuw unless \p signedIndices.
+static cir::PtrOffsetFlags getArraySubscriptFlags(CIRGenFunction &cgf,
+ mlir::Value ptr,
+ bool signedIndices) {
+ if (cgf.getLangOpts().PointerOverflowDefined)
+ return cir::PtrOffsetFlags::None;
+
+ // An OpenMP section map passes `section - host_offset`, which can point
+ // outside the mapped allocation, so a subscript of an array of arrays based
+ // on it cannot be marked inbounds. A base derived from a cir.alloca is
+ // never such a pointer.
+ auto arrayTy = mlir::dyn_cast<cir::ArrayType>(
+ mlir::cast<cir::PointerType>(ptr.getType()).getPointee());
+ bool isSectionSubscript =
+ cgf.getLangOpts().OpenMPIsTargetDevice && arrayTy &&
+ mlir::isa<cir::ArrayType>(arrayTy.getElementType()) &&
+ !isBasedOnAlloca(ptr);
+
+ cir::PtrOffsetFlags flags = cir::PtrOffsetFlags::None;
+ if (!isSectionSubscript)
+ flags |= cir::PtrOffsetFlags::InBounds;
+ if (!signedIndices)
+ flags |= cir::PtrOffsetFlags::NoUnsignedWrap;
+ return flags;
+}
+
static mlir::Value emitArraySubscriptPtr(CIRGenFunction &cgf,
mlir::Location beginLoc,
mlir::Location endLoc, mlir::Value ptr,
mlir::Type eltTy, mlir::Value idx,
- bool shouldDecay) {
+ bool signedIndices, bool shouldDecay) {
CIRGenModule &cgm = cgf.getCIRGenModule();
- // TODO(cir): LLVM codegen emits in bound gep check here, is there anything
- // that would enhance tracking this later in CIR?
assert(!cir::MissingFeatures::emitCheckedInBoundsGEP());
- return cgm.getBuilder().getArrayElement(beginLoc, endLoc, ptr, eltTy, idx,
- shouldDecay);
+ return cgm.getBuilder().getArrayElement(
+ beginLoc, endLoc, ptr, eltTy, idx, shouldDecay,
+ getArraySubscriptFlags(cgf, ptr, signedIndices));
}
static Address emitArraySubscriptPtr(CIRGenFunction &cgf,
mlir::Location beginLoc,
mlir::Location endLoc, Address addr,
QualType eltType, mlir::Value idx,
- mlir::Location loc, bool shouldDecay) {
+ mlir::Location loc, bool signedIndices,
+ bool shouldDecay) {
// Determine the element size of the statically-sized base. This is
// the thing that the indices are expressed in terms of.
@@ -1416,9 +1465,9 @@ static Address emitArraySubscriptPtr(CIRGenFunction &cgf,
getArrayElementAlign(addr.getAlignment(), idx, eltSize);
assert(!cir::MissingFeatures::preservedAccessIndexRegion());
- const mlir::Value eltPtr =
- emitArraySubscriptPtr(cgf, beginLoc, endLoc, addr.getPointer(),
- addr.getElementType(), idx, shouldDecay);
+ const mlir::Value eltPtr = emitArraySubscriptPtr(
+ cgf, beginLoc, endLoc, addr.getPointer(), addr.getElementType(), idx,
+ signedIndices, shouldDecay);
const mlir::Type elementType = cgf.convertTypeForMem(eltType);
return Address(eltPtr, elementType, eltAlign);
}
@@ -1435,6 +1484,8 @@ CIRGenFunction::emitArraySubscriptExpr(const clang::ArraySubscriptExpr *e) {
mlir::Value idxPre = (e->getLHS() == e->getIdx())
? emitScalarExpr(e->getIdx())
: mlir::Value();
+ const bool signedIndices =
+ e->getIdx()->getType()->isSignedIntegerOrEnumerationType();
auto emitIdxAfterBase = [&, idxPre](bool promote) -> mlir::Value {
mlir::Value idx = idxPre;
@@ -1449,10 +1500,7 @@ CIRGenFunction::emitArraySubscriptExpr(const clang::ArraySubscriptExpr *e) {
if (promote) {
// Choose the type we extend or truncate to based on the signedness of the
// index type.
- mlir::Type desiredIdxTy =
- e->getIdx()->getType()->isSignedIntegerOrEnumerationType()
- ? ptrDiffTy
- : uIntPtrTy;
+ mlir::Type desiredIdxTy = signedIndices ? ptrDiffTy : uIntPtrTy;
if (idx.getType() != desiredIdxTy) {
cir::CastKind kind = mlir::isa<cir::BoolType>(idx.getType())
@@ -1496,7 +1544,7 @@ CIRGenFunction::emitArraySubscriptExpr(const clang::ArraySubscriptExpr *e) {
addr = emitArraySubscriptPtr(*this, cgm.getLoc(e->getBeginLoc()),
cgm.getLoc(e->getEndLoc()), addr, e->getType(),
idx, cgm.getLoc(e->getExprLoc()),
- /*shouldDecay=*/false);
+ signedIndices, /*shouldDecay=*/false);
return makeAddrLValue(addr, elementType, lv.getBaseInfo());
}
@@ -1526,7 +1574,7 @@ CIRGenFunction::emitArraySubscriptExpr(const clang::ArraySubscriptExpr *e) {
addr = emitArraySubscriptPtr(*this, cgm.getLoc(e->getBeginLoc()),
cgm.getLoc(e->getEndLoc()), addr, e->getType(),
idx, cgm.getLoc(e->getExprLoc()),
- /*shouldDecay=*/false);
+ signedIndices, /*shouldDecay=*/false);
return makeAddrLValue(addr, vla->getElementType(), LValueBaseInfo());
}
@@ -1543,7 +1591,7 @@ CIRGenFunction::emitArraySubscriptExpr(const clang::ArraySubscriptExpr *e) {
const Address addr = emitArraySubscriptPtr(
*this, cgm.getLoc(array->getBeginLoc()), cgm.getLoc(array->getEndLoc()),
arrayLV.getAddress(), e->getType(), idx, cgm.getLoc(e->getExprLoc()),
- /*shouldDecay=*/true);
+ signedIndices, /*shouldDecay=*/true);
const LValue lv = LValue::makeAddr(addr, e->getType(), LValueBaseInfo());
@@ -1564,7 +1612,7 @@ CIRGenFunction::emitArraySubscriptExpr(const clang::ArraySubscriptExpr *e) {
// Propagate the alignment from the array itself to the result.
const Address addxr = emitArraySubscriptPtr(
*this, cgm.getLoc(e->getBeginLoc()), cgm.getLoc(e->getEndLoc()), ptrAddr,
- e->getType(), idx, cgm.getLoc(e->getExprLoc()),
+ e->getType(), idx, cgm.getLoc(e->getExprLoc()), signedIndices,
/*shouldDecay=*/false);
const LValue lv = LValue::makeAddr(addxr, e->getType(), eltBaseInfo);
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 719ec9ac184ff6..371b74ceb28660 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -2282,8 +2282,13 @@ mlir::LogicalResult CIRToLLVMGetElementOpLowering::matchAndRewrite(
// offset is always zero. The second offset tell us which member it
// will access.
std::array<mlir::LLVM::GEPArg, 2> offset{0, index};
- rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(op, llResultTy, llArrayTy,
- adaptor.getBase(), offset);
+ mlir::LLVM::GEPNoWrapFlags flags = mlir::LLVM::GEPNoWrapFlags::none;
+ if (op.getInBounds())
+ flags |= mlir::LLVM::GEPNoWrapFlags::inbounds;
+ if (op.getNoUnsignedWrap())
+ flags |= mlir::LLVM::GEPNoWrapFlags::nuw;
+ rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(
+ op, llResultTy, llArrayTy, adaptor.getBase(), offset, flags);
return mlir::success();
}
diff --git a/clang/test/CIR/CodeGen/amdgpu-array-addrspace.cpp b/clang/test/CIR/CodeGen/amdgpu-array-addrspace.cpp
index ba7889fab59341..080d3ba1e24966 100644
--- a/clang/test/CIR/CodeGen/amdgpu-array-addrspace.cpp
+++ b/clang/test/CIR/CodeGen/amdgpu-array-addrspace.cpp
@@ -49,11 +49,11 @@ void pass_global_array() {
// CIR-LABEL: cir.func{{.*}} @_Z18index_global_arrayi
// CIR: %[[ARR:.*]] = cir.get_global @globalArr : !cir.ptr<!cir.array<!s32i x 10>, target_address_space(1)>
// CIR-NEXT: %[[FLAT:.*]] = cir.cast address_space %[[ARR]] : !cir.ptr<!cir.array<!s32i x 10>, target_address_space(1)> -> !cir.ptr<!cir.array<!s32i x 10>>
-// CIR: %[[ELEM:.*]] = cir.get_element %[[FLAT]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CIR: %[[ELEM:.*]] = cir.get_element inbounds %[[FLAT]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
// CIR-NEXT: %{{.*}} = cir.load align(4) %[[ELEM]] : !cir.ptr<!s32i>, !s32i
// LLVM-LABEL: define{{.*}} i32 @_Z18index_global_arrayi
-// LLVM: %[[GEP:.*]] = getelementptr [10 x i32], ptr addrspacecast (ptr addrspace(1) @globalArr to ptr), i32 0, i64 %{{.*}}
+// LLVM: %[[GEP:.*]] = getelementptr inbounds [10 x i32], ptr addrspacecast (ptr addrspace(1) @globalArr to ptr), i32 0, i64 %{{.*}}
// LLVM-NEXT: %{{.*}} = load i32, ptr %[[GEP]], align 4
// OGCG-LABEL: define{{.*}} i32 @_Z18index_global_arrayi
diff --git a/clang/test/CIR/CodeGen/array-init-loop-exprs.cpp b/clang/test/CIR/CodeGen/array-init-loop-exprs.cpp
index f5907892fe53e3..083514951f2a53 100644
--- a/clang/test/CIR/CodeGen/array-init-loop-exprs.cpp
+++ b/clang/test/CIR/CodeGen/array-init-loop-exprs.cpp
@@ -65,7 +65,7 @@ struct HasMultiDimArray {
// CIR: %[[ITR_LOAD:.*]] = cir.load {{.*}}%[[ITR_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_NonTrivial>>, !cir.ptr<!rec_NonTrivial>
// CIR: %[[IDX:.*]] = cir.ptr_diff %[[ITR_LOAD]], %[[THIS_ARR_DECAY]] : !cir.ptr<!rec_NonTrivial> -> !s64i
// CIR: %[[IDX_CAST:.*]] = cir.cast integral %[[IDX]] : !s64i -> !u64i
-// CIR: %[[RHS_ELT:.*]] = cir.get_element %[[RHS_ARR]][%[[IDX_CAST]] : !u64i] : !cir.ptr<!cir.array<!rec_NonTrivial x 3>> -> !cir.ptr<!rec_NonTrivial>
+// CIR: %[[RHS_ELT:.*]] = cir.get_element inbounds nuw %[[RHS_ARR]][%[[IDX_CAST]] : !u64i] : !cir.ptr<!cir.array<!rec_NonTrivial x 3>> -> !cir.ptr<!rec_NonTrivial>
// CIR: cir.call @_ZN10NonTrivialC1ERKS_(%[[ITR_LOAD]], %[[RHS_ELT]]) : (!cir.ptr<!rec_NonTrivial> {{.*}}, !cir.ptr<!rec_NonTrivial> {{.*}}) -> ()
// CIR: %[[ONE:.*]] = cir.const #cir.int<1> : !s64i
// CIR: %[[NEXT_ITR:.*]] = cir.ptr_stride %[[ITR_LOAD]], %[[ONE]] : (!cir.ptr<!rec_NonTrivial>, !s64i) -> !cir.ptr<!rec_NonTrivial>
@@ -96,7 +96,7 @@ struct HasMultiDimArray {
// CIR: %[[ITR1_LOAD:.*]] = cir.load {{.*}}%[[ITR1_ALLOCA]] : !cir.ptr<!cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>>>, !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>>
// CIR: %[[IDX1:.*]] = cir.ptr_diff %[[ITR1_LOAD]], %[[THIS_ARR_DECAY]] : !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>> -> !s64i
// CIR: %[[IDX1_CAST:.*]] = cir.cast integral %[[IDX1]] : !s64i -> !u64i
-// CIR: %[[RHS_ELT1:.*]] = cir.get_element %[[RHS_ARR]][%[[IDX1_CAST]] : !u64i] : !cir.ptr<!cir.array<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3> x 2>> -> !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>>
+// CIR: %[[RHS_ELT1:.*]] = cir.get_element inbounds nuw %[[RHS_ARR]][%[[IDX1_CAST]] : !u64i] : !cir.ptr<!cir.array<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3> x 2>> -> !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>>
// CIR: %[[ARR1_DECAY:.*]] = cir.cast array_to_ptrdecay %[[ITR1_LOAD]] : !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>> -> !cir.ptr<!cir.array<!rec_NonTrivial x 4>>
// CIR: cir.store {{.*}}%[[ARR1_DECAY]], %[[ITR2_ALLOCA]] : !cir.ptr<!cir.array<!rec_NonTrivial x 4>>, !cir.ptr<!cir.ptr<!cir.array<!rec_NonTrivial x 4>>>
// CIR: %[[SIZE2_CONST:.*]] = cir.const #cir.int<3> : !s64i
@@ -105,7 +105,7 @@ struct HasMultiDimArray {
// CIR: %[[ITR2_LOAD:.*]] = cir.load {{.*}}%[[ITR2_ALLOCA]] : !cir.ptr<!cir.ptr<!cir.array<!rec_NonTrivial x 4>>>, !cir.ptr<!cir.array<!rec_NonTrivial x 4>>
// CIR: %[[IDX2:.*]] = cir.ptr_diff %[[ITR2_LOAD]], %[[ARR1_DECAY]] : !cir.ptr<!cir.array<!rec_NonTrivial x 4>> -> !s64i
// CIR: %[[IDX2_CAST:.*]] = cir.cast integral %[[IDX2]] : !s64i -> !u64i
-// CIR: %[[RHS_ELT2:.*]] = cir.get_element %[[RHS_ELT1]][%[[IDX2_CAST]] : !u64i] : !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>> -> !cir.ptr<!cir.array<!rec_NonTrivial x 4>>
+// CIR: %[[RHS_ELT2:.*]] = cir.get_element inbounds nuw %[[RHS_ELT1]][%[[IDX2_CAST]] : !u64i] : !cir.ptr<!cir.array<!cir.array<!rec_NonTrivial x 4> x 3>> -> !cir.ptr<!cir.array<!rec_NonTrivial x 4>>
// CIR: %[[ARR2_DECAY:.*]] = cir.cast array_to_ptrdecay %[[ITR2_LOAD]] : !cir.ptr<!cir.array<!rec_NonTrivial x 4>> -> !cir.ptr<!rec_NonTrivial>
// CIR: cir.store {{.*}}%[[ARR2_DECAY]], %[[ITR3_ALLOCA]] : !cir.ptr<!rec_NonTrivial>, !cir.ptr<!cir.ptr<!rec_NonTrivial>>
// CIR: %[[SIZE3_CONST:.*]] = cir.const #cir.int<4> : !s64i
@@ -114,7 +114,7 @@ struct HasMultiDimArray {
// CIR: %[[ITR3_LOAD:.*]] = cir.load {{.*}}%[[ITR3_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_NonTrivial>>, !cir.ptr<!rec_NonTrivial>
// CIR: %[[IDX3:.*]] = cir.ptr_diff %[[ITR3_LOAD]], %[[ARR2_DECAY]] : !cir.ptr<!rec_NonTrivial> -> !s64i
// CIR: %[[IDX3_CAST:.*]] = cir.cast integral %[[IDX3]] : !s64i -> !u64i
-// CIR: %[[RHS_ELT3:.*]] = cir.get_element %[[RHS_ELT2]][%[[IDX3_CAST]] : !u64i] : !cir.ptr<!cir.array<!rec_NonTrivial x 4>> -> !cir.ptr<!rec_NonTrivial>
+// CIR: %[[RHS_ELT3:.*]] = cir.get_element inbounds nuw %[[RHS_ELT2]][%[[IDX3_CAST]] : !u64i] : !cir.ptr<!cir.array<!rec_NonTrivial x 4>> -> !cir.ptr<!rec_NonTrivial>
// CIR: cir.call @_ZN10NonTrivialC1ERKS_(%[[ITR3_LOAD]], %[[RHS_ELT3]]) : (!cir.ptr<!rec_NonTrivial> {{.*}}, !cir.ptr<!rec_NonTrivial> {{.*}}) -> ()
// CIR: %[[ONE:.*]] = cir.const #cir.int<1> : !s64i
// CIR: %[[NEXT_ITR3:.*]] = cir.ptr_stride %[[ITR3_LOAD]], %[[ONE]] : (!cir.ptr<!rec_NonTrivial>, !s64i) -> !cir.ptr<!rec_NonTrivial>
@@ -175,7 +175,7 @@ struct HasMultiDimArray {
// LLVM: %[[START_PTR:.*]] = ptrtoint ptr %[[THIS_ARR_DECAY]] to i64
// LLVM: %[[PTR_DIFF:.*]] = sub i64 %[[ITR_PTR]], %[[START_PTR]]
// LLVM: %[[IDX:.*]] = sdiv exact i64 %[[PTR_DIFF]], 4
-// LLVM: %[[RHS_ELT:.*]] = getelementptr [3 x %struct.NonTrivial], ptr %[[RHS_ARR]], i32 0, i64 %[[IDX]]
+// LLVM: %[[RHS_ELT:.*]] = getelementptr inbounds nuw [3 x %struct.NonTrivial], ptr %[[RHS_ARR]], i32 0, i64 %[[IDX]]
// LLVM: call void @_ZN10NonTrivialC1ERKS_(ptr {{.*}}%[[LOAD_ITR]], ptr {{.*}}%[[RHS_ELT]])
// LLVM: %[[NEXT_ITR:.*]] = getelementptr %struct.NonTrivial, ptr %[[LOAD_ITR]], i64 1
// LLVM: store ptr %[[NEXT_ITR]], ptr %[[ITR_ALLOCA]]
@@ -210,7 +210,7 @@ struct HasMultiDimArray {
// LLVM: %[[ITR1_START_PTR:.*]] = ptrtoint ptr %[[THIS_ARR_DECAY]] to i64
// LLVM: %[[PTR_DIFF:.*]] = sub i64 %[[ITR1_PTR]], %[[ITR1_START_PTR]]
// LLVM: %[[IDX1:.*]] = sdiv exact i64 %[[PTR_DIFF]], 48
-// LLVM: %[[RHS_ELT1:.*]] = getelementptr [2 x [3 x [4 x %struct.NonTrivial]]], ptr %[[RHS_ARR]], i32 0, i64 %[[IDX1]]
+// LLVM: %[[RHS_ELT1:.*]] = getelementptr inbounds nuw [2 x [3 x [4 x %struct.NonTrivial]]], ptr %[[RHS_ARR]], i32 0, i64 %[[IDX1]]
// LLVM: %[[ARR1_DECAY:.*]] = getelementptr [4 x %struct.NonTrivial], ptr %[[ITR1_LOAD]], i32 0
// LLVM: store ptr %[[ARR1_DECAY]], ptr %[[ITR2_ALLOCA]]
// LLVM: %[[END_ITR2:.*]] = getelementptr [4 x %struct.NonTrivial], ptr %[[ARR1_DECAY]], i64 3
@@ -227,7 +227,7 @@ struct HasMultiDimArray {
// LLVM: %[[ITR2_START_PTR:.*]] = ptrtoint ptr %[[ARR1_DECAY]] to i64
// LLVM: %[[PTR_DIFF:.*]] = sub i64 %[[ITR2_PTR]], %[[ITR2_START_PTR]]
// LLVM: %[[IDX2:.*]] = sdiv exact i64 %[[PTR_DIFF]], 16
-// LLVM: %[[RHS_ELT2:.*]] = getelementptr [3 x [4 x %struct.NonTrivial]], ptr %[[RHS_ELT1]], i32 0, i64 %[[IDX2]]
+// LLVM: %[[RHS_ELT2:.*]] = getelementptr inbounds nuw [3 x [4 x %struct.NonTrivial]], ptr %[[RHS_ELT1]], i32 0, i64 %[[IDX2]]
// LLVM: %[[ARR2_DECAY:.*]] = getelementptr %struct.NonTrivial, ptr %[[ITR2_LOAD]], i32 0
// LLVM: store ptr %[[ARR2_DECAY]], ptr %[[ITR3_ALLOCA]]
// LLVM: %[[END_ITR3:.*]] = getelementptr %struct.NonTrivial, ptr %[[ARR2_DECAY]], i64 4
@@ -244,7 +244,7 @@ struct HasMultiDimArray {
// LLVM: %[[ITR3_START_PTR:.*]] = ptrtoint ptr %[[ARR2_DECAY]] to i64
// LLVM: %[[PTR_DIFF:.*]] = sub i64 %[[ITR3_PTR]], %[[ITR3_START_PTR]]
// LLVM: %[[IDX3:.*]] = sdiv exact i64 %[[PTR_DIFF]], 4
-// LLVM: %[[RHS_ELT3:.*]] = getelementptr [4 x %struct.NonTrivial], ptr %[[RHS_ELT2]], i32 0, i64 %[[IDX3]]
+// LLVM: %[[RHS_ELT3:.*]] = getelementptr inbounds nuw [4 x %struct.NonTrivial], ptr %[[RHS_ELT2]], i32 0, i64 %[[IDX3]]
// LLVM: call void @_ZN10NonTrivialC1ERKS_(ptr{{.*}} %[[ITR3_LOAD]], ptr{{.*}} %[[RHS_ELT3]])
// LLVM: %[[NEXT_ITR3:.*]] = getelementptr %struct.NonTrivial, ptr %[[ITR3_LOAD]], i64 1
// LLVM: store ptr %[[NEXT_ITR3]], ptr %[[ITR3_ALLOCA]]
diff --git a/clang/test/CIR/CodeGen/array-subscript-inbounds.c b/clang/test/CIR/CodeGen/array-subscript-inbounds.c
new file mode 100644
index 00000000000000..c3f634c4a71916
--- /dev/null
+++ b/clang/test/CIR/CodeGen/array-subscript-inbounds.c
@@ -0,0 +1,232 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefixes=LLVM,LLVMCIR --input-file=%t-cir.ll %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefixes=LLVM,OGCG --input-file=%t.ll %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fwrapv-pointer -fclangir -emit-cir %s -o %t-wrapv.cir
+// RUN: FileCheck --check-prefix=WRAPV-CIR --input-file=%t-wrapv.cir %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fwrapv-pointer -fclangir -emit-llvm %s -o %t-wrapv-cir.ll
+// RUN: FileCheck --check-prefixes=WRAPV-LLVM,WRAPV-LLVMCIR --input-file=%t-wrapv-cir.ll %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fwrapv-pointer -emit-llvm %s -o %t-wrapv.ll
+// RUN: FileCheck --check-prefixes=WRAPV-LLVM,WRAPV-OGCG --input-file=%t-wrapv.ll %s
+
+int signed_index(int i) {
+ int a[4] = {0};
+ return a[i];
+}
+
+// CIR-LABEL: cir.func {{.*}} @signed_index(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[IDX:.*]] = cir.cast integral %{{.*}} : !s32i -> !s64i
+// CIR: cir.get_element inbounds %[[A]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @signed_index(
+// LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// LLVM: %[[IDX:.*]] = sext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[IDX]]
+
+// WRAPV-CIR-LABEL: cir.func {{.*}} @signed_index(
+// WRAPV-CIR: cir.get_element %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// WRAPV-LLVM-LABEL: define dso_local i32 @signed_index(
+// WRAPV-LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// WRAPV-LLVM: %[[IDX:.*]] = sext i32 %{{.*}} to i64
+// WRAPV-LLVM: getelementptr [4 x i32], ptr %[[A]],
+// WRAPV-LLVMCIR-SAME: i32 0,
+// WRAPV-OGCG-SAME: i64 0,
+// WRAPV-LLVM-SAME: i64 %[[IDX]]
+
+int unsigned_index(unsigned u) {
+ int a[4] = {0};
+ return a[u];
+}
+
+// CIR-LABEL: cir.func {{.*}} @unsigned_index(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[IDX:.*]] = cir.cast integral %{{.*}} : !u32i -> !u64i
+// CIR: cir.get_element inbounds nuw %[[A]][%[[IDX]] : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @unsigned_index(
+// LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// LLVM: %[[IDX:.*]] = zext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds nuw [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[IDX]]
+
+// WRAPV-CIR-LABEL: cir.func {{.*}} @unsigned_index(
+// WRAPV-CIR: cir.get_element %{{.*}}[%{{.*}} : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// WRAPV-LLVM-LABEL: define dso_local i32 @unsigned_index(
+// WRAPV-LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// WRAPV-LLVM: %[[IDX:.*]] = zext i32 %{{.*}} to i64
+// WRAPV-LLVM: getelementptr [4 x i32], ptr %[[A]],
+// WRAPV-LLVMCIR-SAME: i32 0,
+// WRAPV-OGCG-SAME: i64 0,
+// WRAPV-LLVM-SAME: i64 %[[IDX]]
+
+int constant_index(void) {
+ int a[4] = {0};
+ return a[1];
+}
+
+// CIR-LABEL: cir.func {{.*}} @constant_index(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[ONE:.*]] = cir.const #cir.int<1> : !s64i
+// CIR: cir.get_element inbounds %[[A]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @constant_index(
+// LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// LLVM: getelementptr inbounds [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 1
+
+int unsigned_constant_index(void) {
+ int a[4] = {0};
+ return a[1u];
+}
+
+// CIR-LABEL: cir.func {{.*}} @unsigned_constant_index(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[ONE:.*]] = cir.const #cir.int<1> : !u64i
+// CIR: cir.get_element inbounds nuw %[[A]][%[[ONE]] : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @unsigned_constant_index(
+// LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// LLVM: getelementptr inbounds nuw [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 1
+
+enum SignedIdx { NegOne = -1, SZero, SOne };
+enum UnsignedIdx { UZero, UOne };
+
+int enum_index(enum SignedIdx s, enum UnsignedIdx u) {
+ int a[4] = {0};
+ return a[s] + a[u];
+}
+
+// CIR-LABEL: cir.func {{.*}} @enum_index(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[S:.*]] = cir.cast integral %{{.*}} : !s32i -> !s64i
+// CIR: cir.get_element inbounds %[[A]][%[[S]] : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+// CIR: %[[U:.*]] = cir.cast integral %{{.*}} : !u32i -> !u64i
+// CIR: cir.get_element inbounds nuw %[[A]][%[[U]] : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @enum_index(
+// LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// LLVM: %[[S:.*]] = sext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[S]]
+// LLVM: %[[U:.*]] = zext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds nuw [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[U]]
+
+int reversed_subscript(int i, unsigned u) {
+ int a[4] = {0};
+ return i[a] + u[a];
+}
+
+// CIR-LABEL: cir.func {{.*}} @reversed_subscript(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[I:.*]] = cir.cast integral %{{.*}} : !s32i -> !s64i
+// CIR: cir.get_element inbounds %[[A]][%[[I]] : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+// CIR: %[[U:.*]] = cir.cast integral %{{.*}} : !u32i -> !u64i
+// CIR: cir.get_element inbounds nuw %[[A]][%[[U]] : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @reversed_subscript(
+// LLVM: %[[A:.*]] = alloca [4 x i32], align 16
+// LLVM: %[[I:.*]] = sext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[I]]
+// LLVM: %[[U:.*]] = zext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds nuw [4 x i32], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[U]]
+
+int two_dimensions(int i, unsigned j) {
+ int a[3][4] = {{0}};
+ return a[i][j];
+}
+
+// CIR-LABEL: cir.func {{.*}} @two_dimensions(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!cir.array<!s32i x 4> x 3>>
+// CIR: %[[I:.*]] = cir.cast integral %{{.*}} : !s32i -> !s64i
+// CIR: %[[ROW:.*]] = cir.get_element inbounds %[[A]][%[[I]] : !s64i] : !cir.ptr<!cir.array<!cir.array<!s32i x 4> x 3>> -> !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: %[[J:.*]] = cir.cast integral %{{.*}} : !u32i -> !u64i
+// CIR: cir.get_element inbounds nuw %[[ROW]][%[[J]] : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @two_dimensions(
+// LLVM: %[[A:.*]] = alloca [3 x [4 x i32]], align 16
+// LLVM: %[[I:.*]] = sext i32 %{{.*}} to i64
+// LLVM: %[[ROW:.*]] = getelementptr inbounds [3 x [4 x i32]], ptr %[[A]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[I]]
+// LLVM: %[[J:.*]] = zext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds nuw [4 x i32], ptr %[[ROW]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[J]]
+
+int pointer_to_2d(int (*p)[3][4], int i, int j) { return (*p)[i][j]; }
+
+// CIR-LABEL: cir.func {{.*}} @pointer_to_2d(
+// CIR: %[[ROW:.*]] = cir.get_element inbounds %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!s32i x 4> x 3>> -> !cir.ptr<!cir.array<!s32i x 4>>
+// CIR: cir.get_element inbounds %[[ROW]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @pointer_to_2d(
+// LLVM: %[[ROW:.*]] = getelementptr inbounds [3 x [4 x i32]], ptr
+// LLVM: getelementptr inbounds [4 x i32], ptr %[[ROW]],
+
+struct Flex {
+ int n;
+ int a[];
+};
+
+int flexible_member(struct Flex *s, int i) { return s->a[i]; }
+
+// CIR-LABEL: cir.func {{.*}} @flexible_member(
+// CIR: %[[MEMBER:.*]] = cir.get_member %{{.*}}[1] {name = "a"} : !cir.ptr<!rec_Flex> -> !cir.ptr<!cir.array<!s32i x 0>>
+// CIR: %[[IDX:.*]] = cir.cast integral %{{.*}} : !s32i -> !s64i
+// CIR: cir.get_element inbounds %[[MEMBER]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 0>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @flexible_member(
+// LLVM: %[[MEMBER:.*]] = getelementptr inbounds nuw %struct.Flex, ptr %{{.*}}, i32 0, i32 1
+// LLVM: %[[IDX:.*]] = sext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds [0 x i32], ptr %[[MEMBER]],
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[IDX]]
+
+static const int debruijn[32] = {
+ 0, 1, 28, 2, 29, 14, 24, 3, 30, 22, 20, 15, 25, 17, 4, 8,
+ 31, 27, 13, 23, 21, 19, 16, 7, 26, 12, 18, 6, 11, 5, 10, 9};
+
+int count_trailing_zeros(unsigned x) {
+ return debruijn[((x & -x) * 0x077CB531u) >> 27];
+}
+
+// CIR-LABEL: cir.func {{.*}} @count_trailing_zeros(
+// CIR: %[[TABLE:.*]] = cir.get_global @debruijn : !cir.ptr<!cir.array<!s32i x 32>>
+// CIR: %[[IDX:.*]] = cir.cast integral %{{.*}} : !u32i -> !u64i
+// CIR: cir.get_element inbounds nuw %[[TABLE]][%[[IDX]] : !u64i] : !cir.ptr<!cir.array<!s32i x 32>> -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define dso_local i32 @count_trailing_zeros(
+// LLVM: %[[IDX:.*]] = zext i32 %{{.*}} to i64
+// LLVM: getelementptr inbounds nuw [32 x i32], ptr @debruijn,
+// LLVMCIR-SAME: i32 0,
+// OGCG-SAME: i64 0,
+// LLVM-SAME: i64 %[[IDX]]
diff --git a/clang/test/CIR/CodeGen/array-subscript-order.cpp b/clang/test/CIR/CodeGen/array-subscript-order.cpp
index c064b15a06f8b8..8cf5b5d8935748 100644
--- a/clang/test/CIR/CodeGen/array-subscript-order.cpp
+++ b/clang/test/CIR/CodeGen/array-subscript-order.cpp
@@ -57,12 +57,12 @@ int test3(void) {
// CIR-LABEL: cir.func{{.*}} @_Z5test3v
// CIR: %[[BASE:.*]] = cir.call @_Z6getArrv()
// CIR-NEXT: %[[IDX:.*]] = cir.call @_Z5idxFnv()
-// CIR: cir.get_element %[[BASE]][{{.*}}]
+// CIR: cir.get_element inbounds %[[BASE]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
// LLVM-LABEL: define{{.*}} i32 @_Z5test3v()
// LLVM: %[[BASE:.*]] = call {{.*}} ptr @_Z6getArrv()
// LLVM-NEXT: %[[IDX:.*]] = call {{.*}} i32 @_Z5idxFnv()
-// LLVM: getelementptr {{.*}}[4 x i32], ptr %[[BASE]],
+// LLVM: getelementptr inbounds [4 x i32], ptr %[[BASE]],
int test_vec(void) {
return getVec()[idxFn()];
diff --git a/clang/test/CIR/CodeGen/array.cpp b/clang/test/CIR/CodeGen/array.cpp
index 61017e019c207b..c68fbf08845e57 100644
--- a/clang/test/CIR/CodeGen/array.cpp
+++ b/clang/test/CIR/CodeGen/array.cpp
@@ -169,11 +169,11 @@ void func() {
// CIR: %[[INIT:.*]] = cir.alloca "e" {{.*}} init : !cir.ptr<!s32i>
// CIR: %[[INIT_2:.*]] = cir.alloca "e2" {{.*}} init : !cir.ptr<!s32i>
// CIR: %[[IDX:.*]] = cir.const #cir.int<0> : !s64i
-// CIR: %[[ELE_PTR:.*]] = cir.get_element %[[ARR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE_PTR:.*]] = cir.get_element inbounds %[[ARR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
// CIR: %[[TMP:.*]] = cir.load{{.*}} %[[ELE_PTR]] : !cir.ptr<!s32i>, !s32i
// CIR" cir.store %[[TMP]], %[[INIT]] : !s32i, !cir.ptr<!s32i>
// CIR: %[[IDX:.*]] = cir.const #cir.int<1> : !s64i
-// CIR: %[[ELE_PTR:.*]] = cir.get_element %[[ARR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE_PTR:.*]] = cir.get_element inbounds %[[ARR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
// CIR: %[[TMP:.*]] = cir.load{{.*}} %[[ELE_PTR]] : !cir.ptr<!s32i>, !s32i
// CIR" cir.store %[[TMP]], %[[INIT_2]] : !s32i, !cir.ptr<!s32i>
@@ -181,10 +181,10 @@ void func() {
// LLVM-NEXT: %[[ARR:.*]] = alloca [10 x i32], align 16
// LLVM-NEXT: %[[INIT:.*]] = alloca i32, align 4
// LLVM-NEXT: %[[INIT_2:.*]] = alloca i32, align 4
-// LLVM-NEXT: %[[ELE_PTR:.*]] = getelementptr [10 x i32], ptr %[[ARR]], i32 0, i64 0
+// LLVM-NEXT: %[[ELE_PTR:.*]] = getelementptr inbounds [10 x i32], ptr %[[ARR]], i32 0, i64 0
// LLVM-NEXT: %[[TMP_1:.*]] = load i32, ptr %[[ELE_PTR]], align 16
// LLVM-NEXT: store i32 %[[TMP_1]], ptr %[[INIT]], align 4
-// LLVM-NEXT: %[[ELE_PTR:.*]] = getelementptr [10 x i32], ptr %[[ARR]], i32 0, i64 1
+// LLVM-NEXT: %[[ELE_PTR:.*]] = getelementptr inbounds [10 x i32], ptr %[[ARR]], i32 0, i64 1
// LLVM-NEXT: %[[TMP_2:.*]] = load i32, ptr %[[ELE_PTR]], align 4
// LLVM-NEXT: store i32 %[[TMP_2]], ptr %[[INIT_2]], align 4
@@ -229,7 +229,7 @@ void func3() {
// CIR: cir.store{{.*}} %[[IDX_V]], %[[IDX]] : !s32i, !cir.ptr<!s32i>
// CIR: %[[TMP_IDX:.*]] = cir.load{{.*}} %[[IDX]] : !cir.ptr<!s32i>, !s32i
// CIR: %[[TMP_IDX_64:.*]] = cir.cast integral %[[TMP_IDX]] : !s32i -> !s64i
-// CIR: %[[ELE_PTR:.*]] = cir.get_element %[[ARR]][%[[TMP_IDX_64]] : !s64i] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE_PTR:.*]] = cir.get_element inbounds %[[ARR]][%[[TMP_IDX_64]] : !s64i] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!s32i>
// CIR: %[[ELE_TMP:.*]] = cir.load{{.*}} %[[ELE_PTR]] : !cir.ptr<!s32i>, !s32i
// CIR: cir.store{{.*}} %[[ELE_TMP]], %[[INIT]] : !s32i, !cir.ptr<!s32i>
@@ -241,7 +241,7 @@ void func3() {
// LLVM: store i32 1, ptr %[[IDX]], align 4
// LLVM: %[[TMP1:.*]] = load i32, ptr %[[IDX]], align 4
// LLVM: %[[IDX_I64:.*]] = sext i32 %[[TMP1]] to i64
-// LLVM: %[[ELE:.*]] = getelementptr [2 x i32], ptr %[[ARR]], i32 0, i64 %[[IDX_I64]]
+// LLVM: %[[ELE:.*]] = getelementptr inbounds [2 x i32], ptr %[[ARR]], i32 0, i64 %[[IDX_I64]]
// LLVM: %[[TMP2:.*]] = load i32, ptr %[[ELE]], align 4
// LLVM: store i32 %[[TMP2]], ptr %[[INIT]], align 4
@@ -266,9 +266,9 @@ void func4() {
// CIR: %[[CONST:.*]] = cir.get_global @[[FUNC4_ARR]] : !cir.ptr<!cir.array<!cir.array<!s32i x 1> x 2>>
// CIR: cir.copy %[[CONST]] to %[[ARR]] : !cir.ptr<!cir.array<!cir.array<!s32i x 1> x 2>>
// CIR: %[[IDX_1:.*]] = cir.const #cir.int<1> : !s64i
-// CIR: %[[ARR_1:.*]] = cir.get_element %[[ARR]][%[[IDX_1]] : !s64i] : !cir.ptr<!cir.array<!cir.array<!s32i x 1> x 2>> -> !cir.ptr<!cir.array<!s32i x 1>>
+// CIR: %[[ARR_1:.*]] = cir.get_element inbounds %[[ARR]][%[[IDX_1]] : !s64i] : !cir.ptr<!cir.array<!cir.array<!s32i x 1> x 2>> -> !cir.ptr<!cir.array<!s32i x 1>>
// CIR: %[[IDX:.*]] = cir.const #cir.int<0> : !s64i
-// CIR: %[[ELE_0:.*]] = cir.get_element %[[ARR_1]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 1>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE_0:.*]] = cir.get_element inbounds %[[ARR_1]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 1>> -> !cir.ptr<!s32i>
// CIR: %[[TMP:.*]] = cir.load{{.*}} %[[ELE_0]] : !cir.ptr<!s32i>, !s32i
// CIR: cir.store{{.*}} %[[TMP]], %[[INIT]] : !s32i, !cir.ptr<!s32i>
@@ -276,8 +276,8 @@ void func4() {
// LLVM: %[[ARR:.*]] = alloca [2 x [1 x i32]], align 4
// LLVM: %[[INIT:.*]] = alloca i32, align 4
// LLVM: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %[[ARR]], ptr align 4 @[[FUNC4_ARR:.*]], i64 8, i1 false)
-// LLVM: %[[ARR_1:.*]] = getelementptr [2 x [1 x i32]], ptr %[[ARR]], i32 0, i64 1
-// LLVM: %[[ELE_PTR:.*]] = getelementptr [1 x i32], ptr %[[ARR_1]], i32 0, i64 0
+// LLVM: %[[ARR_1:.*]] = getelementptr inbounds [2 x [1 x i32]], ptr %[[ARR]], i32 0, i64 1
+// LLVM: %[[ELE_PTR:.*]] = getelementptr inbounds [1 x i32], ptr %[[ARR_1]], i32 0, i64 0
// LLVM: %[[TMP:.*]] = load i32, ptr %[[ELE_PTR]], align 4
// LLVM: store i32 %[[TMP]], ptr %[[INIT]], align 4
@@ -414,7 +414,7 @@ void func9(int arr[10][5]) {
// CIR: %[[IDX_1:.*]] = cir.const #cir.int<1> : !s64i
// CIR: %[[ARR_1:.*]] = cir.ptr_stride %[[TMP_1]], %[[IDX_1]] : (!cir.ptr<!cir.array<!s32i x 5>>, !s64i) -> !cir.ptr<!cir.array<!s32i x 5>>
// CIR: %[[IDX:.*]] = cir.const #cir.int<2> : !s64i
-// CIR: %[[ARR_1_2:.*]] = cir.get_element %[[ARR_1]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 5>> -> !cir.ptr<!s32i>
+// CIR: %[[ARR_1_2:.*]] = cir.get_element inbounds %[[ARR_1]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 5>> -> !cir.ptr<!s32i>
// CIR: %[[TMP_2:.*]] = cir.load{{.*}} %[[ARR_1_2]] : !cir.ptr<!s32i>, !s32i
// CIR: cir.store{{.*}} %[[TMP_2]], %[[INIT]] : !s32i, !cir.ptr<!s32i>
@@ -424,7 +424,7 @@ void func9(int arr[10][5]) {
// LLVM: store ptr %[[ARG]], ptr %[[ARR]], align 8
// LLVM: %[[TMP_1:.*]] = load ptr, ptr %[[ARR]], align 8
// LLVM: %[[ARR_1:.*]] = getelementptr [5 x i32], ptr %[[TMP_1]], i64 1
-// LLVM: %[[ARR_1_2:.*]] = getelementptr [5 x i32], ptr %[[ARR_1]], i32 0, i64 2
+// LLVM: %[[ARR_1_2:.*]] = getelementptr inbounds [5 x i32], ptr %[[ARR_1]], i32 0, i64 2
// LLVM: %[[TMP_2:.*]] = load i32, ptr %[[ARR_1_2]], align 4
// LLVM: store i32 %[[TMP_2]], ptr %[[INIT]], align 4
@@ -517,7 +517,7 @@ void boolean_index_access(int x) {
// CIR: %[[FOUR:.*]] = cir.const #cir.int<4> : !s32i
// CIR: %[[CMP:.*]] = cir.cmp gt %[[X]], %[[FOUR]] : !s32i
// CIR: %[[IDX:.*]] = cir.cast bool_to_int %[[CMP]] : !cir.bool -> !u64i
-// CIR: %[[ELE:.*]] = cir.get_element %[[ARR_ADDR]][%[[IDX]] : !u64i] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE:.*]] = cir.get_element inbounds nuw %[[ARR_ADDR]][%[[IDX]] : !u64i] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!s32i>
// CIR: %[[VAL:.*]] = cir.load{{.*}} %[[ELE]] : !cir.ptr<!s32i>, !s32i
// CIR: cir.store{{.*}} %[[VAL]], %[[N_ADDR]] : !s32i, !cir.ptr<!s32i>
@@ -527,7 +527,7 @@ void boolean_index_access(int x) {
// LLVM: %[[X:.*]] = load i32, ptr %[[X_ADDR]]
// LLVM: %[[CMP:.*]] = icmp sgt i32 %[[X]], 4
// LLVM: %[[IDX:.*]] = zext i1 %[[CMP]] to i64
-// LLVM: %[[ELE:.*]] = getelementptr [2 x i32], ptr %[[ARR_ADDR]], i32 0, i64 %[[IDX]]
+// LLVM: %[[ELE:.*]] = getelementptr inbounds nuw [2 x i32], ptr %[[ARR_ADDR]], i32 0, i64 %[[IDX]]
// LLVM: %[[VAL:.*]] = load i32, ptr %[[ELE]]
// LLVM: store i32 %[[VAL]], ptr %[[N_ADDR]]
@@ -555,7 +555,7 @@ void bitint_index_access(_BitInt(7) i) {
// CIR: %[[N_ADDR:.*]] = cir.alloca "n" {{.*}} init : !cir.ptr<!s32i>
// CIR: %[[I:.*]] = cir.load{{.*}} %[[I_ADDR]] : !cir.ptr<!cir.int<s, 7, bitint>>, !cir.int<s, 7, bitint>
// CIR: %[[IDX:.*]] = cir.cast integral %[[I]] : !cir.int<s, 7, bitint> -> !s64i
-// CIR: %[[ELE:.*]] = cir.get_element %[[ARR_ADDR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE:.*]] = cir.get_element inbounds %[[ARR_ADDR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
// CIR: %[[VAL:.*]] = cir.load{{.*}} %[[ELE]] : !cir.ptr<!s32i>, !s32i
// CIR: cir.store{{.*}} %[[VAL]], %[[N_ADDR]] : !s32i, !cir.ptr<!s32i>
@@ -566,7 +566,7 @@ void bitint_index_access(_BitInt(7) i) {
// LLVM: %[[I:.*]] = load i8, ptr %[[I_ADDR]]
// LLVM: %[[I_CAST:.*]] = trunc i8 %[[I]] to i7
// LLVM: %[[IDX:.*]] = sext i7 %[[I_CAST]] to i64
-// LLVM: %[[ELE:.*]] = getelementptr [10 x i32], ptr %[[ARR_ADDR]], i32 0, i64 %[[IDX]]
+// LLVM: %[[ELE:.*]] = getelementptr inbounds [10 x i32], ptr %[[ARR_ADDR]], i32 0, i64 %[[IDX]]
// LLVM: %[[VAL:.*]] = load i32, ptr %[[ELE]]
// LLVM: store i32 %[[VAL]], ptr %[[N_ADDR]]
@@ -588,7 +588,7 @@ char access_last_element_of_huge_array() {
// CIR: cir.func{{.*}} @_Z33access_last_element_of_huge_arrayv
// CIR: %[[HUGE:.*]] = cir.get_global @huge : !cir.ptr<!cir.array<!s8i x 2305843009213693951>>
// CIR: %[[IDX:.*]] = cir.const #cir.int<2305843009213693950> : !u64i
-// CIR: %[[ELE:.*]] = cir.get_element %[[HUGE]][%[[IDX]] : !u64i] : !cir.ptr<!cir.array<!s8i x 2305843009213693951>> -> !cir.ptr<!s8i>
+// CIR: %[[ELE:.*]] = cir.get_element inbounds nuw %[[HUGE]][%[[IDX]] : !u64i] : !cir.ptr<!cir.array<!s8i x 2305843009213693951>> -> !cir.ptr<!s8i>
// CIR: cir.load{{.*}} %[[ELE]] : !cir.ptr<!s8i>, !s8i
// LLVM: define{{.*}} @_Z33access_last_element_of_huge_arrayv()
diff --git a/clang/test/CIR/CodeGen/binassign.c b/clang/test/CIR/CodeGen/binassign.c
index 13ef5b7db0555e..719aea80defa29 100644
--- a/clang/test/CIR/CodeGen/binassign.c
+++ b/clang/test/CIR/CodeGen/binassign.c
@@ -124,7 +124,7 @@ int ignore_result_assign() {
// CIR: %[[VAL_5:.*]] = cir.const #cir.int<5> : !s32i
// CIR: cir.store{{.*}} %[[VAL_5]], %[[I]] : !s32i, !cir.ptr<!s32i>
// CIR: %[[VAL_5_64:.*]] = cir.const #cir.int<5> : !s64i
-// CIR: %[[ARR_ELEM:.*]] = cir.get_element %[[ARR]][%[[VAL_5_64]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CIR: %[[ARR_ELEM:.*]] = cir.get_element inbounds %[[ARR]][%[[VAL_5_64]] : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
// CIR: %[[ARR_LOAD:.*]] = cir.load{{.*}} %[[ARR_ELEM]] : !cir.ptr<!s32i>, !s32i
// CIR: cir.store{{.*}} %[[ARR_LOAD]], %[[J]] : !s32i, !cir.ptr<!s32i>
// CIR: %[[NULL:.*]] = cir.const #cir.ptr<null> : !cir.ptr<!s32i>
@@ -155,7 +155,7 @@ int ignore_result_assign() {
// LLVM: store i32 123, ptr %[[I_PTR]]
// LLVM: store i32 123, ptr %[[J_PTR]]
// LLVM: store i32 5, ptr %[[I_PTR]]
-// LLVM: %[[GEP:.*]] = getelementptr [10 x i32], ptr %[[ARR_PTR]], i32 0, i64 5
+// LLVM: %[[GEP:.*]] = getelementptr inbounds [10 x i32], ptr %[[ARR_PTR]], i32 0, i64 5
// LLVM: %[[ARR_VAL:.*]] = load i32, ptr %[[GEP]]
// LLVM: store i32 %[[ARR_VAL]], ptr %[[J_PTR]]
// LLVM: store ptr null, ptr %[[Q_PTR]]
diff --git a/clang/test/CIR/CodeGen/bitint-record-layout.c b/clang/test/CIR/CodeGen/bitint-record-layout.c
index a16a79a4090107..d4c95bee2a49ba 100644
--- a/clang/test/CIR/CodeGen/bitint-record-layout.c
+++ b/clang/test/CIR/CodeGen/bitint-record-layout.c
@@ -202,7 +202,7 @@ _BitInt(128) get_bi(void) { return l128[1].bi; }
// CIR-NEXT: %[[RET_ALLOC:.*]] = cir.alloca "__retval" align(8) : !cir.ptr<!s128i_bitint>
// CIR-NEXT: %[[GET_GLOB:.*]] = cir.get_global @l128 : !cir.ptr<!cir.array<!rec_Last128 x 2>>
// CIR-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s64i
-// CIR-NEXT: %[[ARR_GEP:.*]] = cir.get_element %[[GET_GLOB]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!rec_Last128 x 2>> -> !cir.ptr<!rec_Last128>
+// CIR-NEXT: %[[ARR_GEP:.*]] = cir.get_element inbounds %[[GET_GLOB]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!rec_Last128 x 2>> -> !cir.ptr<!rec_Last128>
// CIR-NEXT: %[[GET_BI:.*]] = cir.get_member %[[ARR_GEP]][2] {name = "bi"} : !cir.ptr<!rec_Last128> -> !cir.ptr<!s128i_bitint>
// CIR-NEXT: %[[LOAD_BI:.*]] = cir.load align(8) %[[GET_BI]] : !cir.ptr<!s128i_bitint>, !s128i_bitint
// CIR-NEXT: cir.store %[[LOAD_BI]], %[[RET_ALLOC]] : !s128i_bitint, !cir.ptr<!s128i_bitint>
@@ -216,10 +216,10 @@ _BitInt(128) get_bi2(void) { return arrMem[1].bi[1]; }
// CIR-NEXT: %[[RET_ALLOC:.*]] = cir.alloca "__retval" align(8) : !cir.ptr<!s128i_bitint>
// CIR-NEXT: %[[GET_GLOB:.*]] = cir.get_global @arrMem : !cir.ptr<!cir.array<!rec_ArrMem x 2>>
// CIR-NEXT: %[[ONE_2:.*]] = cir.const #cir.int<1> : !s64i
-// CIR-NEXT: %[[ARR_GEP:.*]] = cir.get_element %[[GET_GLOB]][%[[ONE_2]] : !s64i] : !cir.ptr<!cir.array<!rec_ArrMem x 2>> -> !cir.ptr<!rec_ArrMem>
+// CIR-NEXT: %[[ARR_GEP:.*]] = cir.get_element inbounds %[[GET_GLOB]][%[[ONE_2]] : !s64i] : !cir.ptr<!cir.array<!rec_ArrMem x 2>> -> !cir.ptr<!rec_ArrMem>
// CIR-NEXT: %[[GET_BI_ARR:.*]] = cir.get_member %[[ARR_GEP]][2] {name = "bi"} : !cir.ptr<!rec_ArrMem> -> !cir.ptr<!cir.array<!s128i_bitint x 2>>
// CIR-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s64i
-// CIR-NEXT: %[[GET_BI_ELT:.*]] = cir.get_element %[[GET_BI_ARR]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!s128i_bitint x 2>> -> !cir.ptr<!s128i_bitint>
+// CIR-NEXT: %[[GET_BI_ELT:.*]] = cir.get_element inbounds %[[GET_BI_ARR]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!s128i_bitint x 2>> -> !cir.ptr<!s128i_bitint>
// CIR-NEXT: %[[LOAD_BI:.*]] = cir.load align(8) %[[GET_BI_ELT]] : !cir.ptr<!s128i_bitint>, !s128i_bitint
// CIR-NEXT: cir.store %[[LOAD_BI]], %[[RET_ALLOC]] : !s128i_bitint, !cir.ptr<!s128i_bitint>
// CIR-NEXT: %[[RET_LOAD:.*]] = cir.load %[[RET_ALLOC]] : !cir.ptr<!s128i_bitint>, !s128i_bitint
diff --git a/clang/test/CIR/CodeGen/bitint-wide.c b/clang/test/CIR/CodeGen/bitint-wide.c
index ac96abc4c92742..6b843845fc05d2 100644
--- a/clang/test/CIR/CodeGen/bitint-wide.c
+++ b/clang/test/CIR/CodeGen/bitint-wide.c
@@ -89,10 +89,10 @@ _BitInt(129) local_arr(void) {
// CIR: cir.copy %[[CONST_ARR]] align(16) to %[[ARR]] align(16) : !cir.ptr<!cir.array<!cir.int<s, 129, bitint> x 2>>
// CIR: %[[THREE:.*]] = cir.const #cir.int<3> : !cir.int<s, 129, bitint>
// CIR: %[[ONE:.*]] = cir.const #cir.int<1> : !s64i
-// CIR: %[[GET_ELT_ONE:.*]] = cir.get_element %[[ARR]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!cir.int<s, 129, bitint> x 2>> -> !cir.ptr<!cir.int<s, 129, bitint>>
+// CIR: %[[GET_ELT_ONE:.*]] = cir.get_element inbounds %[[ARR]][%[[ONE]] : !s64i] : !cir.ptr<!cir.array<!cir.int<s, 129, bitint> x 2>> -> !cir.ptr<!cir.int<s, 129, bitint>>
// CIR: cir.store align(8) %[[THREE]], %[[GET_ELT_ONE]] : !cir.int<s, 129, bitint>, !cir.ptr<!cir.int<s, 129, bitint>>
// CIR: %[[ZERO:.*]] = cir.const #cir.int<0> : !s64i
-// CIR: %[[GET_ELT_ZERO:.*]] = cir.get_element %[[ARR]][%[[ZERO]] : !s64i] : !cir.ptr<!cir.array<!cir.int<s, 129, bitint> x 2>> -> !cir.ptr<!cir.int<s, 129, bitint>>
+// CIR: %[[GET_ELT_ZERO:.*]] = cir.get_element inbounds %[[ARR]][%[[ZERO]] : !s64i] : !cir.ptr<!cir.array<!cir.int<s, 129, bitint> x 2>> -> !cir.ptr<!cir.int<s, 129, bitint>>
// CIR: %[[LOAD_ELT_ZERO:.*]] = cir.load align(16) %[[GET_ELT_ZERO]] : !cir.ptr<!cir.int<s, 129, bitint>>, !cir.int<s, 129, bitint>
// CIR: cir.store %[[LOAD_ELT_ZERO]], %[[RET_ARG]] : !cir.int<s, 129, bitint>, !cir.ptr<!cir.int<s, 129, bitint>>
@@ -100,9 +100,13 @@ _BitInt(129) local_arr(void) {
// LLVM-SAME: ptr{{.*}}%[[RET_ARG:.*]])
// LLVM: %[[ARR:.*]] = alloca [2 x [24 x i8]], align 16
// LLVM: call void @llvm.memcpy.p0.p0.i64(ptr align 16 %[[ARR]], ptr align 16 @__const.local_arr.arr, i64 48, i1 false)
-// LLVM: %[[GET_ELT_ONE:.*]] = getelementptr{{.*}} [2 x [24 x i8]], ptr %[[ARR]], i{{.*}} 0, i64 1
+// LLVM: %[[GET_ELT_ONE:.*]] = getelementptr inbounds [2 x [24 x i8]], ptr %[[ARR]],
+// CIRONLY-SAME: i32 0, i64 1
+// OGCG-SAME: i64 0, i64 1
// LLVM: store i192 3, ptr %[[GET_ELT_ONE]], align 8
-// LLVM: %[[GET_ELT_ZERO:.*]] = getelementptr{{.*}} [2 x [24 x i8]], ptr %[[ARR]], i{{.*}} 0, i64 0
+// LLVM: %[[GET_ELT_ZERO:.*]] = getelementptr inbounds [2 x [24 x i8]], ptr %[[ARR]],
+// CIRONLY-SAME: i32 0, i64 0
+// OGCG-SAME: i64 0, i64 0
// LLVM: %[[LOAD_ELT_ZERO:.*]] = load i192, ptr %[[GET_ELT_ZERO]], align 16
// LLVM: %[[TRUNC_ELT:.*]] = trunc i192 %[[LOAD_ELT_ZERO]] to i129
// LLVM: %[[EXT_ELT:.*]] = sext i129 %[[TRUNC_ELT]] to i192
diff --git a/clang/test/CIR/CodeGen/builtin-cpu-supports.c b/clang/test/CIR/CodeGen/builtin-cpu-supports.c
index eecde11ce26f33..d3ac98aa9aed57 100644
--- a/clang/test/CIR/CodeGen/builtin-cpu-supports.c
+++ b/clang/test/CIR/CodeGen/builtin-cpu-supports.c
@@ -15,7 +15,7 @@ extern void a(const char *);
// CIR-NEXT: %[[CPU_MODEL:.*]] = cir.get_global @__cpu_model : !cir.ptr<!rec_anon_struct>
// CIR-NEXT: %[[CPU_FEAT:.*]] = cir.get_member %[[CPU_MODEL]][3] {name = "__cpu_features"} : !cir.ptr<!rec_anon_struct> -> !cir.ptr<!cir.array<!u32i x 1>>
// CIR-NEXT: %[[ZERO:.*]] = cir.const #cir.int<0> : !u32i
-// CIR-NEXT: %[[FEAT0:.*]] = cir.get_element %[[CPU_FEAT]][%[[ZERO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 1>> -> !cir.ptr<!u32i>
+// CIR-NEXT: %[[FEAT0:.*]] = cir.get_element inbounds %[[CPU_FEAT]][%[[ZERO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 1>> -> !cir.ptr<!u32i>
// CIR-NEXT: %[[LOAD_FEAT0:.*]] = cir.load {{.*}} %[[FEAT0]] : !cir.ptr<!u32i>, !u32i
// CIR-NEXT: %[[MASK:.*]] = cir.const #cir.int<256> : !u32i
// CIR-NEXT: %[[AND:.*]] = cir.and %[[LOAD_FEAT0]], %[[MASK]] : !u32i
@@ -29,7 +29,7 @@ extern void a(const char *);
// CIR-NEXT: cir.scope {
// CIR-NEXT: %[[CPU_FEAT2:.*]] = cir.get_global @__cpu_features2 : !cir.ptr<!cir.array<!u32i x 3>>
// CIR-NEXT: %[[ZERO:.*]] = cir.const #cir.int<0> : !u32i
-// CIR-NEXT: %[[FEAT2_0:.*]] = cir.get_element %[[CPU_FEAT2]][%[[ZERO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
+// CIR-NEXT: %[[FEAT2_0:.*]] = cir.get_element inbounds %[[CPU_FEAT2]][%[[ZERO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
// CIR-NEXT: %[[LOAD_FEAT2_0:.*]] = cir.load {{.*}} %[[FEAT2_0]] : !cir.ptr<!u32i>, !u32i
// CIR-NEXT: %[[MASK:.*]] = cir.const #cir.int<1> : !u32i
// CIR-NEXT: %[[AND:.*]] = cir.and %[[LOAD_FEAT2_0]], %[[MASK]] : !u32i
@@ -98,7 +98,7 @@ int main(void) {
// CIR-NEXT: %[[RETVAL:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!s32i>
// CIR-NEXT: %[[CPU_FEAT2:.*]] = cir.get_global @__cpu_features2 : !cir.ptr<!cir.array<!u32i x 3>>
// CIR-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !u32i
-// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element %[[CPU_FEAT2]][%[[ONE]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
+// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element inbounds %[[CPU_FEAT2]][%[[ONE]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
// CIR-NEXT: %[[FEAT2_ELT_LOAD:.*]] = cir.load {{.*}} %[[FEAT2_ELT]] : !cir.ptr<!u32i>, !u32i
// CIR-NEXT: %[[MASK:.*]] = cir.const #cir.int<2147483648> : !u32i
// CIR-NEXT: %[[AND:.*]] = cir.and %[[FEAT2_ELT_LOAD]], %[[MASK]] : !u32i
@@ -132,7 +132,7 @@ int baseline() { return __builtin_cpu_supports("x86-64"); }
// CIR-NEXT: %[[RETVAL:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!s32i>
// CIR-NEXT: %[[CPU_FEAT2:.*]] = cir.get_global @__cpu_features2 : !cir.ptr<!cir.array<!u32i x 3>>
// CIR-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !u32i
-// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element %[[CPU_FEAT2]][%[[TWO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
+// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element inbounds %[[CPU_FEAT2]][%[[TWO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
// CIR-NEXT: %[[FEAT2_ELT_LOAD:.*]] = cir.load {{.*}} %[[FEAT2_ELT]] : !cir.ptr<!u32i>, !u32i
// CIR-NEXT: %[[MASK:.*]] = cir.const #cir.int<1> : !u32i
// CIR-NEXT: %[[AND:.*]] = cir.and %[[FEAT2_ELT_LOAD]], %[[MASK]] : !u32i
@@ -166,7 +166,7 @@ int v2() { return __builtin_cpu_supports("x86-64-v2"); }
// CIR-NEXT: %[[RETVAL:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!s32i>
// CIR-NEXT: %[[CPU_FEAT2:.*]] = cir.get_global @__cpu_features2 : !cir.ptr<!cir.array<!u32i x 3>>
// CIR-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !u32i
-// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element %[[CPU_FEAT2]][%[[TWO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
+// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element inbounds %[[CPU_FEAT2]][%[[TWO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
// CIR-NEXT: %[[FEAT2_ELT_LOAD:.*]] = cir.load {{.*}} %[[FEAT2_ELT]] : !cir.ptr<!u32i>, !u32i
// CIR-NEXT: %[[MASK:.*]] = cir.const #cir.int<2> : !u32i
// CIR-NEXT: %[[AND:.*]] = cir.and %[[FEAT2_ELT_LOAD]], %[[MASK]] : !u32i
@@ -200,7 +200,7 @@ int v3() { return __builtin_cpu_supports("x86-64-v3"); }
// CIR-NEXT: %[[RETVAL:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!s32i>
// CIR-NEXT: %[[CPU_FEAT2:.*]] = cir.get_global @__cpu_features2 : !cir.ptr<!cir.array<!u32i x 3>>
// CIR-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !u32i
-// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element %[[CPU_FEAT2]][%[[TWO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
+// CIR-NEXT: %[[FEAT2_ELT:.*]] = cir.get_element inbounds %[[CPU_FEAT2]][%[[TWO]] : !u32i] : !cir.ptr<!cir.array<!u32i x 3>> -> !cir.ptr<!u32i>
// CIR-NEXT: %[[FEAT2_ELT_LOAD:.*]] = cir.load {{.*}} %[[FEAT2_ELT]] : !cir.ptr<!u32i>, !u32i
// CIR-NEXT: %[[MASK:.*]] = cir.const #cir.int<4> : !u32i
// CIR-NEXT: %[[AND:.*]] = cir.and %[[FEAT2_ELT_LOAD]], %[[MASK]] : !u32i
diff --git a/clang/test/CIR/CodeGen/complex.cpp b/clang/test/CIR/CodeGen/complex.cpp
index cc027b1e2c4b89..407a6a8b0e6ea9 100644
--- a/clang/test/CIR/CodeGen/complex.cpp
+++ b/clang/test/CIR/CodeGen/complex.cpp
@@ -680,13 +680,13 @@ void foo24() {
// CIR: %[[ARR:.*]] = cir.alloca "arr" {{.*}} : !cir.ptr<!cir.array<!cir.complex<!s32i> x 2>>
// CIR: %[[RESULT:.*]] = cir.alloca "r" {{.*}} init : !cir.ptr<!cir.complex<!s32i>>
// CIR: %[[IDX:.*]] = cir.const #cir.int<1> : !s64i
-// CIR: %[[RESULT_VAL:.*]] = cir.get_element %[[ARR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!cir.complex<!s32i> x 2>> -> !cir.ptr<!cir.complex<!s32i>>
+// CIR: %[[RESULT_VAL:.*]] = cir.get_element inbounds %[[ARR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!cir.complex<!s32i> x 2>> -> !cir.ptr<!cir.complex<!s32i>>
// CIR: %[[TMP:.*]] = cir.load{{.*}} %[[RESULT_VAL]] : !cir.ptr<!cir.complex<!s32i>>, !cir.complex<!s32i>
// CIR: cir.store{{.*}} %[[TMP]], %[[RESULT]] : !cir.complex<!s32i>, !cir.ptr<!cir.complex<!s32i>>
// LLVM: %[[ARR:.*]] = alloca [2 x { i32, i32 }], align 16
// LLVM: %[[RESULT:.*]] = alloca { i32, i32 }, align 4
-// LLVM: %[[RESULT_VAL:.*]] = getelementptr [2 x { i32, i32 }], ptr %[[ARR]], i32 0, i64 1
+// LLVM: %[[RESULT_VAL:.*]] = getelementptr inbounds [2 x { i32, i32 }], ptr %[[ARR]], i32 0, i64 1
// LLVM: %[[TMP:.*]] = load { i32, i32 }, ptr %[[RESULT_VAL]], align 8
// LLVM: store { i32, i32 } %[[TMP]], ptr %[[RESULT]], align 4
diff --git a/clang/test/CIR/CodeGen/const-label-addr.c b/clang/test/CIR/CodeGen/const-label-addr.c
index 756569fe580206..e77275a77c85b9 100644
--- a/clang/test/CIR/CodeGen/const-label-addr.c
+++ b/clang/test/CIR/CodeGen/const-label-addr.c
@@ -71,14 +71,14 @@ void c(int x) {
// CIR: %[[C_TBL:.*]] = cir.get_global @c.tbl
// CIR: [[LABEL_A:.*]]:
// CIR: cir.label "A"
-// CIR: %[[P:.*]] = cir.get_element %[[C_TBL]][%{{.*}}]
+// CIR: %[[P:.*]] = cir.get_element inbounds %[[C_TBL]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.ptr<!void> x 3>> -> !cir.ptr<!cir.ptr<!void>>
// CIR: [[LABEL_B:.*]]:
// CIR: cir.label "B"
// LLVM: define dso_local void @c(i32 noundef %{{.*}})
// LLVM: br label %[[C_A]]
// LLVM: [[C_A]]:
-// LLVM: %[[TARGET:.*]] = getelementptr{{.*}} [3 x ptr], ptr @c.tbl
+// LLVM: %[[TARGET:.*]] = getelementptr inbounds [3 x ptr], ptr @c.tbl,
// LLVM: %[[P:.*]] = load ptr, ptr %[[TARGET]], align 8
// LLVM: br label %[[C_B]]
// LLVM: [[C_B]]:
diff --git a/clang/test/CIR/CodeGen/empty-union.cpp b/clang/test/CIR/CodeGen/empty-union.cpp
index a297cf3a34a033..39221b2d189833 100644
--- a/clang/test/CIR/CodeGen/empty-union.cpp
+++ b/clang/test/CIR/CodeGen/empty-union.cpp
@@ -183,7 +183,7 @@ int getArray() { return leadArr[1].x; }
// CIR: cir.func{{.*}} @_Z8getArrayv()
// CIR: %[[A:.*]] = cir.get_global @leadArr : !cir.ptr<!cir.array<!rec_Leading x 2>>
-// CIR: %[[E:.*]] = cir.get_element %[[A]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!rec_Leading x 2>> -> !cir.ptr<!rec_Leading>
+// CIR: %[[E:.*]] = cir.get_element inbounds %[[A]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!rec_Leading x 2>> -> !cir.ptr<!rec_Leading>
// CIR: %{{.*}} = cir.get_member %[[E]][1] {name = "x"} : !cir.ptr<!rec_Leading> -> !cir.ptr<!s32i>
// LLVM: define dso_local noundef i32 @_Z8getArrayv()
// LLVM: load i32, ptr getelementptr inbounds nuw (i8, ptr @leadArr, i64 12), align 4
diff --git a/clang/test/CIR/CodeGen/no-odr-use.cpp b/clang/test/CIR/CodeGen/no-odr-use.cpp
index 3b37316588a8b4..d7bb9790b93291 100644
--- a/clang/test/CIR/CodeGen/no-odr-use.cpp
+++ b/clang/test/CIR/CodeGen/no-odr-use.cpp
@@ -60,8 +60,8 @@ int f(int i) {
return (n >= 0
// CIR: %[[A:.*]] = cir.get_global @[[F_A]] : !cir.ptr<!rec_A>
// CIR: %[[ARR:.*]] = cir.get_member %[[A]][2] {name = "arr"} : !cir.ptr<!rec_A> -> !cir.ptr<!cir.array<!s32i x 3>>
- // CIR: cir.get_element %[[ARR]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 3>> -> !cir.ptr<!s32i>
- // LLVM: getelementptr [3 x i32], ptr getelementptr inbounds nuw (i8, ptr @[[F_A]], i64 12), i32 0, i64 %{{.*}}
+ // CIR: cir.get_element inbounds %[[ARR]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s32i x 3>> -> !cir.ptr<!s32i>
+ // LLVM: getelementptr inbounds [3 x i32], ptr getelementptr inbounds nuw (i8, ptr @[[F_A]], i64 12), i32 0, i64 %{{.*}}
// OGCG: getelementptr inbounds [3 x i32], ptr getelementptr inbounds nuw (i8, ptr @__const._Z1fi.a, i64 12), i64 0, i64 %{{.*}}
? a.arr[n]
// CIR: cir.ternary
@@ -85,9 +85,9 @@ int f(int i) {
// CIR: %[[N:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!s32i>, !s32i
// CIR: %[[SUB:.*]] = cir.sub nsw %[[TWO]], %[[N]] : !s32i
// CIR: %[[SUB_64:.*]] = cir.cast integral %[[SUB]] : !s32i -> !s64i
- // CIR: cir.get_element %[[Y]][%[[SUB_64]] : !s64i] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!s32i>
+ // CIR: cir.get_element inbounds %[[Y]][%[[SUB_64]] : !s64i] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!s32i>
- // LLVM: getelementptr [2 x i32], ptr getelementptr inbounds nuw ({{.*}} @[[F_A]], i64 4), i32 0, i64 %{{.*}}
+ // LLVM: getelementptr inbounds [2 x i32], ptr getelementptr inbounds nuw (i8, ptr @[[F_A]], i64 4), i32 0, i64 %{{.*}}
// LLVM: load i32
// OGCG: getelementptr inbounds [2 x i32], ptr getelementptr inbounds nuw (i8, ptr @__const._Z1fi.a, i64 4), i64 0, i64 %{{.*}}
@@ -114,10 +114,10 @@ namespace PR42276 {
for (int i = 0; i < sizeof(m) / sizeof(m[0]); ++i)
// CIR-CXX11: %[[M:.*]] = cir.get_global @_ZN7PR422765State1mE.const : !cir.ptr<!cir.array<!rec_anon_struct x 2>>
// CIR-CXX2A: %[[M:.*]] = cir.get_global @_ZN7PR422765State1mE : !cir.ptr<!cir.array<!rec_anon_struct x 2>>
- // CIR: cir.get_element %[[M]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!rec_anon_struct x 2>> -> !cir.ptr<!rec_anon_struct>
+ // CIR: cir.get_element inbounds %[[M]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!rec_anon_struct x 2>> -> !cir.ptr<!rec_anon_struct>
- // LLVM-CXX11: getelementptr [2 x { i64, i64 }], ptr @_ZN7PR422765State1mE.const, i32 0, i64 %{{.*}}
- // LLVM-CXX2A: getelementptr [2 x { i64, i64 }], ptr @_ZN7PR422765State1mE, i32 0, i64 %{{.*}}
+ // LLVM-CXX11: getelementptr inbounds [2 x { i64, i64 }], ptr @_ZN7PR422765State1mE.const, i32 0, i64 %{{.*}}
+ // LLVM-CXX2A: getelementptr inbounds [2 x { i64, i64 }], ptr @_ZN7PR422765State1mE, i32 0, i64 %{{.*}}
// OGCG-CXX11: getelementptr inbounds [2 x { i64, i64 }], ptr @_ZN7PR422765State1mE.const, i64 0, i64 %{{.*}}
// OGCG-CXX2A: getelementptr inbounds [2 x { i64, i64 }], ptr @_ZN7PR422765State1mE, i64 0, i64 %{{.*}}
(this->*m[i])();
@@ -126,9 +126,9 @@ namespace PR42276 {
for (int i = 0; i< sizeof(dms) / sizeof(dms[0]); ++i)
// CIR-CXX11: %[[DMS:.*]] = cir.get_global @_ZN7PR422765State3dmsE.const : !cir.ptr<!cir.array<!s64i x 2>>
// CIR-CXX2A: %[[DMS:.*]] = cir.get_global @_ZN7PR422765State3dmsE : !cir.ptr<!cir.array<!s64i x 2>>
- // CIR: cir.get_element %[[DMS]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s64i x 2>> -> !cir.ptr<!s64i>
- // LLVM-CXX11: getelementptr [2 x i64], ptr @_ZN7PR422765State3dmsE.const, i32 0, i64 %{{.*}}
- // LLVM-CXX2A: getelementptr [2 x i64], ptr @_ZN7PR422765State3dmsE, i32 0, i64 %{{.*}}
+ // CIR: cir.get_element inbounds %[[DMS]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s64i x 2>> -> !cir.ptr<!s64i>
+ // LLVM-CXX11: getelementptr inbounds [2 x i64], ptr @_ZN7PR422765State3dmsE.const, i32 0, i64 %{{.*}}
+ // LLVM-CXX2A: getelementptr inbounds [2 x i64], ptr @_ZN7PR422765State3dmsE, i32 0, i64 %{{.*}}
// OGCG-CXX11: getelementptr inbounds [2 x i64], ptr @_ZN7PR422765State3dmsE.const, i64 0, i64 %{{.*}}
// OGCG-CXX2A: getelementptr inbounds [2 x i64], ptr @_ZN7PR422765State3dmsE, i64 0, i64 %{{.*}}
sum += this->*dms[i];
diff --git a/clang/test/CIR/CodeGen/union.c b/clang/test/CIR/CodeGen/union.c
index ebb3e369e884c6..65937399a1e832 100644
--- a/clang/test/CIR/CodeGen/union.c
+++ b/clang/test/CIR/CodeGen/union.c
@@ -183,7 +183,7 @@ void f3(union U3 u) {
// CIR-NEXT: %[[ZERO:.*]] = cir.const #cir.int<0> : !s8i
// CIR-NEXT: %[[C_PTR:.*]] = cir.get_member %[[U]][0] {name = "c"} : !cir.ptr<!rec_U3> -> !cir.ptr<!cir.array<!s8i x 5>>
// CIR-NEXT: %[[IDX:.*]] = cir.const #cir.int<2> : !s64i
-// CIR-NEXT: %[[ELEM_PTR:.*]] = cir.get_element %[[C_PTR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s8i x 5>> -> !cir.ptr<!s8i>
+// CIR-NEXT: %[[ELEM_PTR:.*]] = cir.get_element inbounds %[[C_PTR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s8i x 5>> -> !cir.ptr<!s8i>
// CIR-NEXT: cir.store{{.*}} %[[ZERO]], %[[ELEM_PTR]] : !s8i, !cir.ptr<!s8i>
// CIR-NEXT: cir.return
@@ -192,7 +192,7 @@ void f3(union U3 u) {
// LLVM-NEXT: store i40 %[[ARG]], ptr %[[COERCE]], align 8
// LLVM-NEXT: %[[U:.*]] = alloca %union.U3, align 1
// LLVM-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 1 %[[U]], ptr align 8 %[[COERCE]], i64 5, i1 false)
-// LLVM-NEXT: %[[ELEM_PTR:.*]] = getelementptr [5 x i8], ptr %[[U]], i32 0, i64 2
+// LLVM-NEXT: %[[ELEM_PTR:.*]] = getelementptr inbounds [5 x i8], ptr %[[U]], i32 0, i64 2
// LLVM-NEXT: store i8 0, ptr %[[ELEM_PTR]], align 1
// LLVM-NEXT: ret void
@@ -217,7 +217,7 @@ void f5(union U4 u) {
// CIR-NEXT: %[[CHAR_VAL:.*]] = cir.const #cir.int<65> : !s8i
// CIR-NEXT: %[[C_PTR:.*]] = cir.get_member %[[U]][0] {name = "c"} : !cir.ptr<!rec_U4> -> !cir.ptr<!cir.array<!s8i x 5>>
// CIR-NEXT: %[[IDX:.*]] = cir.const #cir.int<4> : !s64i
-// CIR-NEXT: %[[ELEM_PTR:.*]] = cir.get_element %[[C_PTR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s8i x 5>> -> !cir.ptr<!s8i>
+// CIR-NEXT: %[[ELEM_PTR:.*]] = cir.get_element inbounds %[[C_PTR]][%[[IDX]] : !s64i] : !cir.ptr<!cir.array<!s8i x 5>> -> !cir.ptr<!s8i>
// CIR-NEXT: cir.store{{.*}} %[[CHAR_VAL]], %[[ELEM_PTR]] : !s8i, !cir.ptr<!s8i>
// CIR-NEXT: cir.return
@@ -226,7 +226,7 @@ void f5(union U4 u) {
// LLVM-NEXT: store i64 %[[ARG]], ptr %[[COERCE]], align 8
// LLVM-NEXT: %[[U:.*]] = alloca %union.U4, align 4
// LLVM-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %[[U]], ptr align 8 %[[COERCE]], i64 8, i1 false)
-// LLVM-NEXT: %[[ELEM_PTR:.*]] = getelementptr [5 x i8], ptr %[[U]], i32 0, i64 4
+// LLVM-NEXT: %[[ELEM_PTR:.*]] = getelementptr inbounds [5 x i8], ptr %[[U]], i32 0, i64 4
// LLVM-NEXT: store i8 65, ptr %[[ELEM_PTR]], align 4
// LLVM-NEXT: ret void
diff --git a/clang/test/CIR/CodeGenBuiltins/builtin-constant-p.c b/clang/test/CIR/CodeGenBuiltins/builtin-constant-p.c
index f87f7805f4aa31..6e559ba340dc5b 100644
--- a/clang/test/CIR/CodeGenBuiltins/builtin-constant-p.c
+++ b/clang/test/CIR/CodeGenBuiltins/builtin-constant-p.c
@@ -168,7 +168,7 @@ int test6(void) {
// CIR: cir.func {{.*}} @test6() -> !s32i
// CIR: %[[ARR:.*]] = cir.get_global @arr : !cir.ptr<!cir.array<!s32i x 3>>
// CIR: %[[TWO:.*]] = cir.const #cir.int<2> : !s64i
-// CIR: %[[ELE_PTR:.*]] = cir.get_element %[[ARR]][%[[TWO]] : !s64i] : !cir.ptr<!cir.array<!s32i x 3>> -> !cir.ptr<!s32i>
+// CIR: %[[ELE_PTR:.*]] = cir.get_element inbounds %[[ARR]][%[[TWO]] : !s64i] : !cir.ptr<!cir.array<!s32i x 3>> -> !cir.ptr<!s32i>
// CIR: %[[ELE:.*]] = cir.load{{.*}} %[[ELE_PTR]] : !cir.ptr<!s32i>, !s32i
// CIR: %[[IS_CONSTANT:.*]] = cir.is_constant %[[ELE]] : !s32i -> !cir.bool
diff --git a/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip b/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip
index 6c1ddeb3256245..f7859a649862a0 100644
--- a/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip
+++ b/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip
@@ -157,24 +157,24 @@ __device__ void func_name() { take(__func__); }
// CIR-SPV-LABEL: cir.func{{.*}} @_Z9subscripti
// CIR-SPV: %[[G:.*]] = cir.get_global @".str.2" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(1)>
// CIR-SPV: %[[C:.*]] = cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(1)> -> !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)>
-// CIR-SPV: cir.get_element %[[C]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)>
+// CIR-SPV: cir.get_element inbounds %[[C]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)>
// CIR-GCN-LABEL: cir.func{{.*}} @_Z9subscripti
// CIR-GCN: %[[G:.*]] = cir.get_global @".str.2" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)>
// CIR-GCN: %[[C:.*]] = cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!cir.array<!s8i x 4>>
-// CIR-GCN: cir.get_element %[[C]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>> -> !cir.ptr<!s8i>
+// CIR-GCN: cir.get_element inbounds %[[C]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s8i x 4>> -> !cir.ptr<!s8i>
// LLVM-SPV-LABEL: define{{.*}} i8 @_Z9subscripti
-// LLVM-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4))
+// LLVM-SPV: getelementptr inbounds [4 x i8], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4))
// OGCG-SPV-LABEL: define{{.*}} i8 @_Z9subscripti
-// OGCG-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4))
+// OGCG-SPV: getelementptr inbounds [4 x i8], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4))
// LLVM-GCN-LABEL: define{{.*}} i8 @_Z9subscripti
-// LLVM-GCN: getelementptr {{.*}}[4 x i8], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr)
+// LLVM-GCN: getelementptr inbounds [4 x i8], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr)
// OGCG-GCN-LABEL: define{{.*}} i8 @_Z9subscripti
-// OGCG-GCN: getelementptr {{.*}}[4 x i8], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr)
+// OGCG-GCN: getelementptr inbounds [4 x i8], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr)
__device__ char subscript(int i) { return "abc"[i]; }
// Binding a reference to a literal goes through constant emission and reuses
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c
index 828ee49b2a29ec..b910743f16e203 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c
+++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c
@@ -1088,7 +1088,7 @@ void copy_member_of_array_element_member() {
for(int i = 0; i < 5; ++i);
// CHECK-NEXT: %[[GETINNER:.*]] = cir.get_member %[[OUTER]][0] {name = "inner"} : !cir.ptr<!rec_OuterTy> -> !cir.ptr<!cir.array<!rec_InnerTy x 4>>
// CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s64i
- // CHECK-NEXT: %[[GET_ELT:.*]] = cir.get_element %[[GETINNER]][%[[TWO]] : !s64i] : !cir.ptr<!cir.array<!rec_InnerTy x 4>> -> !cir.ptr<!rec_InnerTy>
+ // CHECK-NEXT: %[[GET_ELT:.*]] = cir.get_element inbounds %[[GETINNER]][%[[TWO]] : !s64i] : !cir.ptr<!cir.array<!rec_InnerTy x 4>> -> !cir.ptr<!rec_InnerTy>
// CHECK-NEXT: %[[GETB:.*]] = cir.get_member %[[GET_ELT]][1] {name = "b"} : !cir.ptr<!rec_InnerTy> -> !cir.ptr<!s32i>
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETB]] : !cir.ptr<!s32i>) dataClause(acc_copy) name("outer.inner[2].b") -> !cir.ptr<!s32i>
// CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
diff --git a/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl b/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl
index c0d19ea60a98ef..5f3b7f7fe4276b 100644
--- a/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl
+++ b/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl
@@ -70,23 +70,23 @@ void call_take(void) { take("hi"); }
// CIR-SPV-LABEL: cir.func{{.*}} @subscript
// CIR-SPV: %[[G:.*]] = cir.get_global @".str.1" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(2)>
-// CIR-SPV: cir.get_element %[[G]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(2)> -> !cir.ptr<!s8i, target_address_space(2)>
+// CIR-SPV: cir.get_element inbounds %[[G]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(2)> -> !cir.ptr<!s8i, target_address_space(2)>
// CIR-GCN-LABEL: cir.func{{.*}} @subscript
// CIR-GCN: %[[G:.*]] = cir.get_global @".str.1" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)>
-// CIR-GCN: cir.get_element %[[G]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)>
+// CIR-GCN: cir.get_element inbounds %[[G]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)>
// LLVM-SPV-LABEL: define{{.*}} i8 @subscript
-// LLVM-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(2) @.str.1
+// LLVM-SPV: getelementptr inbounds [4 x i8], ptr addrspace(2) @.str.1
// OGCG-SPV-LABEL: define{{.*}} i8 @subscript
-// OGCG-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(2) @.str.1
+// OGCG-SPV: getelementptr inbounds [4 x i8], ptr addrspace(2) @.str.1
// LLVM-GCN-LABEL: define{{.*}} i8 @subscript
-// LLVM-GCN: getelementptr {{.*}}[4 x i8], ptr addrspace(4) @.str.1
+// LLVM-GCN: getelementptr inbounds [4 x i8], ptr addrspace(4) @.str.1
// OGCG-GCN-LABEL: define{{.*}} i8 @subscript
-// OGCG-GCN: getelementptr {{.*}}[4 x i8], ptr addrspace(4) @.str.1
+// OGCG-GCN: getelementptr inbounds [4 x i8], ptr addrspace(4) @.str.1
char subscript(int i) { return "abc"[i]; }
// CIR-SPV-LABEL: cir.func{{.*}} @func_name
diff --git a/clang/test/CIR/CodeGenOpenMP/device-array-subscript.c b/clang/test/CIR/CodeGenOpenMP/device-array-subscript.c
new file mode 100644
index 00000000000000..8adcf70f321801
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/device-array-subscript.c
@@ -0,0 +1,104 @@
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefixes=LLVM,LLVMCIR
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -emit-llvm %s -o %t.ll
+// RUN: FileCheck --input-file=%t.ll %s -check-prefixes=LLVM,OGCG
+
+#pragma omp declare target
+
+double param2d(double (*p)[5][10], int i, int j) { return (*p)[i][j]; }
+
+// CIR-LABEL: cir.func {{.*}} @param2d(
+// CIR: cir.get_element %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>> -> !cir.ptr<!cir.array<!cir.double x 10>>
+// CIR: cir.get_element inbounds %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.double x 10>> -> !cir.ptr<!cir.double>
+
+// LLVM-LABEL: define hidden double @param2d(
+// LLVM: %[[ROW:.*]] = getelementptr [5 x [10 x double]], ptr
+// LLVM: getelementptr inbounds [10 x double], ptr %[[ROW]],
+
+double uparam2d(double (*p)[5][10], unsigned i, int j) { return (*p)[i][j]; }
+
+// CIR-LABEL: cir.func {{.*}} @uparam2d(
+// CIR: cir.get_element nuw %{{.*}}[%{{.*}} : !u64i] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>> -> !cir.ptr<!cir.array<!cir.double x 10>>
+// CIR: cir.get_element inbounds %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.double x 10>> -> !cir.ptr<!cir.double>
+
+// LLVM-LABEL: define hidden double @uparam2d(
+// LLVM: %[[ROW:.*]] = getelementptr nuw [5 x [10 x double]], ptr
+// LLVM: getelementptr inbounds [10 x double], ptr %[[ROW]],
+
+double param3d(double (*p)[2][5][10], int i, int j, int k) {
+ return (*p)[i][j][k];
+}
+
+// CIR-LABEL: cir.func {{.*}} @param3d(
+// CIR: %[[PLANE:.*]] = cir.get_element %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 10> x 5> x 2>> -> !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>
+// CIR: %[[ROW:.*]] = cir.get_element %[[PLANE]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>> -> !cir.ptr<!cir.array<!cir.double x 10>>
+// CIR: cir.get_element inbounds %[[ROW]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.double x 10>> -> !cir.ptr<!cir.double>
+
+// LLVM-LABEL: define hidden double @param3d(
+// LLVM: %[[PLANE:.*]] = getelementptr [2 x [5 x [10 x double]]], ptr
+// LLVM: %[[ROW:.*]] = getelementptr [5 x [10 x double]], ptr %[[PLANE]],
+// LLVM: getelementptr inbounds [10 x double], ptr %[[ROW]],
+
+double local2d(int i, int j) {
+ double a[5][10] = {{0}};
+ return a[i][j];
+}
+
+// CIR-LABEL: cir.func {{.*}} @local2d(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>, target_address_space(5)>
+// CIR: %[[FLAT:.*]] = cir.cast address_space %[[A]] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>, target_address_space(5)> -> !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>
+// CIR: cir.get_element inbounds %[[FLAT]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>> -> !cir.ptr<!cir.array<!cir.double x 10>>
+// CIR: cir.get_element inbounds %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.double x 10>> -> !cir.ptr<!cir.double>
+
+// LLVM-LABEL: define hidden double @local2d(
+// LLVMCIR: alloca [5 x [10 x double]], align 8, addrspace(5)
+// LLVMCIR: %[[ROW:.*]] = getelementptr inbounds [5 x [10 x double]], ptr
+// OGCG: call align 8 ptr @__kmpc_alloc_shared(i64 400)
+// OGCG: %[[ROW:.*]] = getelementptr [5 x [10 x double]], ptr
+// LLVM: getelementptr inbounds [10 x double], ptr %[[ROW]],
+
+double local3d(int i, int j, int k) {
+ double a[2][5][10] = {{{0}}};
+ return a[i][j][k];
+}
+
+// CIR-LABEL: cir.func {{.*}} @local3d(
+// CIR: %[[PLANE:.*]] = cir.get_element inbounds %{{.*}}[%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 10> x 5> x 2>> -> !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>
+// CIR: %[[ROW:.*]] = cir.get_element inbounds %[[PLANE]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>> -> !cir.ptr<!cir.array<!cir.double x 10>>
+// CIR: cir.get_element inbounds %[[ROW]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.double x 10>> -> !cir.ptr<!cir.double>
+
+// LLVM-LABEL: define hidden double @local3d(
+// LLVMCIR: alloca [2 x [5 x [10 x double]]], align 8, addrspace(5)
+// LLVMCIR: %[[PLANE:.*]] = getelementptr inbounds [2 x [5 x [10 x double]]], ptr
+// LLVMCIR: %[[ROW:.*]] = getelementptr inbounds [5 x [10 x double]], ptr %[[PLANE]],
+// OGCG: call align 8 ptr @__kmpc_alloc_shared(i64 800)
+// OGCG: %[[PLANE:.*]] = getelementptr [2 x [5 x [10 x double]]], ptr
+// OGCG: %[[ROW:.*]] = getelementptr [5 x [10 x double]], ptr %[[PLANE]],
+// LLVM: getelementptr inbounds [10 x double], ptr %[[ROW]],
+
+double stride3d(int i, int j, int k) {
+ double a[2][5][10] = {{{0}}};
+ return (a + i)[0][j][k];
+}
+
+// CIR-LABEL: cir.func {{.*}} @stride3d(
+// CIR: %[[A:.*]] = cir.alloca "a" {{.*}} : !cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 10> x 5> x 2>, target_address_space(5)>
+// CIR: %[[FLAT:.*]] = cir.cast address_space %[[A]] : !cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 10> x 5> x 2>, target_address_space(5)> -> !cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 10> x 5> x 2>>
+// CIR: %[[DECAY:.*]] = cir.cast array_to_ptrdecay %[[FLAT]] : !cir.ptr<!cir.array<!cir.array<!cir.array<!cir.double x 10> x 5> x 2>> -> !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>
+// CIR: %[[ADD:.*]] = cir.ptr_stride %[[DECAY]], %{{.*}} : (!cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>, !s32i) -> !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>
+// CIR: %[[PLANE:.*]] = cir.ptr_stride %[[ADD]], %{{.*}} : (!cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>, !s64i) -> !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>>
+// CIR: cir.get_element inbounds %[[PLANE]][%{{.*}} : !s64i] : !cir.ptr<!cir.array<!cir.array<!cir.double x 10> x 5>> -> !cir.ptr<!cir.array<!cir.double x 10>>
+
+// LLVM-LABEL: define hidden double @stride3d(
+// LLVMCIR: alloca [2 x [5 x [10 x double]]], align 8, addrspace(5)
+// LLVMCIR: %[[ROW:.*]] = getelementptr inbounds [5 x [10 x double]], ptr
+// OGCG: call align 8 ptr @__kmpc_alloc_shared(i64 800)
+// OGCG: %[[ROW:.*]] = getelementptr [5 x [10 x double]], ptr %{{.*}}, i64 0, i64
+// LLVM: getelementptr inbounds [10 x double], ptr %[[ROW]],
+
+#pragma omp end declare target
diff --git a/clang/test/CIR/IR/array.cir b/clang/test/CIR/IR/array.cir
index 70c38d09cd8df5..fefcf6d43f0ad8 100644
--- a/clang/test/CIR/IR/array.cir
+++ b/clang/test/CIR/IR/array.cir
@@ -1,6 +1,8 @@
// RUN: cir-opt %s --verify-roundtrip | FileCheck %s
!s32i = !cir.int<s, 32>
+!s64i = !cir.int<s, 64>
+!u64i = !cir.int<u, 64>
module {
@@ -62,4 +64,20 @@ cir.func @func3(%arg0: !cir.ptr<!cir.array<!s32i x 10>>) {
// CHECK: cir.return
// CHECK: }
+cir.func @get_element(%arg0: !cir.ptr<!cir.array<!s32i x 10>>, %arg1: !s64i, %arg2: !u64i) {
+ %0 = cir.get_element %arg0[%arg1 : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+ %1 = cir.get_element inbounds %arg0[%arg1 : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+ %2 = cir.get_element nuw %arg0[%arg2 : !u64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+ %3 = cir.get_element inbounds nuw %arg0[%arg2 : !u64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+ cir.return
+}
+
+// CHECK: cir.func{{.*}} @get_element(%arg0: !cir.ptr<!cir.array<!s32i x 10>>, %arg1: !s64i, %arg2: !u64i) {
+// CHECK: %0 = cir.get_element %arg0[%arg1 : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CHECK: %1 = cir.get_element inbounds %arg0[%arg1 : !s64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CHECK: %2 = cir.get_element nuw %arg0[%arg2 : !u64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CHECK: %3 = cir.get_element inbounds nuw %arg0[%arg2 : !u64i] : !cir.ptr<!cir.array<!s32i x 10>> -> !cir.ptr<!s32i>
+// CHECK: cir.return
+// CHECK: }
+
}
diff --git a/clang/test/CIR/Lowering/array.cpp b/clang/test/CIR/Lowering/array.cpp
index 0225be03906c98..441be59415f5c9 100644
--- a/clang/test/CIR/Lowering/array.cpp
+++ b/clang/test/CIR/Lowering/array.cpp
@@ -51,10 +51,10 @@ void func() {
// CHECK-NEXT: %[[ARR_ALLOCA:.*]] = alloca [10 x i32], align 16
// CHECK-NEXT: %[[INIT:.*]] = alloca i32, align 4
// CHECK-NEXT: %[[INIT_2:.*]] = alloca i32, align 4
-// CHECK-NEXT: %[[ELE_PTR:.*]] = getelementptr [10 x i32], ptr %[[ARR_ALLOCA]], i32 0, i64 0
+// CHECK-NEXT: %[[ELE_PTR:.*]] = getelementptr inbounds [10 x i32], ptr %[[ARR_ALLOCA]], i32 0, i64 0
// CHECK-NEXT: %[[TMP:.*]] = load i32, ptr %[[ELE_PTR]], align 16
// CHECK-NEXT: store i32 %[[TMP]], ptr %[[INIT]], align 4
-// CHECK-NEXT: %[[ELE_PTR:.*]] = getelementptr [10 x i32], ptr %[[ARR_ALLOCA]], i32 0, i64 1
+// CHECK-NEXT: %[[ELE_PTR:.*]] = getelementptr inbounds [10 x i32], ptr %[[ARR_ALLOCA]], i32 0, i64 1
// CHECK-NEXT: %[[TMP:.*]] = load i32, ptr %[[ELE_PTR]], align 4
// CHECK-NEXT: store i32 %[[TMP]], ptr %[[INIT_2]], align 4
@@ -82,8 +82,8 @@ void func4() {
// CHECK: %[[ARR_ALLOCA:.*]] = alloca [2 x [1 x i32]], align 4
// CHECK: %[[INIT:.*]] = alloca i32, align 4
// CHECK: call void @llvm.memcpy{{.*}}(ptr align 4 %[[ARR_ALLOCA]], ptr align 4 @[[FUNC4_ARR]], i64 8, i1 false)
-// CHECK: %[[ARR_1:.*]] = getelementptr [2 x [1 x i32]], ptr %[[ARR_ALLOCA]], i32 0, i64 1
-// CHECK: %[[ELE_PTR:.*]] = getelementptr [1 x i32], ptr %[[ARR_1]], i32 0, i64 0
+// CHECK: %[[ARR_1:.*]] = getelementptr inbounds [2 x [1 x i32]], ptr %[[ARR_ALLOCA]], i32 0, i64 1
+// CHECK: %[[ELE_PTR:.*]] = getelementptr inbounds [1 x i32], ptr %[[ARR_1]], i32 0, i64 0
// CHECK: %[[TMP:.*]] = load i32, ptr %[[ELE_PTR]], align 4
// CHECK: store i32 %[[TMP]], ptr %[[INIT]], align 4
diff --git a/clang/test/CIR/Lowering/get-element.cir b/clang/test/CIR/Lowering/get-element.cir
new file mode 100644
index 00000000000000..7d3b34334be5d2
--- /dev/null
+++ b/clang/test/CIR/Lowering/get-element.cir
@@ -0,0 +1,40 @@
+// RUN: cir-translate -cir-to-llvmir --disable-cc-lowering -o %t.ll %s
+// RUN: FileCheck --input-file=%t.ll -check-prefix=LLVM %s
+
+!s32i = !cir.int<s, 32>
+!s64i = !cir.int<s, 64>
+!u64i = !cir.int<u, 64>
+
+module attributes {cir.triple = "x86_64-unknown-linux-gnu"} {
+ cir.func @plain(%arg0: !cir.ptr<!cir.array<!s32i x 4>>, %arg1: !s64i) -> !cir.ptr<!s32i> {
+ %0 = cir.get_element %arg0[%arg1 : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+ cir.return %0 : !cir.ptr<!s32i>
+ }
+
+ cir.func @inbounds(%arg0: !cir.ptr<!cir.array<!s32i x 4>>, %arg1: !s64i) -> !cir.ptr<!s32i> {
+ %0 = cir.get_element inbounds %arg0[%arg1 : !s64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+ cir.return %0 : !cir.ptr<!s32i>
+ }
+
+ cir.func @nuw(%arg0: !cir.ptr<!cir.array<!s32i x 4>>, %arg1: !u64i) -> !cir.ptr<!s32i> {
+ %0 = cir.get_element nuw %arg0[%arg1 : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+ cir.return %0 : !cir.ptr<!s32i>
+ }
+
+ cir.func @inbounds_nuw(%arg0: !cir.ptr<!cir.array<!s32i x 4>>, %arg1: !u64i) -> !cir.ptr<!s32i> {
+ %0 = cir.get_element inbounds nuw %arg0[%arg1 : !u64i] : !cir.ptr<!cir.array<!s32i x 4>> -> !cir.ptr<!s32i>
+ cir.return %0 : !cir.ptr<!s32i>
+ }
+}
+
+// LLVM-LABEL: define ptr @plain(
+// LLVM: getelementptr [4 x i32], ptr %{{.*}}, i32 0, i64 %{{.*}}
+
+// LLVM-LABEL: define ptr @inbounds(
+// LLVM: getelementptr inbounds [4 x i32], ptr %{{.*}}, i32 0, i64 %{{.*}}
+
+// LLVM-LABEL: define ptr @nuw(
+// LLVM: getelementptr nuw [4 x i32], ptr %{{.*}}, i32 0, i64 %{{.*}}
+
+// LLVM-LABEL: define ptr @inbounds_nuw(
+// LLVM: getelementptr inbounds nuw [4 x i32], ptr %{{.*}}, i32 0, i64 %{{.*}}
More information about the cfe-commits
mailing list