)]}'
{
  "commit": "0a600c34c8c1fe87c9661b6020e5044b24da3dc7",
  "tree": "c17db2b168026611bed5494a230c4cd337b8eaaa",
  "parents": [
    "05ad0d46325732e2f7759cb93c94f3e15b41d110"
  ],
  "author": {
    "name": "Guray Ozen",
    "email": "guray.ozen@gmail.com",
    "time": "Tue Feb 13 09:50:34 2024 +0100"
  },
  "committer": {
    "name": "GitHub",
    "email": "noreply@github.com",
    "time": "Tue Feb 13 09:50:34 2024 +0100"
  },
  "message": "[mlir][nvgpu] Make `phaseParity` of `mbarrier.try_wait` `i1` (#81460)\n\nCurrently, `phaseParity` argument of `nvgpu.mbarrier.try_wait.parity` is\r\nindex. This can cause a problem if it\u0027s passed any value different than\r\n0 or 1. Because the PTX instruction only accepts even or odd phase. This\r\nPR makes phaseParity argument i1 to avoid misuse.\r\n\r\nHere is the information from PTX doc:\r\n\r\n```\r\nThe .parity variant of the instructions test for the completion of the phase indicated \r\nby the operand phaseParity, which is the integer parity of either the current phase or \r\nthe immediately preceding phase of the mbarrier object. An even phase has integer \r\nparity 0 and an odd phase has integer parity of 1. So the valid values of phaseParity \r\noperand are 0 and 1.\r\n```\r\nSee for more information:\r\n\r\nhttps://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-test-wait-mbarrier-try-wait",
  "tree_diff": [
    {
      "type": "modify",
      "old_id": "a0c0d4cfd8714bae8ce57a9776811acefaa955df",
      "old_mode": 33188,
      "old_path": "mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td",
      "new_id": "dda8f31e688fe9bf9951cc3e35210be716c651e1",
      "new_mode": 33188,
      "new_path": "mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td"
    },
    {
      "type": "modify",
      "old_id": "5080956a4589828b9db18b2ecfa7bc9155730762",
      "old_mode": 33188,
      "old_path": "mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp",
      "new_id": "9b5d19ebd783a92d0eac75de2bc152788bfda71b",
      "new_mode": 33188,
      "new_path": "mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp"
    },
    {
      "type": "modify",
      "old_id": "c81742233e6d0ec0a3bf2cd7410af5e63d350f2a",
      "old_mode": 33188,
      "old_path": "mlir/lib/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp",
      "new_id": "1635297a5447d4f92c03c5c4486954c71a13e927",
      "new_mode": 33188,
      "new_path": "mlir/lib/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp"
    },
    {
      "type": "modify",
      "old_id": "09a873fa331d596364612092c848a615d6f3a8d5",
      "old_mode": 33188,
      "old_path": "mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir",
      "new_id": "dbf8ead49f78db8f7dd13986564e2060e9cbcc25",
      "new_mode": 33188,
      "new_path": "mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir"
    },
    {
      "type": "modify",
      "old_id": "29e300a992d3ad7ebe9f576e9e8e194e89b03d14",
      "old_mode": 33188,
      "old_path": "mlir/test/Dialect/NVGPU/tmaload-transform.mlir",
      "new_id": "40acd82cd0558513ca943917db9d577fd2cc92e8",
      "new_mode": 33188,
      "new_path": "mlir/test/Dialect/NVGPU/tmaload-transform.mlir"
    },
    {
      "type": "modify",
      "old_id": "35ca0ee8677cca3b338cb763394d13f98d598c73",
      "old_mode": 33188,
      "old_path": "mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir",
      "new_id": "51bcf459d83d5ca9285a756f99aefd8fc5f2bc6d",
      "new_mode": 33188,
      "new_path": "mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir"
    },
    {
      "type": "modify",
      "old_id": "5a10bbba26d8cf974f03bdbf12d7dc7d92c3bd2d",
      "old_mode": 33188,
      "old_path": "mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir",
      "new_id": "85bdb38d67f0fe95563a899a3c9d9867646e3760",
      "new_mode": 33188,
      "new_path": "mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir"
    },
    {
      "type": "modify",
      "old_id": "9c5aacf96b0d69c8e59574e347482eac3fc6fedf",
      "old_mode": 33188,
      "old_path": "mlir/test/Integration/GPU/CUDA/sm90/tma_load_128x64_swizzle128b.mlir",
      "new_id": "b50772f8249fb745e7f1a7f07868dc68d13a1f44",
      "new_mode": 33188,
      "new_path": "mlir/test/Integration/GPU/CUDA/sm90/tma_load_128x64_swizzle128b.mlir"
    },
    {
      "type": "modify",
      "old_id": "536e71d260f5683027761e4ef146d9cfb46f38ec",
      "old_mode": 33188,
      "old_path": "mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir",
      "new_id": "65e5fc0aff6aa3ef54e38963a10a4e08506738c7",
      "new_mode": 33188,
      "new_path": "mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir"
    },
    {
      "type": "modify",
      "old_id": "aee265e3faf175177c3e35c554490381ad88573e",
      "old_mode": 33188,
      "old_path": "mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir",
      "new_id": "2e59b7234e53df41051f54d96ffe524ac61eafa2",
      "new_mode": 33188,
      "new_path": "mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir"
    }
  ]
}