[all-commits] [llvm/llvm-project] b5d694: [mlir][nvvm] Introduce `nvvm.barrier` OP (#81487)

[Guray Ozen via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: b5d694ba14524e0161421b13c875747d5fa917de
      https://github.com/llvm/llvm-project/commit/b5d694ba14524e0161421b13c875747d5fa917de
  Author: Guray Ozen <guray.ozen at gmail.com>
  Date:   2024-02-14 (Wed, 14 Feb 2024)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/test/Dialect/LLVMIR/nvvm.mlir
    M mlir/test/Target/LLVMIR/nvvmir.mlir

  Log Message:
  -----------
  [mlir][nvvm] Introduce `nvvm.barrier` OP (#81487)

This PR that introduces the `nvvm.barrier` OP to the NVVM dialect.
Currently, NVVM only supports the `nvvm.barrier0`, which synchronizes
all threads using barrier resource 0.

The new `nvvm.barrier` has two essential arguments: the barrier resource
and the number of threads. This added flexibility allows for selective
synchronization of threads within a CTA, aligning with the capabilities
provided by LLVM intrinsics or the PTX model.

I think we can deprecate `nvvm.barrier0` in favor of the more generic
`nvvm.barrier`.

```
// Equivalent to nvvm.barrier0 (or __syncthreads() in CUDA)
nvvm.barrier

// Synchronize all threads using the 3rd barrier resource.
nvvm.barrier id = 3

// Synchronize %numberOfThreads threads using the 3rd barrier resource.
nvvm.barrier id = 3 number_of_threads = %numberOfThreads
```