[all-commits] [llvm/llvm-project] 0a600c: [mlir][nvgpu] Make `phaseParity` of `mbarrier.try_...

[Guray Ozen via All-commits]

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

  Changed paths:
    M mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
    M mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
    M mlir/lib/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp
    M mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
    M mlir/test/Dialect/NVGPU/tmaload-transform.mlir
    M mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir
    M mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir
    M mlir/test/Integration/GPU/CUDA/sm90/tma_load_128x64_swizzle128b.mlir
    M mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
    M mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir

  Log Message:
  -----------
  [mlir][nvgpu] Make `phaseParity` of `mbarrier.try_wait` `i1` (#81460)

Currently, `phaseParity` argument of `nvgpu.mbarrier.try_wait.parity` is
index. This can cause a problem if it's passed any value different than
0 or 1. Because the PTX instruction only accepts even or odd phase. This
PR makes phaseParity argument i1 to avoid misuse.

Here is the information from PTX doc:

```
The .parity variant of the instructions test for the completion of the phase indicated 
by the operand phaseParity, which is the integer parity of either the current phase or 
the immediately preceding phase of the mbarrier object. An even phase has integer 
parity 0 and an odd phase has integer parity of 1. So the valid values of phaseParity 
operand are 0 and 1.
```
See for more information:

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-test-wait-mbarrier-try-wait