[all-commits] [llvm/llvm-project] 9a553d: [MLIR][NVVM] Add NVVMRequiresSM op traits (#126886)

[Srinivasa Ravi via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 9a553d3766aacb69e884823da92dedff264e3f0f
      https://github.com/llvm/llvm-project/commit/9a553d3766aacb69e884823da92dedff264e3f0f
  Author: Srinivasa Ravi <srinivasar at nvidia.com>
  Date:   2025-05-21 (Wed, 21 May 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
    M mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
    M mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
    M mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    A mlir/include/mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.h
    A mlir/include/mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.td
    M mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
    M mlir/lib/Dialect/LLVMIR/CMakeLists.txt
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    A mlir/lib/Dialect/LLVMIR/IR/NVVMRequiresSMTraits.cpp
    A mlir/test/Dialect/LLVMIR/nvvm-check-targetSM.mlir
    M mlir/test/lib/Dialect/Test/CMakeLists.txt
    M mlir/test/lib/Dialect/Test/TestOps.h
    M mlir/test/lib/Dialect/Test/TestOps.td
    M utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

  Log Message:
  -----------
  [MLIR][NVVM] Add NVVMRequiresSM op traits (#126886)

Motivation:
Currently, the NVVMOps are not verified against the supported SM
architectures. This can manifest as an ISel failure in the NVPTX LLVM
backend during CodeGen to PTX ISA. This PR addresses this issue by
adding verifier checks for Target-SM architectures in the NVVM Dialect
itself, thereby catching the errors early on.

Summary:
* Parametric traits named `NVVMRequiresSM` and `NVVMRequiresSMa` are
added to facilitate the version checks for typical and arch-accelerated
versions respectively.
* These traits can be attached to any NVVM Op to enable the checks for
the particular Op. (example shown below)
* An attribute interface called named `TargetAttrVerifyInterface` is
added to the GPU dialect which any target attribute seeking to perform
target-verification on the module can implement.
* The checks are performed by the `NVVMTargetAttr` (implementing the
`TargetAttrVerifyInterface` interface) when called from the GPU module
verifier where it walks through the module and performs the checks for
Ops with the `NVVMRequiresSM` traits.
* A few Ops in `NVVMOps.td` have been updated to serve as examples.

Example Usage:
```
       def NVVM_ReduxOp : NVVM_Op<"redux.sync"> {...} 
 ----> def NVVM_ReduxOp : NVVM_Op<"redux.sync", [NVVMRequiresSM<80>]> {...}

       def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned"> {...}
 ----> def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", [NVVMRequiresSMa<[90]>]> {...}
```

---------

Co-authored-by: Guray Ozen <guray.ozen at gmail.com>



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