[RFC] Should `vector.transfer_read`/`write` keep `in_bounds`? Measurements on the masking alternative

MLIR
1

Summary

While reviewing #215340, @banach-space raised a question bigger than that patch:

We have effectively been walking away from in_bounds in favour of masking. Given SVE’s
solid HW support for predication, that’s what we’ve been using instead. But even more
broadly, I am not aware of in_bounds providing anything that masking wouldn’t. So, instead
of making the logic around in_bounds more complex, I’d rather see in_bounds completely
removed and instead rely on masks.

This affects the whole dialect, not one patch, so I am bringing it here with measurements.

@banach-space: I moved this out of the PR thread rather than let it stall there. I have tried to state your proposal the way you made it. Please correct me if I got any of it wrong: the idea is yours, the framing is mine.

Short version. Masking replaces in_bounds completely on hardware that has predication. On hardware that does not, LLVM turns every masked load into one branch per lane, and the code gets about ten times bigger. Only 8 of the 25 in-tree LLVM targets implement masked load/store at all; the other 17 fall back to that per-lane expansion. So removing in_bounds today would make generated code worse on most targets.

I am not arguing against removing it. I am arguing against removing it first. The proposal in Sec.8 is about ordering.

My interest, stated up front: I have an open PR that adds in_bounds inference, so I am not a neutral person. That is why this post is measurements and source links rather than argument. Everything below can be re-run.

1. What in_bounds actually does

in_bounds is not just a hint. It picks which lowering you get. From the op description in VectorOps.td:

A vector.transfer_read can be lowered to a simple load if all dimensions are specified to
be within bounds and no mask was specified.

And the code that enforces it, LowerVectorTransfer.cpp:459:

// Out-of-bounds dims are handled by MaterializeTransferMask.
if (read.hasOutOfBoundsDim())
  return rewriter.notifyMatchFailure(read, "out-of-bounds needs mask");

So “remove in_bounds and rely on masks” means “always take the masked path”.

2. What the masked path costs, per target

This is the main point, and it is decided in the LLVM backend, not in MLIR.

llvm/lib/Transforms/Scalar/ScalarizeMaskedMemIntrin.cpp says what it does:

This pass replaces masked memory intrinsics - when unsupported by the target

  • with a chain of basic blocks, that deal with the elements one-by-one if the
    appropriate mask bit is set.

It asks the target through TTI.isLegalMaskedLoad / isLegalMaskedStore. If the target says no, you get one conditional block per lane.

The default answer is no. TargetTransformInfoImpl.h:363:

virtual bool isLegalMaskedLoad(Type *DataType, Align Alignment,
                               unsigned AddressSpace,
                               TTI::MaskKind MaskKind) const {
  return false;
}

BasicTTIImpl does not override this, so any target that does not implement it inherits false. 8 of 25 in-tree targets override it: AArch64, ARM, Hexagon, NVPTX, PowerPC, RISCV, VE, X86. The other 17, including AMDGPU, WebAssembly, SPIR-V, SystemZ, MIPS, LoongArch and SPARC, always scalarize.

And even the 8 that implement it only allow it under conditions:

target masked load/store allowed when
AArch64 isSVEorStreamingSVEAvailable(). Fixed-width vectors also need to be 64- or 128-bit, or useSVEForFixedLengthVectors (AArch64TargetTransformInfo.h:311)
X86 ST->hasAVX() (X86TargetTransformInfo.cpp:6690)
ARM hasMVEIntegerOps() and EnableMaskedLoadStores (ARMTargetTransformInfo.cpp:1115)
RISC-V hasVInstructions(). Fixed-width also needs useRVVForFixedLengthVectors() (RISCVTargetTransformInfo.h:281)
the other 17 never

The part worth noticing: this is not “x86 versus Arm”. It is a split inside AArch64. SVE has predication and NEON does not, so the same architecture lands on both sides.

3. Measurement A: the mechanism on its own

Hand-written LLVM IR, no MLIR involved. The only difference between the two runs is one target flag.

declare <4 x float> @llvm.masked.load.v4f32.p0(ptr, i32, <4 x i1>, <4 x float>)
declare void @llvm.masked.store.v4f32.p0(<4 x float>, ptr, i32, <4 x i1>)

define <4 x float> @masked_load(ptr %p, <4 x i1> %m, <4 x float> %pass) {
  %r = call <4 x float> @llvm.masked.load.v4f32.p0(ptr %p, i32 4, <4 x i1> %m, <4 x float> %pass)
  ret <4 x float> %r
}
define void @masked_store(<4 x float> %v, ptr %p, <4 x i1> %m) {
  call void @llvm.masked.store.v4f32.p0(<4 x float> %v, ptr %p, i32 4, <4 x i1> %m)
  ret void
}
define <4 x float> @plain_load(ptr %p) {
  %r = load <4 x float>, ptr %p, align 4
  ret <4 x float> %r
}

llc -mtriple=aarch64-unknown-linux-gnu             masked.ll -o neon.s   # NEON
llc -mtriple=aarch64-unknown-linux-gnu -mattr=+sve masked.ll -o sve.s    # SVE

Instructions and branches in each function:

NEON SVE
masked_load 25 instrs, 7 branches 7 instrs, 0 branches
masked_store 23 instrs, 7 branches 6 instrs, 0 branches
plain_load 2 instrs, 0 branches 2 instrs, 0 branches

On NEON you get one conditional block per lane, exactly as the pass description says:

	tbnz	w8, #0, .LBB0_5
	tbnz	w8, #1, .LBB0_6
...
.LBB0_5:                                // %cond.load
	ld1	{ v1.s }[0], [x0]
.LBB0_6:                                // %cond.load1
	add	x9, x0, #4
	ld1	{ v1.s }[1], [x9]

On SVE the mask is just a predicate register:

	cmpne	p1.s, p0/z, z0.s, #0
	ld1w	{ z0.s }, p1/z, [x0]

4. Measurement B: the same thing through MLIR

Same vector.transfer_read twice, differing only in the attribute, put through the standard lowering pipeline:

mlir-opt ib.mlir -lower-affine -convert-vector-to-scf -convert-scf-to-cf \
  -convert-vector-to-llvm -finalize-memref-to-llvm -convert-func-to-llvm \
  -convert-arith-to-llvm -convert-cf-to-llvm -convert-ub-to-llvm \
  -reconcile-unrealized-casts -o ib-llvm.mlir
mlir-translate --mlir-to-llvmir ib-llvm.mlir -o ib.ll
llc -mtriple=aarch64-unknown-linux-gnu             ib.ll -o ib-neon.s
llc -mtriple=aarch64-unknown-linux-gnu -mattr=+sve ib.ll -o ib-sve.s

func.func @with_in_bounds(%m: memref<1000xf32>, %i: index, %p: f32) -> vector<4xf32> {
  %v = vector.transfer_read %m[%i], %p {in_bounds = [true]} : memref<1000xf32>, vector<4xf32>
  return %v : vector<4xf32>
}
func.func @without_in_bounds(%m: memref<1000xf32>, %i: index, %p: f32) -> vector<4xf32> {
  %v = vector.transfer_read %m[%i], %p : memref<1000xf32>, vector<4xf32>
  return %v : vector<4xf32>
}
function NEON SVE
with_in_bounds 3 instrs, 0 branches 3 instrs, 0 branches
without_in_bounds 32 instrs, 7 branches 13 instrs, 0 branches

Two things come out of this. The second one surprised me.

  1. On NEON, dropping in_bounds makes this about 10x bigger and adds seven branches. The four %cond.load blocks are there, same as Measurement A.
  2. Even on SVE it goes from 3 instructions to 13. The load itself is fine. But the mask still has to be built from the memref size before every transfer. Predication makes the memory access cheap. It does not make the mask free.

5. Measurement C: a real kernel

PolyBench 4.2.1 jacobi-2d, MEDIUM size (N=250, TSTEPS=100), unmodified. Both arms are the same MLIR; the only difference is whether in_bounds was inferred. The two arms run interleaved in one process, their outputs are compared bit-for-bit before any timing, and an A-vs-A run gives the noise floor.

AArch64 NEON, vector width 4 masked with in_bounds
kernel instructions 317 85
kernel branches 108 12
llvm.masked.* in the module 14 0
runtime 14.79 ms 2.91 ms (-80%)

For contrast, on AVX-512, where a masked load really is one instruction, the same kernel still gets 13-32% faster depending on vector width, and drops from 88 to 59 instructions. The reason is the same as point 2 above: this kernel does 12 masked memory operations per iteration and builds a mask for each one.

What this does not show. It is one kernel. The change only helps when the trailing dimension is not divisible by the vector width and the loop bound is smaller than that dimension. That describes stencils, not dense linear algebra: 1 of the 7 PolyBench kernels I transcribed. The AArch64 numbers were measured on Apple silicon, which has NEON and no SVE; the AVX-512 numbers on a c7i.8xlarge. Nothing in this post was timed on SVE hardware. The SVE columns in Measurements A and B are static instruction counts from cross-compiling, not timings.

6. Where in_bounds is used today

@banach-space is right that the project has been moving toward masking. But the move is only
half done. createReadOrMaskedRead / createWriteOrMaskedWrite take a
useInBoundsInsteadOfMasking flag, and on main today:

caller which mode
Linalg vectorizer, when the user gives vector sizes (Vectorization.cpp:1866, :2068) masking
Linalg pack / unpack / insert-slice, static shapes (:1699, :1808, :3012) in_bounds
-affine-super-vectorize (SuperVectorize.cpp:1270, :1329) in_bounds

So removing the attribute also means migrating those callers.

How much code is involved, counted on upstream/main @ 7214eb03c282:

count
.cpp/.h files under mlir/lib using the attribute or its accessors 28
files under mlir/include 5
test files mentioning in_bounds 95
times it appears in tests ~1,810

Plus the op description itself, and any out-of-tree code that relies on it.

7. The part of the argument I think is correct

in_bounds really has caused subtle bugs. I found two of them myself, both in the same
fifteen-line function:

  • #213506: the fold ignored scalable
    dimensions. A vector<[4]xf32> read from a memref<4xf32> was marked in-bounds and lowered
    to an unpredicated load. I confirmed it as a segfault on a Graviton3 against a guard page.
  • The same check accepted negative constant indices, which contradicts the spec line saying
    in_bounds covers the starting point too.

An attribute that claims a safety property has to be got right in every folder and every lowering. A mask is safe because of what it is, not because someone proved something. That is a real argument for removal and I do not want to talk it away.

8. Proposal

Not “keep in_bounds”. An ordering:

  1. in_bounds can go once llvm.masked.load/store lower well on targets without predication. Right now they do not, and that is an LLVM backend problem rather than an MLIR one. Removing the attribute before that just moves the cost onto those 17 targets.
  2. The migration in Sec.6 has to finish first. The callers still passing useInBoundsInsteadOfMasking=true need to move over, and that is where the remaining semantic questions actually are.
  3. Separately from both, there is a question of whether MLIR should keep inferring in_bounds while it still exists. My open PR does that inside a folder, which @matthias-springer rightly objected to because of compile-time cost. Both he and @banach-space prefer an opt-in pass, and that is what I plan to write.

Questions

  1. Is the target-support argument in Sec.2 the right way to look at this? Or is the intent to accept slower code on the other 17 targets in exchange for a simpler dialect?
  2. Is there a plan to improve masked lowering on targets without predication that I have missed? If ScalarizeMaskedMemIntrin is expected to get better, the ordering changes.
  3. Point 2 in Sec.4, the cost of building the mask even where masking is native, suggests in_bounds` still buys something on SVE and AVX-512, just less. Is that overhead considered acceptable, or is there a lowering that avoids rebuilding the mask for every transfer?
  4. If removal does go ahead, should the inference work stop now, or is an opt-in pass still worth having in the meantime?

Reproducing

Measurements A and B are two small files and four llc runs. The IR and the exact commands are above, so they should be reproducible from this post alone. Measurement C, the PolyBench kernels, the A/B timing harness and the fold survey are here:
https://github.com/dhairyashilRG/mlir-inbounds-harness

If any number here does not reproduce, please say so and I will correct the post.

2

+1 on removing in_bounds , You can also analyze the mask tensor to find out which dimensions it’s always true for making in_bounds effectively redundant. It’s also not useful because it doesn’t say “how much of this load is out of bounds”. You can have a load of 127 elements and use a vector<128>, but if you split this load up into 16 * vector<8> loads, you get all of them in_bounds = false with in_bounds, but with masking you can prove that the first 15 loads are in bounds while the last one is.

3

Thanks for sharing all these details @dhairyashilRG :smiley:

I’m in favor of removing in_bounds, too. I agree that masking is a cleaner abstraction. The in_bounds flag has caused issues in understanding the interaction with LLVM which has its own inbounds attr for GEPs (e.g., #202118).

Regarding performance, let me try to raise a couple of things I’m curious about, so maybe we can discuss them here. I see two potentially interesting questions:

  1. Today ScalarizeMaskedMemIntrin falls back to one conditional block per lane when isLegalMaskedLoad/Store is false. This is what makes the NEON case explode from 3 instructions to 32 and introduces 7 branches. I’m curious whether there are ways to make this fallback more efficient on targets without native predication, rather than relying on the per-lane branch expansion.
  2. Even on targets with native predication, the measurements show that mask construction itself can be significant: SVE goes from 3 to 13 instructions. I’m curious whether equivalent masks could be hoisted or CSE’d when several transfers share the same bounds, so that the mask is computed once and reused across transfers. Maybe something similar already exists upstream?

The plan you propose in §8 seems very reasonable to me. But I’ll leave this aspect to more experienced folks to evaluate further.

I am open to helping with the implementation of this RFC.

4

Thanks for putting this together, @dhairyashilRG! This has been on the agenda for a long time!

+1 to removing or demoting (see below) in_bounds. I also think it is harder than it looks, which is why it is still pending. A few points on the plan and execution:

  1. in_bounds inference is quite load bearing: Beyond the callers you listed, some folders set the attribute silently (e.g., foldTransferInBoundsAttribute) and codegen quality currently depends on in_bounds being set. We can definitely drop or move the in_bounds attribute inference into an opt-in pass but we would still need a way to produce the same code quality in MLIR. That’s the non-trivial part.

  2. The mask lowering gap is in MLIR: We cannot rely on LLVM alone here. SPIR-V and potentially other non-LLVM targets would also require a mask-to-predication lowering. One option is to migrate the existing in_bounds lowering, generateInBoundsCheck in VectorToSCF to masking, so that the scf.if it emits is driven by masks instead. Having multiple lowerings for masked operations were part of the original masking plan, so this would align well with that.

  3. Masking does not fully subsume in_bounds semantics: A mask states which lanes are active or not but it doesn’t state whether those lanes, active or inactive, are in bounds or out of bounds. We can use masking to mask out elements that are perfectly in bounds. There is a difference between:

  %mask = vector.create_mask %n : vector<8xi1>
  %v = vector.transfer_read %buf[%c0], %pad, %mask : memref<?xf32>, vector<8xf32>

and

  %mask = vector.create_mask %n : vector<8xi1>
  %v = vector.transfer_read %buf[%c0], %pad, %mask {in_bounds} : memref<?xf32>, vector<8xf32>

The latter can be scalarized without predication, because the attribute guarantees all the accessed elements are inside the buffer. The former cannot.

So, I see the required work more as a demotion than a complete removal. The in_bounds attribute on transfer ops should become a hint, similar to what we have in other places in MLIR and LLVM.

Hopefully this helps sketch out the execution plan!

Diego

5

Thanks all. Three things, one of them a correction to something I was about to assert.

@dcaballe, about your sec.8 correction first: you’re right that I framed the lowering gap as an LLVM problem and that’s incomplete. VectorToSPIRV has no handling for transfer or masked ops at all, so it isn’t that a masked path would regress there, it’s that one would have to be built. generateInBoundsCheck looks like the right migration target, and its doc comment says it already accounts for a mask on 1-D transfers, so that reads more like a migration than a rewrite.

@Groverkss, I went to verify the split example and it doesn’t reproduce as stated, though the reason is interesting. Unrolling a vector<4x4> read of a memref<3x3xf32> into 2x2 slices gives [true,true], [true,false], [false,true] and [false,false], per-slice correct, with no -canonicalize in the pipeline. VectorUnroll.cpp:188 does pass the parent attribute verbatim, but the greedy driver folds each new op and foldTransferInBoundsAttribute re-derives it per slice. With a memref<?x?xf32> control, where the folder can prove nothing, all four slices come out bare.

Your abstraction point stands: a per-dimension boolean can’t say how much is out of bounds. But this makes @dcaballe’s first point concrete. That recovery is done entirely by the folder we’re proposing to move into an opt-in pass. Move it, and unrolled transfers stop getting per-slice bounds by default. That seems like a good test case for whatever preserves code quality.

@ FedericoBruzzone, on your second question, most of it already works: vector.create_mask is Pure, so -cse dedupes identical masks and -loop-invariant-code-motion hoists loop-invariant ones (both checked). vector::eliminateVectorMasks removes all-true masks but its only in-tree caller is a test pass, so wiring that up looks like the first concrete step. What’s missing is dedup of masks that are equivalent but not structurally identical.

On demotion to a hint: agreed, that’s a better conclusion than the one I posted. I have #215340 open, which I’m restructuring from the folder into the opt-in pass @ matthiasspringer and @ banach-space asked for. I’ll continue on that basis unless you’d rather it wait for the demotion design.

6

@FedericoBruzzone (tagging separately as only two were allowed above)

7

Thank you for bringing this up, @dhairyashilRG !

I think that we are broadly aligned, so let me share my view on some finer points.

Agreed, that’s a subtle, but also a significant difference IMO.

However, from my experience that is not how we have been differentiating between the two and the very existence of useInBoundsInsteadOfMasking in createWriteOrMaskedWrite is a very good example. We have been mixing both concepts and I would use this effort to revert that.

Thank you for all the analysis, @dhairyashilRG :folded_hands: IMHO, one important point is missing though. Specifically, regarding this remark:

Not necessarily and I think that that’s key :slight_smile: Let me suggest the following evolution:

  • in_bounds = {true} → %mask = "all true" → no mask!

(provided there is no other source of information/direction requiring a mask).

EXAMPLE 1 (in_bounds)
Before:

  %v = vector.transfer_read %buf[%c0], %pad, in_bounds= [true] : memref<?xf32>, vector<8xf32>

After in_bounds is removed (mask is used instead of in_bounds):

  %mask = vector.constant_mask [8] : vector<8xi1>
  %v = vector.transfer_read %buf[%c0], %pad, %mask : memref<?xf32>, vector<8xf32>

After canonicalization (no mask):

  %v = vector.transfer_read %buf[%c0], %pad : memref<?xf32>, vector<8xf32>

Examples like the above should not be lowered to llvm.masked.*. If that is not the case today, we need to fix that.

EXAMPLE 2 (in_bounds + mask)
Before:

  %mask = vector.create_mask %n : vector<8xi1>
  %v = vector.transfer_read %buf[%c0], %pad, in_bounds= [true] : memref<?xf32>, vector<8xf32>

After in_bounds is removed (mask was already present, so it’s preserved):

  %mask = vector.create_mask %n : vector<8xi1>
  %v = vector.transfer_read %buf[%c0], %pad : memref<?xf32>, vector<8xf32>

Examples like the above must be lowered to llvm.masked.* unless there’s a mechanism to prove that the mask is either ALL_TRUE or ALL_FALSE.

Basically, I argue that whatever mechanism is used today to prove that in_bounds should be true, can be used to prove that the mask should be ALL_TRUE and hence avoided. At least in the cases in which in_bounds = [true] is crucial to avoid llvm.masked.load/store (IIUC, that only applies when there is no implicit mask present in the original IR).

Btw, I don’t believe that we have ever clarified/documented the relationship between in_bounds and mask in cases like this (in_bounds and mask are present):

  %mask = vector.create_mask %n : vector<8xi1>
  %v = vector.transfer_read %buf[%c0], %pad, in_bounds= [false] : memref<?xf32>, vector<8xf32>

The comparison between SVE and Neon is tricky - those are two different ISAs and (effectively) LLVM backends. I guess that your point was more about comparing ISA with predication (SVE) vs ISA without predication (Neon) than SVE and Neon specifically.

I do agree that we need to make sure that there’s no regression in terms of instruction count and performance.

I am not sure what you meant by that.

For SVE we have put a lot of effort into avoiding predication when it’s not needed (e.g. when masks are ALL_TRUE) - computing the mask is not free. Is that what you had in mind? If yes, then let me paraphrase my point above:

  • Lets take the logic that we have today to show that in_bounds = {true} and re-use that to show that %mask = "all true" and hence avoid masking altoghether.

I hope that this is helpful. @dhairyashilRG , do my points agree with your observations?

-Andrzej

8

I was only talking about dynamic memref cases here. We should design for dynamic cases.

9

Thanks, @banach-space Yes, your points agree with my observations.

On the masked path. You are right, and my framing was too strong. I wrote that removing in_bounds means always taking the masked path. That is only true if nothing proves the mask is all-true. Your Example 1 is the case I did not account for. The instruction counts in my first post came from hand-written LLVM IR calling llvm.masked.load/llvm.masked.store directly, so they measure the cost of reaching those intrinsics, not the cost of removing in_bounds. They are the cost of removal without an all-true proof.

On “predication makes the memory access cheap, it does not make the mask free”. Agree. The predicated load is one instruction, but the predicate still has to be produced, and when the bound depends on the loop it is produced every iteration.

On reusing the in_bounds proof to show a mask is all-true. I agree, and part of it exists. eliminateVectorMasks in VectorMaskElimination.cpp does this for vector.create_mask: operands that are constants or constant multiples of vscale are checked directly, and for the remaining operands it computes a lower bound with ScalableValueBoundsConstraintSet. If every dimension is at least the mask dimension size it replaces the create_mask with an all-true vector.constant_mask, which the transfer and vector.mask folds then remove. Two things stop it from being the mechanism today.

First, it is scalable-only: if (!vscaleRange) return; is the first statement in the function, under a TODO reading “Support fixed-size case. This is less likely to be useful as for fixed-size code dimensions are all static so masks tend to fold away.” That holds when the mask operands are static. It does not hold when the mask bound comes from a loop induction variable.

Second, its only in-tree caller is TestVectorTransforms.cpp, so even the scalable path does not run in any pipeline.

So the mechanism needs a fixed-size path and a home in a pipeline. The proof itself is the logic in #215340: isProvablyInBounds uses ValueBoundsConstraintSet to bound the transfer index rather than the mask dimension, but it is the same question asked about a different value. I plan to take both pieces on, in that order, once #215340 settles.

On useInBoundsInsteadOfMasking. Agreed, and it is worth being concrete about the scale. There are seven call sites: hard-coded true at two in SuperVectorize.cpp, hard-coded false at two in the Linalg vectorizer, and set from “were vector sizes supplied” at three more there. Undoing the conflation means each of those has to say which of the two properties it
actually wants.

On in_bounds = [false] together with an explicit mask. It is not documented, but the intent is: VectorOps.td carries a TODO directly above transfer_read, “Tighten semantics so that masks and inbounds can’t be used simultaneously within the same transfer op.” Removing in_bounds resolves that TODO rather than requiring a definition for the combination.

On SVE versus Neon. Yes, predication versus no predication was the axis I meant.

@Groverkss, understood. With that restriction your example holds and my static counterexample was beside the point. The dynamic control I reported is the confirmation, not the refutation: on memref<?x?xf32> every unrolled slice comes out with no in_bounds at all, because the folder can prove nothing, so the attribute carries no per-slice information there.

One gap on that path. UnrollCreateMaskPattern already splits a create_mask per slice, computing each local size as min(max(originalSize[d] - offset[d], 0), unrolledSize[d]) with arith.subi, arith.maxsi and arith.minsi, all of which implement ValueBoundsOpInterface, so the per-slice all-true proof is within reach of the mechanism above. But UnrollTransferReadPattern and UnrollTransferWritePattern both return failure when the transfer carries a mask operand, so a masked transfer is not unrolled today. That is the piece that has to exist before masks can replace in_bounds under unrolling.

10

Thanks for the detailed measurements and follow-up, @dhairyashilRG.

I ran an NVIDIA/NVPTX datapoint for this RFC, which grew out of llvm-project #215340. The reproducers use dynamic memrefs in A/B/C, following @Groverkss’s point that the migration must handle dynamic cases. The kernel-ABI control in D has static memref shapes but a dynamic index. The checkout was llvm-project f0d41abb33b6; mlir-opt, mlir-translate, and llc were rebuilt from it.
(The tree had one unrelated uncommitted NVGPU shared-memory change; none of the four pipelines in run.sh invokes nvgpu-optimize-shared-memory.)

Reproducer bundle and commands: mlir-inbounds-nvptx.

LLVM_SOURCE_ROOT=/path/to/llvm-project BIN=/path/to/llvm-project/build/bin ./run.sh

Current GPU conversion requires unmasked, in-bounds transfers.

I compared the same 16x8x16 f16 vector.contract tile in three forms:

input form nvgpu.ldmatrix nvgpu.mma.sync remaining vector.transfer_read remaining vector.contract
in_bounds = [true, true] 3 1 0 0
explicit runtime mask, no in_bounds 0 0 3 1
no in_bounds, no mask 0 0 3 1

The second row models the proposed mask-only form. The third is not a proposed replacement: it shows what remains if that mask can be proved all-true and eliminated. Even then, GPU conversion still needs an independent bounds proof. On dynamic memrefs, the current patterns see an unproved out-of-bounds dimension and leave the transfers and contraction in place, not because of a mask.
Proving a mask all-true does not by itself prove the transfer indices are in bounds. Where source
dimensions and indices have usable bounds in IR, rederivation is plausible; #215340 illustrates
this for static source dimensions. Its pass declines dynamic source dimensions, and in the third
row nothing relates the dynamic memref extents to the tile. No index-only analysis can recover a
fact absent from the IR. For this shape the boundedness guarantee must be carried in some IR form and used by GPU conversion before in_bounds is removed.

The second row likewise has no in_bounds; it does not rely on the mask-plus-in_bounds
combination that VectorOps.td marks for future restriction.

I also carried the positive row through the complete mlir-opt -> mlir-translate -> llc path. The resulting PTX has three ldmatrix.sync instructions and one mma.sync, including ldmatrix.sync.aligned.m8n8.x4.shared.b16 and mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16.
The full sequence is reproducible with the bundle.

This follows the current conversion predicates: transferReadSupportsMMAMatrixType
(VectorToGPU.cpp:174), transferWriteSupportsMMAMatrixType (:200),
CombineTransferReadOpTranspose (:506), and nvgpu::canLowerToWarpMatrixOperation
(MMAUtils.cpp:270,297) each reject a transfer with a mask or an out-of-bounds dimension.

Runtime masks introduce per-lane control flow in current NVPTX lowering.

Rows combining a mask with in_bounds are deliberate controls for the current IR: they separate
mask cost from bounds checks and are not the proposed replacement form.

The PTX instruction and branch counts below are static totals, including prologue and epilogue.
Compare corresponding rows; these counts are not runtime measurements.

For a 4xf32 transfer over a dynamic memref, I get the following counts with sm_80:

function PTX instructions branches loads
read_in_bounds 10 0 4
read_maybe_oob (no in_bounds, no mask) 28 4 4
read_explicit_mask_in_bounds 27 4 4
read_explicit_mask_maybe_oob 35 4 4
read_all_true_mask_in_bounds (constant_mask [4]) 10 0 4
read_partial_constant_mask_in_bounds (constant_mask [3]) 10 0 3
write_in_bounds 10 0 0
write_maybe_oob 22 4 0

Unlike the RFC’s direct llvm.masked.* intrinsic control, this comparison starts at MLIR
vector.transfer_read/write and follows the lowering through LLVM IR to PTX.

The instruction and branch counts in this full 4xf32 matrix are identical for sm_70, sm_80,
and sm_90.

The smaller 8xf32 control gives 15 instructions and 0 branches with in_bounds, 49 instructions
and 8 branches with a runtime mask, and 15 instructions and 0 branches with an all-true constant
mask.

The constant-mask controls complement @banach-space’s all-true example: in these in-bounds rows, even a partially active constant mask remains straight-line, whereas a runtime mask produces one predicate and conditional branch per lane.

The NVPTX TTI legality checks explain this split. Both isLegalMaskedLoad and
isLegalMaskedStore require MaskKind::ConstantMask; runtime masks are rejected and fall back to generic scalarization. The store path is narrower still: it requires at least 32-byte alignment, target support for 256-bit accesses, and exactly eight 32-bit or four 64-bit elements. This support was added by 17852deda7fb ([NVPTX] Lower LLVM masked vector loads and stores to PTX, 2025-11-25).

I repeated the read comparison with a real gpu.func ... kernel ABI and global address space:

kernel PTX instructions branches ld.global
kernel_in_bounds 14 0 4
kernel_maybe_oob 31 4 4

Neither row uses ld.global.v4.b32; both contain four scalar ld.global.b32 because alignment of the descriptor-provided base pointer cannot be proved. The difference therefore isolates the
control-flow overhead rather than a change in vector load width.

The boundary arguments in these reproducers are uniform, so the branches do not imply warp
divergence here. A real tiled GEMM can derive its boundary from block or thread indices; when that condition varies across threads, the same control flow can diverge.

Implication for the proposal.

As @dhairyashilRG already noted in his September 2 reply, vector::eliminateVectorMasks returns early for fixed-size vectors and is called only by a test pass; he also identified the masked
transfer bailouts in VectorUnroll.cpp:162,217. PR #221595 proposes fixed-size support, but it
is still open, and pipeline integration is not yet present in-tree. On the NVIDIA path, the
mask-free third row makes the proof/recovery gap a prerequisite for reaching Tensor Core operations, not merely for reducing masked-load overhead. The current masked-transfer bailout remains a separate obstacle when unrolling masked transfers to hardware-sized vectors.

These are static PTX results from llc -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80. I did not have ptxas or NVIDIA hardware available, so I do not have SASS or runtime measurements.

This does not argue for retaining in_bounds permanently; it adds NVIDIA prerequisites to the
ordering already proposed in this thread. For this path, the sequence should preserve or rederive
boundedness and make GPU conversion use it first, support masked-transfer unrolling where the
migration needs it, and only then remove in_bounds.

Is there another way to carry or recover boundedness before GPU conversion that I have missed?

11

Thank you for this summary, @vaisman .

Unfortunately, there is a lot of fine detail pertinent to Nvidia dialects, GPU architecture and lowering, which I don’t follow. I will defer to experts in the area, @dcaballe and @grypp , to comment.

I have two high-level questions though:

  • In practice/production, what is generating those masked vector.transfer_read + vector.transfer_write Ops for you? In many cases, the easiest solution lies at the very start of compilation/lowering.
  • Could you provide a high-level example (not involving Nvidia dialects) that would illustrate the problem that you are referring to?

Thank you,
-Andrzej

12

For the mask case you should be setting in_bounds as true if your mask is masking out OOB reads. If you supply the mask + no inbounds, it could easily mean that the mask isn’t handling OOB accesses. (If we do remove in_bounds, this would be the assumption on transfer_read that all OOB accesses must be masked out by the user).

In IREE, we had a pass to simplify / fold arith and proving a mask all true was trying to constant fold a i1 vector to a splat, which generalized well.

I think that’s what @banach-space mentioned in his earlier reply:

Not necessarily and I think that that’s key :slight_smile: Let me suggest the following evolution:

  • in_bounds = {true} → %mask = "all true" → no mask!

(provided there is no other source of information/direction requiring a mask).

If you say in_bounds is true, then the mask can be used to find what is OOB.

Maybe it’s just something that can be clarified in the docs and fixed in any folders that are not obeying this.

If we remove the in_bounds attribute, we have to add a assumption that all OOB accesses are undefined and need to be modeled by masking. (Note that this is easy for a tensor compiler to do, not sure for others, because for a tensor compiler, when you vectorize a tensor, you already know if something can be out of bounds or not and mask accordingly).

13

@ Groverkss, I checked what the lowerings do today when a transfer has both a mask and in_bounds = true. They already follow your reading: the mask covers the out-of-bounds lanes.

I tried two paths: -convert-vector-to-llvm, and -convert-vector-to-scf=full-unroll followed by -convert-vector-to-llvm. Both give llvm.intr.masked.load driven by the mask, and no plain load. In the SCF path, in_bounds = true removes the per-row scf.if bounds checks, but each unrolled row keeps its part of the mask, so rows past the bound get an all-false mask. A masked-off lane is never read.

So making the docs say this would describe what already happens. Two things disagree with it:

  • @dcaballe’s reading, where in_bounds on a masked transfer means every lane, active or not, is inside the buffer.
  • The TODO above both transfer ops in VectorOps.td: “Tighten semantics so that masks and inbounds can’t be used simultaneously within the same transfer op”. I quoted this TODO in my last reply as the intent. The lowerings follow your reading instead, so the TODO needs to change.

@vaisman, I think Groverkss’s last paragraph also answers your question. If out-of-bounds accesses are undefined unless masked, a transfer with no mask is in bounds by definition, and the third row of your table would carry that fact on its own.

The cost is that it changes the meaning of transfers that have no mask and no in_bounds today, and not only tensor compilers create them. -affine-super-vectorize="virtual-vector-size=8" on a 100-element copy loop gives:

affine.for %arg2 = 0 to 100 step 8 {
  %1 = vector.transfer_read %arg0[%arg2], %0 : memref<100xf32>, vector<8xf32>
  vector.transfer_write %1, %arg1[%arg2] : vector<8xf32>, memref<100xf32>
}

The last iteration touches elements 96 to 103 of a 100-element buffer. Today this is safe, because it lowers to llvm.intr.masked.load and llvm.intr.masked.store. Under the new rule both would be undefined, and the write would be an out-of-bounds store.

These ops come from createReadOrMaskedRead and createWriteOrMaskedWrite. Both return without building a mask when useInBoundsInsteadOfMasking is true, and both call sites in SuperVectorize.cpp pass true. So those call sites would have to emit masks before the rule changes.

@dcaballe, is there a case where in_bounds on a masked transfer needs to cover the masked-off lanes too? The two paths I tried don’t rely on it. If you and Groverkss agree on one reading, I can send a patch that puts it in the op docs and replaces the TODO.

Separately, #221595, which adds fixed-size support to eliminateVectorMasks as discussed earlier in this thread, is open for review.

14

Thanks, @dhairyashilRG and @Groverkss. Agreed on the mask-free third row: it records the observed convert-vector-to-gpu result for transfers with neither a mask nor in_bounds; it does not refute the proposed rule that out-of-bounds accesses must be masked. The 100-element -affine-super-vectorize copy is a concrete existing tail case that would need updating to preserve its behavior before such a rule change.

15

Thank you for the updates!

(extra emphasis added by me)

I don’t quite understand the issue. Surely the fix is to set useInBoundsInsteadOfMasking to false to make sure that masks are indeed generated. Why wouldn’t that work? Please share a reproducer.

Removing in_bounds would effectively mean hard-coding that flag to true. That’s the path I suggest exploring.

I am not aware of such a path.

What update are you proposing? Wouldn’t removing in_bounds resolve this? And I see no arguments for keeping in_bounds in this thread.

Separately, #221595, which adds fixed-size support to eliminateVectorMasks as discussed earlier in this thread, is open for review.

Reviewed and approved.

Thank you for pushing on this!

16

You’re right, that’s the fix. I should have said so, rather than only saying that these call sites have to change. But I tried it, and flipping the flag isn’t enough on its own today. The mask code in the two helpers doesn’t handle what SuperVectorize passes in:

  • createReadOrMaskedRead builds the mask from the memref size and ignores the read indices.
  • createWriteOrMaskedWrite does subtract an index, but the wrong one when the memref has more dimensions than the vector.

(#202766 added custom indices to the read helper and a permutation map to both, but the mask code stayed as it was.)

Here’s the reproducer. I changed both call sites in SuperVectorize.cpp (lines 1270 and 1329) to false and ran mlir-opt -affine-super-vectorize="virtual-vector-size=8" on:

func.func @copy(%src: memref<100xf32>, %dst: memref<100xf32>) {
  affine.for %i = 0 to 100 {
    %v = affine.load %src[%i] : memref<100xf32>
    affine.store %v, %dst[%i] : memref<100xf32>
  }
  return
}

Output:

%1 = vector.create_mask %c100 : vector<8xi1>
%2 = vector.mask %1 { vector.transfer_read %arg0[%arg2], %0 {in_bounds = [true]} ... }
%3 = arith.subi %c100_0, %arg2 : index
%4 = vector.create_mask %3 : vector<8xi1>
vector.mask %4 { vector.transfer_write %2, %arg1[%arg2] {in_bounds = [true]} ... }

The write mask is 100 - %arg2, which is right. The read mask is just 100, so every lane is on, and the read is also marked in_bounds = [true]. After -lower-affine -lower-vector-mask -convert-vector-to-llvm the read is a plain llvm.load, and on the last iteration it reads elements 96 to 103 of a 100-element memref.

It’s worse with a 2-D memref and a 1-D vector. The read fails to verify (“'vector.create_mask' op must specify an operand for each result vector dimension”), and the write mask uses the row index where it needs the column index. 12 of the 19 tests in test/Dialect/Affine/SuperVectorize fail, and 11 of those fail on invalid IR.

I fixed the mask code locally so it uses the indices and follows the permutation map, then set the flag to false again. Now all the output verifies, the masks are right in the 1-D, 2-D and transposed cases I looked at, and the Linalg and Vector tests still pass. 10 SuperVectorize tests fail only because their CHECK lines don’t expect masks. The 11th is vectorize_2d_inbounds.mlir, which checks that an evenly divisible 512x512 copy lowers to a plain load and store. With the flag flipped it gets masked ones.

That last one is where your Example 1 comes in, and it works. On a 128-element copy the vectorizer now adds masks where today it adds none, and -test-eliminate-vector-masks proves them all-true and removes them, so you get the same plain load and store as today. (On main that pass only does this with a vscale range. #221595 lifts that.)

So the order I have in mind: first fix the mask code in the two helpers, then switch SuperVectorize to false once mask elimination runs in a real pipeline, so loops that divide evenly don’t get slower. I’ll send the first part as a PR.

I want to make sure I read this right. With true the helpers never create a mask, so I’d have expected removing in_bounds to mean hard-coding it to false: always create the mask, then drop the ones that are all-true. Did you mean false, or am I missing something?

Yes, it would. The doc change was only meant for the time before in_bounds is gone, so I’ll drop it.

Agreed, and I’m not arguing for keeping it. My only concern is the order of the steps before it goes.

One correction to post 9, since you quoted it: there are 12 calls to these two helpers, not 7. Two pass true (SuperVectorize), five always get false (two explicitly, three through the default), and five set it based on whether vector sizes were given. It doesn’t change the point.

And thanks for reviewing and approving #221595.

17

Thank you for your detailed reply, @dhairyashilRG !

SG!

I meant false, sorry for the confusion!

Last, but not least, this is tons of great work, thank you for taking care of it :folded_hands:

18

Thanks @banach-space, and sorry for the delayed follow-up. This answers your two questions from #11 and corrects one point from my earlier post.

Correction. Under the rule Groverkss proposed in #12, every out-of-bounds lane must be masked off; otherwise the access is undefined. An unmasked transfer can therefore be treated as in-bounds. The mask-free third row from my earlier post (A3 in the reproducer) is not a blocker under that rule.

What produces the masks? A plain upstream pipeline: a dynamic linalg.matmul, tiled to [16, 8, 16], peeled across all three loops, and vectorized with vector_sizes and create_named_contraction. In the reproducer, vectorization initially masks all eight contractions.
canonicalize removes the mask from the full main tile, but seven genuinely partial remainder masks remain.

eliminateVectorMasks, including the fixed-size support from #221595, cannot prove these masks all-true, so the following canonicalization leaves them in place.

Example without NVIDIA dialects. The remainder is ordinary Vector IR:

%mask = vector.create_mask %m, %n, %k : vector<16x8x16xi1>
%r = vector.mask %mask {
  vector.contract ...
} : vector<16x8x16xi1> -> vector<16x8xf16>

Those partial masks therefore reach convert-vector-to-gpu. The two MMA preparation patterns ignore the enclosing vector.mask: they can insert additional operations into its region, producing invalid IR, or only swap the operands, effectively exchanging the m and n iteration dimensions under the unchanged mask. The latter still verifies for square shapes but applies the bounds to the wrong dimensions. This happens on both the vendor-neutral gpu.subgroup_mma path and the nvgpu.mma.sync path.

Running -lower-vector-mask first does not avoid the failure. It moves the masks on vector.transfer_read and vector.transfer_write to mask operands, but vector.contract has no mask operand and remains inside vector.mask.

On the unpatched pipeline, conversion fails when it reaches the first masked remainder, so it produces no usable converted result for the function. With #226732, the preparation patterns bail out on masked contractions: the eligible main tile lowers to one nvgpu.mma.sync in H and one gpu.subgroup_mma_compute in I, while the partial remainders stay in Vector IR.

Why this matters for the RFC. This seems relevant to the ordering in the RFC. Before changing the transfer semantics and making producers represent potentially out-of-bounds accesses with masks, downstream patterns that can see region-masked operations need either to preserve the enclosing vector.mask or to bail out safely.

The -eliminate-vector-masks pass proposed in #226517 can turn provably all-true vector.create_mask operations into constant masks that canonicalization can remove. Genuinely partial masks still reach downstream patterns.

Canonicalization has the same blind spot: the transfer RAW/WAW folds treat an operation inside vector.mask as unmasked, which can silently change the masked-off lanes. I have a fix for that as well.

I found the same representation mismatch in Vector unrolling. Transfers with a mask operand already bail out, but an enclosing vector.mask is not recognized, so unrolling places multiple slices inside its region and produces invalid IR. I have a bail-out fix ready in the meantime.

As a next step, I would like to look at unrolling masked transfers in both mask forms, slicing the mask per tile with vector.extract_strided_slice, as UnrollGatherPattern already does for
its mask operand.

@dhairyashilRG, would that fit with the helper work you outlined, or would you prefer to keep it separate?

Reproducers and exact commands, sections H and I: mlir-inbounds-nvptx

19

Thank you for your detailed reply, @vaisman !

Sounds like we are all aligned - in_bounds is to be removed. First, however, any functionality that depends on in_bounds needs to be ported to rely on masks instead. In particular the GPU path as highlighted by Vasili - thank you for looking into that!

@vaisman , IIUC, you are basically reporting two issues related to vector.mask:

Issue 1: -convert-vector-to-gpu incorrectly handles vector.mask` regions.

Issue 2: -lower-vector-mask does not remove masks from vector.contract - is that what you’d expect?

Mindful that there’s quite a few issues here, I created #227821 to track all of this. @dhairyashilRG and @vaisman , are you OK to use that ticket to track the outstanding items?

Thank you all for helping to improve Vector :folded_hands:
-Andrzej

20

Yes, #227821 ([mlir][vector] Improve masking + remove in_bounds from Xfer Ops · Issue #227821 · llvm/llvm-project · GitHub) works for me. Here is where my three items stand:

Separate is better. My work is in the two helpers and their callers, and unrolling is in VectorUnroll.cpp, so they don’t overlap. I’m not working on unrolling, so please go ahead.

The switch will make SuperVectorize emit vector.mask regions, so I’ll check it against the fold issue you describe before I send it.