[Mlir-commits] [mlir] [MLIR][NVVM] Fix the lowering of legacy mbar.arrive_drop (PR #222916)

[Dhairyashil R G]

dhairyashilRG wrote:

@durga4github  Thanks for the pointer, I've updated the description (email as well).

No, I didn't hit this in a real workload. I was checking, for each mbarrier op, whether what it lowers to is selectable at sm_80. `arrive_drop` was the only one where a legacy sm_80 intrinsic exists but the op never uses it, the other ops either already lower to something selectable at sm_80, or have no legacy form at all so their sm_90 requirement is real.

It does reproduce:

```mlir
llvm.func @k(%b: !llvm.ptr<3>) {
  %0 = nvvm.mbarrier.arrive_drop %b : !llvm.ptr<3> -> i64
  llvm.return
}
```

```
$ mlir-translate -mlir-to-llvmir t.mlir | llc -mtriple=nvptx64 -mcpu=sm_80
LLVM ERROR: Cannot select: intrinsic %llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta
```
So it's the same gap as the `arrive` one, and as far as I can tell the last of its kind here. It's a small fix, happy to close it if you'd rather not take it.



https://github.com/llvm/llvm-project/pull/222916