[Mlir-commits] [mlir] 202ece6 - [mlir][spirv] Add in-bounds access chain conversion (#216984)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Aug 19 05:59:06 PDT 2026
Author: Hsiangkai Wang
Date: 2026-08-19T13:59:01+01:00
New Revision: 202ece62657cf3e795de3a954e10c79bd44be372
URL: https://github.com/llvm/llvm-project/commit/202ece62657cf3e795de3a954e10c79bd44be372
DIFF: https://github.com/llvm/llvm-project/commit/202ece62657cf3e795de3a954e10c79bd44be372.diff
LOG: [mlir][spirv] Add in-bounds access chain conversion (#216984)
Use spirv.InBoundsAccessChain for static StorageBuffer accesses whose
linearized range fits the declared SPIR-V object.
Keep plain access chains for dynamic layouts and packed sub-16-bit
storage.
Added:
Modified:
mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
mlir/test/Conversion/ConvertToSPIRV/argmax-kernel.mlir
mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules.mlir
mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
mlir/test/Conversion/ConvertToSPIRV/memref.mlir
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
mlir/test/Conversion/MemRefToSPIRV/atomic.mlir
mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir
mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir
mlir/test/Conversion/SCFToSPIRV/for.mlir
mlir/test/Conversion/SCFToSPIRV/if.mlir
mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h b/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
index 577fa2c44496b..acb41b74384dc 100644
--- a/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
+++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
@@ -198,6 +198,13 @@ Value getElementPtr(const SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr, ValueRange indices,
Location loc, OpBuilder &builder);
+/// As above, with the number of contiguous memref elements accessed through
+/// the pointer. This lets vector conversions retain their full access range.
+Value getElementPtr(const SPIRVTypeConverter &typeConverter,
+ MemRefType baseType, Value basePtr, ValueRange indices,
+ Location loc, OpBuilder &builder,
+ uint64_t accessElementCount);
+
// GetElementPtr implementation for Kernel/OpenCL flavored SPIR-V.
Value getOpenCLElementPtr(const SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr,
@@ -208,6 +215,13 @@ Value getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr,
ValueRange indices, Location loc, OpBuilder &builder);
+/// As above, with the number of contiguous memref elements accessed through
+/// the pointer.
+Value getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
+ MemRefType baseType, Value basePtr,
+ ValueRange indices, Location loc, OpBuilder &builder,
+ uint64_t accessElementCount);
+
// Find the largest factor of size among {2,3,4} for the lowest dimension of
// the target shape.
int getComputeVectorSize(int64_t size);
diff --git a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
index 78693e924c4d9..0808d13620be0 100644
--- a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
+++ b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
@@ -743,9 +743,9 @@ struct VectorLoadOpConverter final
const auto &typeConverter = *getTypeConverter<SPIRVTypeConverter>();
auto loc = loadOp.getLoc();
- Value accessChain =
- spirv::getElementPtr(typeConverter, memrefType, adaptor.getBase(),
- adaptor.getIndices(), loc, rewriter);
+ Value accessChain = spirv::getElementPtr(
+ typeConverter, memrefType, adaptor.getBase(), adaptor.getIndices(), loc,
+ rewriter, loadOp.getVectorType().getNumElements());
if (!accessChain)
return rewriter.notifyMatchFailure(
loadOp, "failed to get memref element pointer");
@@ -809,9 +809,9 @@ struct VectorStoreOpConverter final
const auto &typeConverter = *getTypeConverter<SPIRVTypeConverter>();
auto loc = storeOp.getLoc();
- Value accessChain =
- spirv::getElementPtr(typeConverter, memrefType, adaptor.getBase(),
- adaptor.getIndices(), loc, rewriter);
+ Value accessChain = spirv::getElementPtr(
+ typeConverter, memrefType, adaptor.getBase(), adaptor.getIndices(), loc,
+ rewriter, storeOp.getVectorType().getNumElements());
if (!accessChain)
return rewriter.notifyMatchFailure(
storeOp, "failed to get memref element pointer");
diff --git a/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp b/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
index 1c8a22fb35639..c320245408d67 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
@@ -1316,6 +1316,54 @@ static std::optional<uint64_t> getMaxLinearizedIndex(ArrayRef<int64_t> shape,
return maxLinearIndex;
}
+static std::optional<uint64_t> getStorageBufferElementCount(Value basePtr) {
+ auto pointerType = dyn_cast<spirv::PointerType>(basePtr.getType());
+ if (!pointerType ||
+ pointerType.getStorageClass() != spirv::StorageClass::StorageBuffer)
+ return std::nullopt;
+
+ Type pointeeType = pointerType.getPointeeType();
+ if (auto structType = dyn_cast<spirv::StructType>(pointeeType)) {
+ if (structType.getNumElements() != 1)
+ return std::nullopt;
+ pointeeType = structType.getElementType(0);
+ }
+ auto arrayType = dyn_cast<spirv::ArrayType>(pointeeType);
+ if (!arrayType)
+ return std::nullopt;
+ return arrayType.getNumElements();
+}
+
+static bool shouldEmitInBoundsAccessChain(MemRefType baseType, Value basePtr,
+ ArrayRef<int64_t> strides,
+ int64_t offset,
+ uint64_t accessElementCount) {
+ // Sub-16-bit integer memrefs may be stored using a wider SPIR-V array element
+ // than the source element. Keep a plain access chain so later bitwidth
+ // emulation can adjust the final index in storage-element units.
+ if (auto integerType = dyn_cast<IntegerType>(baseType.getElementType()))
+ if (integerType.getWidth() < 16)
+ return false;
+
+ std::optional<uint64_t> maxSourceElementIndex =
+ getMaxLinearizedIndex(baseType.getShape(), strides, offset);
+ std::optional<uint64_t> storageElementCount =
+ getStorageBufferElementCount(basePtr);
+ if (!maxSourceElementIndex || !storageElementCount)
+ return false;
+
+ if (accessElementCount == 0 || accessElementCount > *storageElementCount)
+ return false;
+
+ // `InBoundsAccessChain` requires the computed pointer to stay within the
+ // SPIR-V base object. Dynamic index validity is assumed from the source
+ // operation/caller contract; for vector accesses, `accessElementCount` only
+ // rejects widths that cannot fit in the fixed StorageBuffer object at all.
+ // The static proof here is that the memref layout's linear index space maps
+ // into that same object.
+ return *maxSourceElementIndex < *storageElementCount;
+}
+
} // namespace
//===----------------------------------------------------------------------===//
@@ -1431,7 +1479,8 @@ Value mlir::spirv::linearizeIndex(ValueRange indices, ArrayRef<int64_t> strides,
Value mlir::spirv::getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr,
ValueRange indices, Location loc,
- OpBuilder &builder) {
+ OpBuilder &builder,
+ uint64_t accessElementCount) {
// Get base and offset of the MemRefType and verify they are static.
int64_t offset;
@@ -1462,9 +1511,21 @@ Value mlir::spirv::getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
// Interface memrefs are wrapped in a struct: index to its first elem.
if (isa<spirv::StructType>(pointeeType))
linearizedIndices.insert(linearizedIndices.begin(), zero);
+ if (shouldEmitInBoundsAccessChain(baseType, basePtr, strides, offset,
+ accessElementCount))
+ return spirv::InBoundsAccessChainOp::create(builder, loc, basePtr,
+ linearizedIndices);
return spirv::AccessChainOp::create(builder, loc, basePtr, linearizedIndices);
}
+Value mlir::spirv::getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
+ MemRefType baseType, Value basePtr,
+ ValueRange indices, Location loc,
+ OpBuilder &builder) {
+ return getVulkanElementPtr(typeConverter, baseType, basePtr, indices, loc,
+ builder, /*accessElementCount=*/1);
+}
+
Value mlir::spirv::getOpenCLElementPtr(const SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr,
ValueRange indices, Location loc,
@@ -1506,7 +1567,8 @@ Value mlir::spirv::getOpenCLElementPtr(const SPIRVTypeConverter &typeConverter,
Value mlir::spirv::getElementPtr(const SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr,
ValueRange indices, Location loc,
- OpBuilder &builder) {
+ OpBuilder &builder,
+ uint64_t accessElementCount) {
if (typeConverter.allows(spirv::Capability::Kernel)) {
return getOpenCLElementPtr(typeConverter, baseType, basePtr, indices, loc,
@@ -1514,7 +1576,15 @@ Value mlir::spirv::getElementPtr(const SPIRVTypeConverter &typeConverter,
}
return getVulkanElementPtr(typeConverter, baseType, basePtr, indices, loc,
- builder);
+ builder, accessElementCount);
+}
+
+Value mlir::spirv::getElementPtr(const SPIRVTypeConverter &typeConverter,
+ MemRefType baseType, Value basePtr,
+ ValueRange indices, Location loc,
+ OpBuilder &builder) {
+ return getElementPtr(typeConverter, baseType, basePtr, indices, loc, builder,
+ /*accessElementCount=*/1);
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/test/Conversion/ConvertToSPIRV/argmax-kernel.mlir b/mlir/test/Conversion/ConvertToSPIRV/argmax-kernel.mlir
index 81ac2a04f3362..c300987c31543 100644
--- a/mlir/test/Conversion/ConvertToSPIRV/argmax-kernel.mlir
+++ b/mlir/test/Conversion/ConvertToSPIRV/argmax-kernel.mlir
@@ -18,7 +18,7 @@ module attributes {
// CHECK: %[[ADDRESSLOCALINVOCATIONID:.*]] = spirv.mlir.addressof @[[$LOCALINVOCATIONIDVAR]]
// CHECK: %[[LOCALINVOCATIONID:.*]] = spirv.Load "Input" %[[ADDRESSLOCALINVOCATIONID]]
// CHECK: %[[LOCALINVOCATIONIDX:.*]] = spirv.CompositeExtract %[[LOCALINVOCATIONID]]{{\[}}0 : i32{{\]}}
- // CHECK: %[[AC0:.*]] = spirv.AccessChain %[[ARG0]][%[[C0]], %[[LOCALINVOCATIONIDX]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+ // CHECK: %[[AC0:.*]] = spirv.InBoundsAccessChain %[[ARG0]][%[[C0]], %[[LOCALINVOCATIONIDX]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[LOAD0:.*]] = spirv.Load "StorageBuffer" %[[AC0]] : f32
// CHECK: %[[FUNC0:.*]] = spirv.Variable : !spirv.ptr<i32, Function>
// CHECK: %[[FUNC1:.*]] = spirv.Variable : !spirv.ptr<f32, Function>
@@ -41,7 +41,7 @@ module attributes {
// CHECK: ^[[BODY]]:
// CHECK: %[[MUL:.*]] = spirv.IMul %[[INDVAR0]], %[[C32]] : i32
// CHECK: %[[ADD:.*]] = spirv.IAdd %[[MUL]], %[[LOCALINVOCATIONIDX]] : i32
- // CHECK: %[[AC1:.*]] = spirv.AccessChain %[[ARG0]][%[[C0]], %[[ADD]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+ // CHECK: %[[AC1:.*]] = spirv.InBoundsAccessChain %[[ARG0]][%[[C0]], %[[ADD]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[LOAD1:.*]] = spirv.Load "StorageBuffer" %[[AC1]] : f32
// CHECK: %[[OGT:.*]] = spirv.FOrdGreaterThan %[[LOAD1]], %[[INDVAR2]] : f32
// CHECK: %[[SELECT0:.*]] = spirv.Select %[[OGT]], %[[ADD]], %[[INDVAR1]] : i1, i32
@@ -82,7 +82,7 @@ module attributes {
// CHECK: spirv.mlir.selection {
// CHECK: spirv.BranchConditional %[[EQ]], ^[[TRUE:.*]], ^[[FALSE:.*]]
// CHECK: ^[[TRUE]]:
- // CHECK: %[[AC2:.*]] = spirv.AccessChain %[[ARG1]][%[[C0]], %[[C0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
+ // CHECK: %[[AC2:.*]] = spirv.InBoundsAccessChain %[[ARG1]][%[[C0]], %[[C0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: spirv.Store "StorageBuffer" %[[AC2]], %[[LANE_RES]] : i32
// CHECK: spirv.Branch ^[[FALSE]]
// CHECK: ^[[FALSE]]:
diff --git a/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules.mlir b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules.mlir
index 96ad107d35817..eb020f82d50df 100644
--- a/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules.mlir
+++ b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules.mlir
@@ -81,7 +81,7 @@ module attributes {
// CHECK: spirv.Constant dense<0>
%idx0 = arith.constant 0 : index
%vec0 = arith.constant dense<[0, 0]> : vector<2xi32>
- // CHECK: spirv.AccessChain
+ // CHECK: spirv.InBoundsAccessChain
// CHECK: spirv.Load "StorageBuffer"
%val = memref.load %arg0[%idx0] : memref<2xi32>
// CHECK: spirv.CompositeInsert
@@ -89,7 +89,7 @@ module attributes {
// CHECK: spirv.VectorShuffle
%shuffle = vector.shuffle %vec, %vec[3, 2, 1, 0] : vector<2xi32>, vector<2xi32>
%res = vector.extract %shuffle[%idx0] : i32 from vector<4xi32>
- // CHECK: spirv.AccessChain
+ // CHECK: spirv.InBoundsAccessChain
// CHECK: spirv.Store "StorageBuffer"
memref.store %res, %arg1[%idx0]: memref<4xi32>
// CHECK: spirv.Return
diff --git a/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir b/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
index 2b5c82e5bd877..05493240824bd 100644
--- a/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
+++ b/mlir/test/Conversion/ConvertToSPIRV/gpu.mlir
@@ -55,12 +55,12 @@ module attributes {
gpu.module @kernels {
gpu.func @load_store(%arg0: memref<12x4xf32, #spirv.storage_class<StorageBuffer>>, %arg1: memref<12x4xf32, #spirv.storage_class<StorageBuffer>>, %arg2: memref<12x4xf32, #spirv.storage_class<StorageBuffer>>, %arg3: index, %arg4: index, %arg5: index, %arg6: index) kernel
attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
- // CHECK: %[[PTR1:.*]] = spirv.AccessChain %[[ARG0]]
+ // CHECK: %[[PTR1:.*]] = spirv.InBoundsAccessChain %[[ARG0]]
// CHECK-NEXT: spirv.Load "StorageBuffer" %[[PTR1]]
- // CHECK: %[[PTR2:.*]] = spirv.AccessChain %[[ARG1]]
+ // CHECK: %[[PTR2:.*]] = spirv.InBoundsAccessChain %[[ARG1]]
// CHECK-NEXT: spirv.Load "StorageBuffer" %[[PTR2]]
// CHECK: spirv.FAdd
- // CHECK: %[[PTR3:.*]] = spirv.AccessChain %[[ARG2]]
+ // CHECK: %[[PTR3:.*]] = spirv.InBoundsAccessChain %[[ARG2]]
// CHECK-NEXT: spirv.Store "StorageBuffer" %[[PTR3]]
%0 = gpu.block_id x
%1 = gpu.block_id y
diff --git a/mlir/test/Conversion/ConvertToSPIRV/memref.mlir b/mlir/test/Conversion/ConvertToSPIRV/memref.mlir
index f7d17d1ad16da..b8724ae8b2f5b 100644
--- a/mlir/test/Conversion/ConvertToSPIRV/memref.mlir
+++ b/mlir/test/Conversion/ConvertToSPIRV/memref.mlir
@@ -8,9 +8,9 @@ module attributes {
// CHECK-LABEL: @load_store_float_rank_zero
// CHECK-SAME: %[[ARG0:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>, %[[ARG1:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: %[[CST0:.*]] = spirv.Constant 0 : i32
-// CHECK: %[[AC0:.*]] = spirv.AccessChain %[[ARG0]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC0:.*]] = spirv.InBoundsAccessChain %[[ARG0]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[LOAD:.*]] = spirv.Load "StorageBuffer" %[[AC0]] : f32
-// CHECK: %[[AC1:.*]] = spirv.AccessChain %[[ARG1]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC1:.*]] = spirv.InBoundsAccessChain %[[ARG1]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: spirv.Store "StorageBuffer" %[[AC1]], %[[LOAD]] : f32
// CHECK: spirv.Return
func.func @load_store_float_rank_zero(%arg0: memref<f32>, %arg1: memref<f32>) {
@@ -22,9 +22,9 @@ func.func @load_store_float_rank_zero(%arg0: memref<f32>, %arg1: memref<f32>) {
// CHECK-LABEL: @load_store_int_rank_one
// CHECK-SAME: %[[ARG0:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer>, %[[ARG1:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer>, %[[ARG2:.*]]: i32
// CHECK: %[[CST0:.*]] = spirv.Constant 0 : i32
-// CHECK: %[[AC0:.*]] = spirv.AccessChain %[[ARG0]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC0:.*]] = spirv.InBoundsAccessChain %[[ARG0]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[LOAD:.*]] = spirv.Load "StorageBuffer" %[[AC0]] : i32
-// CHECK: %[[AC1:.*]] = spirv.AccessChain %[[ARG1]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC1:.*]] = spirv.InBoundsAccessChain %[[ARG1]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: spirv.Store "StorageBuffer" %[[AC1]], %[[LOAD]] : i32
// CHECK: spirv.Return
func.func @load_store_int_rank_one(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2 : index) {
@@ -36,9 +36,9 @@ func.func @load_store_int_rank_one(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %
// CHECK-LABEL: @load_store_larger_memref
// CHECK-SAME: %[[ARG0:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<8 x i32, stride=4> [0])>, StorageBuffer>, %[[ARG1:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<8 x i32, stride=4> [0])>, StorageBuffer>, %[[ARG2:.*]]: i32
// CHECK: %[[CST0:.*]] = spirv.Constant 0 : i32
-// CHECK: %[[AC0:.*]] = spirv.AccessChain %[[ARG0]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<8 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC0:.*]] = spirv.InBoundsAccessChain %[[ARG0]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<8 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[LOAD:.*]] = spirv.Load "StorageBuffer" %[[AC0]] : i32
-// CHECK: %[[AC1:.*]] = spirv.AccessChain %[[ARG1]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<8 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC1:.*]] = spirv.InBoundsAccessChain %[[ARG1]][%[[CST0]], %[[ARG2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<8 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: spirv.Store "StorageBuffer" %[[AC1]], %[[LOAD]] : i32
// CHECK: spirv.Return
func.func @load_store_larger_memref(%arg0: memref<8xi32>, %arg1: memref<8xi32>, %arg2 : index) {
@@ -51,9 +51,9 @@ func.func @load_store_larger_memref(%arg0: memref<8xi32>, %arg1: memref<8xi32>,
// CHECK-LABEL: @load_store_vector
// CHECK-SAME: %[[ARG0:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<1 x vector<4xi32>, stride=16> [0])>, StorageBuffer>, %[[ARG1:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<1 x vector<4xi32>, stride=16> [0])>, StorageBuffer>
// CHECK: %[[CST0:.*]] = spirv.Constant 0 : i32
-// CHECK: %[[AC0:.*]] = spirv.AccessChain %[[ARG0]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x vector<4xi32>, stride=16> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC0:.*]] = spirv.InBoundsAccessChain %[[ARG0]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x vector<4xi32>, stride=16> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[LOAD:.*]] = spirv.Load "StorageBuffer" %[[AC0]] : vector<4xi32>
-// CHECK: %[[AC1:.*]] = spirv.AccessChain %[[ARG1]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x vector<4xi32>, stride=16> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[AC1:.*]] = spirv.InBoundsAccessChain %[[ARG1]][%[[CST0]], %[[CST0]]] : !spirv.ptr<!spirv.struct<(!spirv.array<1 x vector<4xi32>, stride=16> [0])>, StorageBuffer>, i32, i32
// CHECK: spirv.Store "StorageBuffer" %[[AC1]], %[[LOAD]] : vector<4xi32>
// CHECK: spirv.Return
func.func @load_store_vector(%arg0: memref<vector<4xi32>>, %arg1: memref<vector<4xi32>>) {
diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
index 4339799ccd5ea..9861acf429eb4 100644
--- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
@@ -63,15 +63,15 @@ module attributes {
// CHECK: %[[STRIDE1_1:.*]] = spirv.Constant 4 : i32
// CHECK: %[[UPDATE1_1:.*]] = spirv.IMul %[[INDEX1]], %[[STRIDE1_1]] : i32
// CHECK: %[[OFFSET1_2:.*]] = spirv.IAdd %[[INDEX2]], %[[UPDATE1_1]] : i32
- // CHECK: %[[PTR1:.*]] = spirv.AccessChain %[[ARG0]]{{\[}}%[[ZERO]], %[[OFFSET1_2]]{{\]}}
+ // CHECK: %[[PTR1:.*]] = spirv.InBoundsAccessChain %[[ARG0]]{{\[}}%[[ZERO]], %[[OFFSET1_2]]{{\]}}
// CHECK-NEXT: %[[VAL1:.*]] = spirv.Load "StorageBuffer" %[[PTR1]]
%14 = memref.load %arg0[%12, %13] : memref<12x4xf32, #spirv.storage_class<StorageBuffer>>
- // CHECK: %[[PTR2:.*]] = spirv.AccessChain %[[ARG1]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
+ // CHECK: %[[PTR2:.*]] = spirv.InBoundsAccessChain %[[ARG1]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
// CHECK-NEXT: %[[VAL2:.*]] = spirv.Load "StorageBuffer" %[[PTR2]]
%15 = memref.load %arg1[%12, %13] : memref<12x4xf32, #spirv.storage_class<StorageBuffer>>
// CHECK: %[[VAL3:.*]] = spirv.FAdd %[[VAL1]], %[[VAL2]]
%16 = arith.addf %14, %15 : f32
- // CHECK: %[[PTR3:.*]] = spirv.AccessChain %[[ARG2]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
+ // CHECK: %[[PTR3:.*]] = spirv.InBoundsAccessChain %[[ARG2]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
// CHECK-NEXT: spirv.Store "StorageBuffer" %[[PTR3]], %[[VAL3]]
memref.store %16, %arg2[%12, %13] : memref<12x4xf32, #spirv.storage_class<StorageBuffer>>
gpu.return
diff --git a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
index 983747be57995..0e7509f694eab 100644
--- a/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/lookup-target-env.mlir
@@ -7,7 +7,7 @@ module attributes {gpu.container_module} {
// CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
%c0 = arith.constant 0 : index
- // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
+ // CHECK: %[[PTR:.*]] = spirv.InBoundsAccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
// CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f32
%0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>
// CHECK: spirv.Return
@@ -30,7 +30,7 @@ module attributes {gpu.container_module} {
// CHECK-SAME: %[[ARG:.*]]: !spirv.ptr<!spirv.struct<(!spirv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>})
gpu.func @load_kernel(%arg0: memref<12x4xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
%c0 = arith.constant 0 : index
- // CHECK: %[[PTR:.*]] = spirv.AccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
+ // CHECK: %[[PTR:.*]] = spirv.InBoundsAccessChain %[[ARG]]{{\[}}{{%.*}}, {{%.*}}{{\]}}
// CHECK-NEXT: {{%.*}} = spirv.Load "StorageBuffer" %[[PTR]] : f32
%0 = memref.load %arg0[%c0, %c0] : memref<12x4xf32>
// CHECK: spirv.Return
diff --git a/mlir/test/Conversion/MemRefToSPIRV/atomic.mlir b/mlir/test/Conversion/MemRefToSPIRV/atomic.mlir
index fa416512aa144..d37aa302933f8 100644
--- a/mlir/test/Conversion/MemRefToSPIRV/atomic.mlir
+++ b/mlir/test/Conversion/MemRefToSPIRV/atomic.mlir
@@ -5,7 +5,7 @@ module attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader
// CHECK: func.func @atomic_addi_storage_buffer
// CHECK-SAME: (%[[VAL:.+]]: i32,
func.func @atomic_addi_storage_buffer(%value: i32, %memref: memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>, %i0: index, %i1: index, %i2: index) -> i32 {
- // CHECK: %[[AC:.+]] = spirv.AccessChain
+ // CHECK: %[[AC:.+]] = spirv.InBoundsAccessChain
// CHECK: %[[ATOMIC:.+]] = spirv.AtomicIAdd <Device> <AcquireRelease|UniformMemory> %[[AC]], %[[VAL]] : !spirv.ptr<i32, StorageBuffer>
// CHECK: return %[[ATOMIC]]
%0 = memref.atomic_rmw "addi" %value, %memref[%i0, %i1, %i2] : (i32, memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>) -> i32
@@ -25,7 +25,7 @@ func.func @atomic_maxs_workgroup(%value: i32, %memref: memref<2x3x4xi32, #spirv.
// CHECK: func.func @atomic_maxu_storage_buffer
// CHECK-SAME: (%[[VAL:.+]]: i32,
func.func @atomic_maxu_storage_buffer(%value: i32, %memref: memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>, %i0: index, %i1: index, %i2: index) -> i32 {
- // CHECK: %[[AC:.+]] = spirv.AccessChain
+ // CHECK: %[[AC:.+]] = spirv.InBoundsAccessChain
// CHECK: %[[ATOMIC:.+]] = spirv.AtomicUMax <Device> <AcquireRelease|UniformMemory> %[[AC]], %[[VAL]] : !spirv.ptr<i32, StorageBuffer>
// CHECK: return %[[ATOMIC]]
%0 = memref.atomic_rmw "maxu" %value, %memref[%i0, %i1, %i2] : (i32, memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>) -> i32
@@ -45,7 +45,7 @@ func.func @atomic_mins_workgroup(%value: i32, %memref: memref<2x3x4xi32, #spirv.
// CHECK: func.func @atomic_minu_storage_buffer
// CHECK-SAME: (%[[VAL:.+]]: i32,
func.func @atomic_minu_storage_buffer(%value: i32, %memref: memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>, %i0: index, %i1: index, %i2: index) -> i32 {
- // CHECK: %[[AC:.+]] = spirv.AccessChain
+ // CHECK: %[[AC:.+]] = spirv.InBoundsAccessChain
// CHECK: %[[ATOMIC:.+]] = spirv.AtomicUMin <Device> <AcquireRelease|UniformMemory> %[[AC]], %[[VAL]] : !spirv.ptr<i32, StorageBuffer>
// CHECK: return %[[ATOMIC]]
%0 = memref.atomic_rmw "minu" %value, %memref[%i0, %i1, %i2] : (i32, memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>) -> i32
@@ -65,7 +65,7 @@ func.func @atomic_ori_workgroup(%value: i32, %memref: memref<2x3x4xi32, #spirv.s
// CHECK: func.func @atomic_andi_storage_buffer
// CHECK-SAME: (%[[VAL:.+]]: i32,
func.func @atomic_andi_storage_buffer(%value: i32, %memref: memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>, %i0: index, %i1: index, %i2: index) -> i32 {
- // CHECK: %[[AC:.+]] = spirv.AccessChain
+ // CHECK: %[[AC:.+]] = spirv.InBoundsAccessChain
// CHECK: %[[ATOMIC:.+]] = spirv.AtomicAnd <Device> <AcquireRelease|UniformMemory> %[[AC]], %[[VAL]] : !spirv.ptr<i32, StorageBuffer>
// CHECK: return %[[ATOMIC]]
%0 = memref.atomic_rmw "andi" %value, %memref[%i0, %i1, %i2] : (i32, memref<2x3x4xi32, #spirv.storage_class<StorageBuffer>>) -> i32
@@ -156,7 +156,7 @@ module attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader
// CHECK: func.func @atomic_addf_storage_buffer
// CHECK-SAME: (%[[VAL:.+]]: f32,
func.func @atomic_addf_storage_buffer(%value: f32, %memref: memref<2x3x4xf32, #spirv.storage_class<StorageBuffer>>, %i0: index, %i1: index, %i2: index) -> f32 {
- // CHECK: %[[AC:.+]] = spirv.AccessChain
+ // CHECK: %[[AC:.+]] = spirv.InBoundsAccessChain
// CHECK: %[[ATOMIC:.+]] = spirv.EXT.AtomicFAdd <Device> <AcquireRelease|UniformMemory> %[[AC]], %[[VAL]] : !spirv.ptr<f32, StorageBuffer>
// CHECK: return %[[ATOMIC]]
%0 = memref.atomic_rmw "addf" %value, %memref[%i0, %i1, %i2] : (f32, memref<2x3x4xf32, #spirv.storage_class<StorageBuffer>>) -> f32
diff --git a/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir b/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir
index 8679bf807240b..a5d51dfa20bbe 100644
--- a/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir
+++ b/mlir/test/Conversion/MemRefToSPIRV/bitwidth-emulation.mlir
@@ -83,13 +83,13 @@ func.func @load_i64(%arg0: memref<10xi64, #spirv.storage_class<StorageBuffer>>,
// CHECK: %[[ARG1_CAST:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] : index to i32
// CHECK: %[[ZERO:.+]] = spirv.Constant 0 : i32
// CHECK-NOT: spirv.SDiv
- // CHECK: %[[PTR:.+]] = spirv.AccessChain %{{.+}}[%[[ZERO]], %[[ARG1_CAST]]] : {{.+}}, i32, i32
+ // CHECK: %[[PTR:.+]] = spirv.InBoundsAccessChain %{{.+}}[%[[ZERO]], %[[ARG1_CAST]]] : {{.+}}, i32, i32
// CHECK: spirv.Load "StorageBuffer" %[[PTR]] : i64
// CHECK-NOT: spirv.ShiftRightArithmetic
// INDEX64: %[[ARG1_CAST:.+]] = builtin.unrealized_conversion_cast %{{.+}} : index to i64
// INDEX64: %[[ZERO:.+]] = spirv.Constant 0 : i64
- // INDEX64: %[[PTR:.+]] = spirv.AccessChain %{{.+}}[%[[ZERO]], %[[ARG1_CAST]]] : {{.+}}, i64, i64
+ // INDEX64: %[[PTR:.+]] = spirv.InBoundsAccessChain %{{.+}}[%[[ZERO]], %[[ARG1_CAST]]] : {{.+}}, i64, i64
// INDEX64: spirv.Load "StorageBuffer" %[[PTR]] : i64
%0 = memref.load %arg0[%index] : memref<10xi64, #spirv.storage_class<StorageBuffer>>
return %0: i64
@@ -179,12 +179,12 @@ func.func @store_i64(%arg0: memref<10xi64, #spirv.storage_class<StorageBuffer>>,
// CHECK-DAG: %[[ARG0_CAST:.+]] = builtin.unrealized_conversion_cast %[[ARG0]]
// CHECK: %[[ZERO:.+]] = spirv.Constant 0 : i32
// CHECK-NOT: spirv.AtomicAnd
- // CHECK: %[[PTR:.+]] = spirv.AccessChain %[[ARG0_CAST]][%[[ZERO]], %[[ARG1_CAST]]] : {{.+}}, i32, i32
+ // CHECK: %[[PTR:.+]] = spirv.InBoundsAccessChain %[[ARG0_CAST]][%[[ZERO]], %[[ARG1_CAST]]] : {{.+}}, i32, i32
// CHECK: spirv.Store "StorageBuffer" %[[PTR]], %[[ARG2]] : i64
// CHECK-NOT: spirv.AtomicOr
// INDEX64: %[[ZERO:.+]] = spirv.Constant 0 : i64
- // INDEX64: %[[PTR:.+]] = spirv.AccessChain %{{.+}}[%[[ZERO]], %{{.+}}] : {{.+}}, i64, i64
+ // INDEX64: %[[PTR:.+]] = spirv.InBoundsAccessChain %{{.+}}[%[[ZERO]], %{{.+}}] : {{.+}}, i64, i64
// INDEX64: spirv.Store "StorageBuffer" %[[PTR]], %{{.+}} : i64
memref.store %value, %arg0[%index] : memref<10xi64, #spirv.storage_class<StorageBuffer>>
return
diff --git a/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir b/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir
index ebddeadf3a31a..495b5bfc87a3c 100644
--- a/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir
+++ b/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir
@@ -24,12 +24,12 @@ func.func @load_store_zero_rank_float(%arg0: memref<f32, #spirv.storage_class<St
// CHECK-DAG: [[ARG0:%.*]] = builtin.unrealized_conversion_cast %[[OARG0]] : memref<f32, #spirv.storage_class<StorageBuffer>> to !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>
// CHECK-DAG: [[ARG1:%.*]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref<f32, #spirv.storage_class<StorageBuffer>> to !spirv.ptr<!spirv.struct<(!spirv.array<1 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: [[ZERO:%.*]] = spirv.Constant 0 : i32
- // CHECK: spirv.AccessChain [[ARG0]][
+ // CHECK: spirv.InBoundsAccessChain [[ARG0]][
// CHECK-SAME: [[ZERO]], [[ZERO]]
// CHECK-SAME: ] :
// CHECK: spirv.Load "StorageBuffer" %{{.*}} : f32
%0 = memref.load %arg0[] : memref<f32, #spirv.storage_class<StorageBuffer>>
- // CHECK: spirv.AccessChain [[ARG1]][
+ // CHECK: spirv.InBoundsAccessChain [[ARG1]][
// CHECK-SAME: [[ZERO]], [[ZERO]]
// CHECK-SAME: ] :
// CHECK: spirv.Store "StorageBuffer" %{{.*}} : f32
@@ -43,12 +43,12 @@ func.func @load_store_zero_rank_int(%arg0: memref<i32, #spirv.storage_class<Stor
// CHECK-DAG: [[ARG0:%.*]] = builtin.unrealized_conversion_cast %[[OARG0]] : memref<i32, #spirv.storage_class<StorageBuffer>> to !spirv.ptr<!spirv.struct<(!spirv.array<1 x i32, stride=4> [0])>, StorageBuffer>
// CHECK-DAG: [[ARG1:%.*]] = builtin.unrealized_conversion_cast %[[OARG1]] : memref<i32, #spirv.storage_class<StorageBuffer>> to !spirv.ptr<!spirv.struct<(!spirv.array<1 x i32, stride=4> [0])>, StorageBuffer>
// CHECK: [[ZERO:%.*]] = spirv.Constant 0 : i32
- // CHECK: spirv.AccessChain [[ARG0]][
+ // CHECK: spirv.InBoundsAccessChain [[ARG0]][
// CHECK-SAME: [[ZERO]], [[ZERO]]
// CHECK-SAME: ] :
// CHECK: spirv.Load "StorageBuffer" %{{.*}} : i32
%0 = memref.load %arg0[] : memref<i32, #spirv.storage_class<StorageBuffer>>
- // CHECK: spirv.AccessChain [[ARG1]][
+ // CHECK: spirv.InBoundsAccessChain [[ARG1]][
// CHECK-SAME: [[ZERO]], [[ZERO]]
// CHECK-SAME: ] :
// CHECK: spirv.Store "StorageBuffer" %{{.*}} : i32
@@ -231,7 +231,7 @@ func.func @static_linearized_index(
// CHECK: %[[STRIDE:.+]] = spirv.Constant 4 : i32
// CHECK: %[[OFFSET:.+]] = spirv.IMul %{{.*}}, %[[STRIDE]] {no_signed_wrap, no_unsigned_wrap} : i32
// CHECK: %[[LINEAR:.+]] = spirv.IAdd %{{.*}}, %[[OFFSET]] {no_signed_wrap, no_unsigned_wrap} : i32
- // CHECK: spirv.AccessChain {{.*}}[%{{.*}}, %[[LINEAR]]]
+ // CHECK: spirv.InBoundsAccessChain {{.*}}[%{{.*}}, %[[LINEAR]]]
%0 = memref.load %arg0[%row, %column] : memref<2x4xf32, #spirv.storage_class<StorageBuffer>>
return %0 : f32
}
@@ -247,7 +247,7 @@ func.func @unsigned_only_linearized_index(
// CHECK-NOT: no_signed_wrap
// CHECK: %[[LINEAR:.+]] = spirv.IAdd %{{.*}}, %[[OFFSET]] {no_unsigned_wrap} : i32
// CHECK-NOT: no_signed_wrap
- // CHECK: spirv.AccessChain {{.*}}[%{{.*}}, %[[LINEAR]]]
+ // CHECK: spirv.InBoundsAccessChain {{.*}}[%{{.*}}, %[[LINEAR]]]
%0 = memref.load %arg0[%row, %column] : memref<2x1073741825xf32, #spirv.storage_class<StorageBuffer>>
return %0 : f32
}
diff --git a/mlir/test/Conversion/SCFToSPIRV/for.mlir b/mlir/test/Conversion/SCFToSPIRV/for.mlir
index 660f5eeaf93e6..20e92c9c6bdc9 100644
--- a/mlir/test/Conversion/SCFToSPIRV/for.mlir
+++ b/mlir/test/Conversion/SCFToSPIRV/for.mlir
@@ -20,9 +20,9 @@ func.func @loop_kernel(%arg2 : memref<10xf32, #spirv.storage_class<StorageBuffer
// CHECK: spirv.BranchConditional %[[CMP]], ^[[BODY:.*]], ^[[MERGE:.*]]
// CHECK: ^[[BODY]]:
// CHECK: %[[ZERO1:.*]] = spirv.Constant 0 : i32
- // CHECK: spirv.AccessChain {{%.*}}{{\[}}%[[ZERO1]], %[[INDVAR]]{{\]}}
+ // CHECK: spirv.InBoundsAccessChain {{%.*}}{{\[}}%[[ZERO1]], %[[INDVAR]]{{\]}}
// CHECK: %[[ZERO2:.*]] = spirv.Constant 0 : i32
- // CHECK: spirv.AccessChain {{%.*}}[%[[ZERO2]], %[[INDVAR]]]
+ // CHECK: spirv.InBoundsAccessChain {{%.*}}[%[[ZERO2]], %[[INDVAR]]]
// CHECK: %[[INCREMENT:.*]] = spirv.IAdd %[[INDVAR]], %[[STEP]] : i32
// CHECK: spirv.Branch ^[[HEADER]](%[[INCREMENT]] : i32)
// CHECK: ^[[MERGE]]
diff --git a/mlir/test/Conversion/SCFToSPIRV/if.mlir b/mlir/test/Conversion/SCFToSPIRV/if.mlir
index 0b3df9a533302..361b32783d2ee 100644
--- a/mlir/test/Conversion/SCFToSPIRV/if.mlir
+++ b/mlir/test/Conversion/SCFToSPIRV/if.mlir
@@ -139,7 +139,7 @@ func.func @simple_if_yield_type_change(%arg2 : memref<10xf32, #spirv.storage_cla
// CHECK: spirv.mlir.merge
// CHECK-NEXT: }
// CHECK: %[[OUT:.*]] = spirv.Load "Function" %[[VAR]] : !spirv.ptr<!spirv.struct<(!spirv.array<10 x f32, stride=4> [0])>, StorageBuffer>
- // CHECK: %[[ADD:.*]] = spirv.AccessChain %[[OUT]][{{%.*}}, {{%.*}}] : !spirv.ptr<!spirv.struct<(!spirv.array<10 x f32, stride=4> [0])>, StorageBuffer>
+ // CHECK: %[[ADD:.*]] = spirv.InBoundsAccessChain %[[OUT]][{{%.*}}, {{%.*}}] : !spirv.ptr<!spirv.struct<(!spirv.array<10 x f32, stride=4> [0])>, StorageBuffer>
// CHECK: spirv.Store "StorageBuffer" %[[ADD]], {{%.*}} : f32
// CHECK: spirv.Return
%i = arith.constant 0 : index
diff --git a/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir b/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir
index f904dd9d35c37..d8e187e919bf1 100644
--- a/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir
+++ b/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir
@@ -1117,7 +1117,7 @@ module attributes {
// CHECK: %[[CST1:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST2:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST3:.+]] = spirv.Constant 1 : i32
-// CHECK: %[[S4:.+]] = spirv.AccessChain %[[S0]][%[[CST1]], %[[S1]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[S4:.+]] = spirv.InBoundsAccessChain %[[S0]][%[[CST1]], %[[S1]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[S5:.+]] = spirv.Bitcast %[[S4]] : !spirv.ptr<f32, StorageBuffer> to !spirv.ptr<vector<4xf32>, StorageBuffer>
// CHECK: %[[R0:.+]] = spirv.Load "StorageBuffer" %[[S5]] : vector<4xf32>
// CHECK: return %[[R0]] : vector<4xf32>
@@ -1128,6 +1128,17 @@ func.func @vector_load(%arg0 : memref<4xf32, #spirv.storage_class<StorageBuffer>
return %0: vector<4xf32>
}
+// Dynamic StorageBuffer layouts do not identify a fixed base object, so keep
+// the pointer calculation conservative even though the vector access itself
+// has in-bounds source semantics.
+// CHECK-LABEL: @vector_load_dynamic
+// CHECK: spirv.AccessChain {{.*}} !spirv.ptr<!spirv.struct<(!spirv.rtarray<f32, stride=4> [0])>, StorageBuffer>
+func.func @vector_load_dynamic(%arg0 : memref<?xf32, #spirv.storage_class<StorageBuffer>>) -> vector<4xf32> {
+ %idx = arith.constant 0 : index
+ %0 = vector.load %arg0[%idx] : memref<?xf32, #spirv.storage_class<StorageBuffer>>, vector<4xf32>
+ return %0: vector<4xf32>
+}
+
// CHECK-LABEL: @vector_load_single_elem
// CHECK-SAME: (%[[ARG0:.*]]: memref<4xf32, #spirv.storage_class<StorageBuffer>>)
@@ -1137,7 +1148,7 @@ func.func @vector_load(%arg0 : memref<4xf32, #spirv.storage_class<StorageBuffer>
// CHECK: %[[CST1:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST2:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST3:.+]] = spirv.Constant 1 : i32
-// CHECK: %[[S4:.+]] = spirv.AccessChain %[[S0]][%[[CST1]], %[[S1]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[S4:.+]] = spirv.InBoundsAccessChain %[[S0]][%[[CST1]], %[[S1]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[S5:.+]] = spirv.Load "StorageBuffer" %[[S4]] : f32
// CHECK: %[[R0:.+]] = builtin.unrealized_conversion_cast %[[S5]] : f32 to vector<1xf32>
// CHECK: return %[[R0]] : vector<1xf32>
@@ -1170,7 +1181,7 @@ func.func @vector_load_aligned(%arg0 : memref<4xf32, #spirv.storage_class<Storag
// CHECK: %[[S3:.+]] = spirv.IMul %[[S1]], %[[CST4]] : i32
// CHECK: %[[CST1:.+]] = spirv.Constant 1 : i32
// CHECK: %[[S6:.+]] = spirv.IAdd %[[S2]], %[[S3]] : i32
-// CHECK: %[[S7:.+]] = spirv.AccessChain %[[S0]][%[[CST0_1]], %[[S6]]] : !spirv.ptr<!spirv.struct<(!spirv.array<16 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[S7:.+]] = spirv.InBoundsAccessChain %[[S0]][%[[CST0_1]], %[[S6]]] : !spirv.ptr<!spirv.struct<(!spirv.array<16 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[S8:.+]] = spirv.Bitcast %[[S7]] : !spirv.ptr<f32, StorageBuffer> to !spirv.ptr<vector<4xf32>, StorageBuffer>
// CHECK: %[[R0:.+]] = spirv.Load "StorageBuffer" %[[S8]] : vector<4xf32>
// CHECK: return %[[R0]] : vector<4xf32>
@@ -1190,7 +1201,7 @@ func.func @vector_load_2d(%arg0 : memref<4x4xf32, #spirv.storage_class<StorageBu
// CHECK: %[[CST1:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST2:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST3:.+]] = spirv.Constant 1 : i32
-// CHECK: %[[S4:.+]] = spirv.AccessChain %[[S0]][%[[CST1]], %[[S1]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[S4:.+]] = spirv.InBoundsAccessChain %[[S0]][%[[CST1]], %[[S1]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[S5:.+]] = spirv.Bitcast %[[S4]] : !spirv.ptr<f32, StorageBuffer> to !spirv.ptr<vector<4xf32>, StorageBuffer>
// CHECK: spirv.Store "StorageBuffer" %[[S5]], %[[ARG1]] : vector<4xf32>
func.func @vector_store(%arg0 : memref<4xf32, #spirv.storage_class<StorageBuffer>>, %arg1 : vector<4xf32>) {
@@ -1218,7 +1229,7 @@ func.func @vector_store_aligned(%arg0 : memref<4xf32, #spirv.storage_class<Stora
// CHECK: %[[CST1:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST2:.+]] = spirv.Constant 0 : i32
// CHECK: %[[CST3:.+]] = spirv.Constant 1 : i32
-// CHECK: %[[S4:.+]] = spirv.AccessChain %[[S0]][%[[CST1]], %[[S2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
+// CHECK: %[[S4:.+]] = spirv.InBoundsAccessChain %[[S0]][%[[CST1]], %[[S2]]] : !spirv.ptr<!spirv.struct<(!spirv.array<4 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
// CHECK: spirv.Store "StorageBuffer" %[[S4]], %[[S1]] : f32
func.func @vector_store_single_elem(%arg0 : memref<4xf32, #spirv.storage_class<StorageBuffer>>, %arg1 : vector<1xf32>) {
%idx = arith.constant 0 : index
@@ -1240,7 +1251,7 @@ func.func @vector_store_single_elem(%arg0 : memref<4xf32, #spirv.storage_class<S
// CHECK: %[[S3:.+]] = spirv.IMul %[[S1]], %[[CST4]] : i32
// CHECK: %[[CST1:.+]] = spirv.Constant 1 : i32
// CHECK: %[[S6:.+]] = spirv.IAdd %[[S2]], %[[S3]] : i32
-// CHECK: %[[S7:.+]] = spirv.AccessChain %[[S0]][%[[CST0_1]], %[[S6]]] : !spirv.ptr<!spirv.struct<(!spirv.array<16 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
+// CHECK: %[[S7:.+]] = spirv.InBoundsAccessChain %[[S0]][%[[CST0_1]], %[[S6]]] : !spirv.ptr<!spirv.struct<(!spirv.array<16 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
// CHECK: %[[S8:.+]] = spirv.Bitcast %[[S7]] : !spirv.ptr<f32, StorageBuffer> to !spirv.ptr<vector<4xf32>, StorageBuffer>
// CHECK: spirv.Store "StorageBuffer" %[[S8]], %[[ARG1]] : vector<4xf32>
func.func @vector_store_2d(%arg0 : memref<4x4xf32, #spirv.storage_class<StorageBuffer>>, %arg1 : vector<4xf32>) {
More information about the Mlir-commits
mailing list