[all-commits] [llvm/llvm-project] bc7e39: [MLIR][NVGPU] Add `mbarrier.get` Op (#133221)

[Guray Ozen via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: bc7e3915e1a0a4e859115dec38c3fdff5e43fcf7
      https://github.com/llvm/llvm-project/commit/bc7e3915e1a0a4e859115dec38c3fdff5e43fcf7
  Author: Guray Ozen <guray.ozen at gmail.com>
  Date:   2025-03-27 (Thu, 27 Mar 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
    M mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
    M mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir

  Log Message:
  -----------
  [MLIR][NVGPU] Add `mbarrier.get` Op (#133221)

The `mbarrier.create` op can create multiple mbarrier objects, and other
mbarrier-related ops can access an mbarrier using a dynamic SSA value.
This is especially useful when using mbarriers in dynamic loops.

This PR adds the `mbarrier.get` op, which returns a pointer to a
specific mbarrier object from a group of barriers created by the
nvgpu.mbarrier.create operation. It is useful when composing the NVGPU
and NVVM dialects.

Example:
```
%mbars = nvgpu.mbarrier.create 
   -> !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>, num_barriers = 10>

%mbar_pointer = nvgpu.mbarrier.get %mbars[%c2] 
  : !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>> 
  -> i32
 ```



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications