Skip to content

[MLIR][NVVM] Declare InferIntRangeInterface for RangeableRegisterOp #122263

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Jan 10, 2025

Conversation

grypp
Copy link
Member

@grypp grypp commented Jan 9, 2025

SpecialRangeableRegister NVVM operations may have a range attribute set. When this attribute is present, it becomes possible to determine their range.

This PR declares the InferIntRangeInterface for all SpecialRangeableRegister operations and extracts range information from the range attribute.

…gisterOp`

`SpecialRangeableRegister` NVVM operations may have a `range` attribute set. When this attribute is present, it becomes possible to determine their range.

This PR declares the `InferIntRangeInterface` for all `SpecialRangeableRegister` operations and extracts range information from the `range` attribute.
@llvmbot
Copy link
Member

llvmbot commented Jan 9, 2025

@llvm/pr-subscribers-mlir-llvm

@llvm/pr-subscribers-mlir

Author: Guray Ozen (grypp)

Changes

SpecialRangeableRegister NVVM operations may have a range attribute set. When this attribute is present, it becomes possible to determine their range.

This PR declares the InferIntRangeInterface for all SpecialRangeableRegister operations and extracts range information from the range attribute.


Full diff: https://github.com/llvm/llvm-project/pull/122263.diff

3 Files Affected:

  • (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h (+1)
  • (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+18-2)
  • (added) mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir (+28)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
index 4fd00ff929bd70..50d1a39126ea3e 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
@@ -19,6 +19,7 @@
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
+#include "mlir/Interfaces/InferIntRangeInterface.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "llvm/IR/IntrinsicsNVPTX.h"
 
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index a2d2102b59dece..d0d720e664ce3b 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -18,6 +18,7 @@ include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
 include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
+include "mlir/Interfaces/InferIntRangeInterface.td"
 
 def LLVM_PointerGeneric : LLVM_PointerInAddressSpace<0>;
 def LLVM_PointerGlobal : LLVM_PointerInAddressSpace<1>;
@@ -134,8 +135,8 @@ class NVVM_SpecialRegisterOp<string mnemonic, list<Trait> traits = []> :
   let assemblyFormat = "attr-dict `:` type($res)";
 }
 
-class NVVM_SpecialRangeableRegisterOp<string mnemonic, list<Trait> traits = []> :
-  NVVM_SpecialRegisterOp<mnemonic, traits> {
+class NVVM_SpecialRangeableRegisterOp<string mnemonic> :
+  NVVM_SpecialRegisterOp<mnemonic, [DeclareOpInterfaceMethods<InferIntRangeInterface, ["inferResultRanges"]>]> {
   let arguments = (ins OptionalAttr<LLVM_ConstantRangeAttr>:$range);
   let assemblyFormat = "(`range` $range^)? attr-dict `:` type($res)";
   let llvmBuilder = baseLlvmBuilder # setRangeRetAttrCode # baseLlvmBuilderCoda;
@@ -147,6 +148,21 @@ class NVVM_SpecialRangeableRegisterOp<string mnemonic, list<Trait> traits = []>
       build($_builder, $_state, resultType, ::mlir::LLVM::ConstantRangeAttr{});
     }]>
   ];
+
+  // Define this method for the InferIntRangeInterface.
+  let extraClassDefinition = [{
+    // Infer the result ranges based on the range attribute.
+    void $cppClass::inferResultRanges(
+        ArrayRef<::mlir::ConstantIntRanges> argRanges,
+        SetIntRangeFn setResultRanges) {
+        if (auto rangeAttr = getOperation()->getAttrOfType<LLVM::ConstantRangeAttr>("range")) {
+          setResultRanges(getResult(), 
+                          {rangeAttr.getLower(), rangeAttr.getUpper(),
+                          rangeAttr.getLower(), rangeAttr.getUpper()});
+        }
+    }
+  }];
+
 }
 
 //===----------------------------------------------------------------------===//
diff --git a/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir b/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir
new file mode 100644
index 00000000000000..7014ff8aaa80b7
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/nvvm-test-range.mlir
@@ -0,0 +1,28 @@
+// RUN: mlir-opt -int-range-optimizations -canonicalize %s | FileCheck %s
+gpu.module @module{
+    gpu.func @kernel_1() kernel {
+        %tidx = nvvm.read.ptx.sreg.tid.x range <i32, 0, 32> : i32
+        %tidy = nvvm.read.ptx.sreg.tid.y range <i32, 0, 128> : i32
+        %tidz = nvvm.read.ptx.sreg.tid.z range <i32, 0, 4> : i32
+        %c64 = arith.constant 64 : i32
+        
+        %1 = arith.cmpi sgt, %tidx, %c64 : i32
+        scf.if %1 {
+            gpu.printf "threadidx"
+        }
+        %2 = arith.cmpi sgt, %tidy, %c64 : i32
+        scf.if %2 {
+            gpu.printf "threadidy"
+        }
+        %3 = arith.cmpi sgt, %tidz, %c64 : i32
+        scf.if %3 {
+            gpu.printf "threadidz"
+        }
+        gpu.return
+    }
+}
+
+// CHECK-LABEL: gpu.func @kernel_1
+// CHECK-NOT: gpu.printf "threadidx"
+// CHECK: gpu.printf "threadidy"
+// CHECK-NOT: gpu.printf "threadidz"
\ No newline at end of file

@grypp grypp requested review from krzysz00 and durga4github January 9, 2025 12:08
Copy link
Contributor

@Dinistro Dinistro left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Dropped a few drive by comments

Copy link

github-actions bot commented Jan 9, 2025

✅ With the latest revision this PR passed the C/C++ code formatter.

@grypp grypp merged commit 66e41a1 into llvm:main Jan 10, 2025
8 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented Jan 10, 2025

LLVM Buildbot has detected a new failure on builder flang-aarch64-libcxx running on linaro-flang-aarch64-libcxx while building mlir at step 5 "build-unified-tree".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/89/builds/14051

Here is the relevant piece of the build log for the reference
Step 5 (build-unified-tree) failure: build (failure)
...
94.903 [2197/15/5097] Copying clang's ppc_wrappers/smmintrin.h...
94.903 [2197/14/5098] Copying clang's ppc_wrappers/nmmintrin.h...
94.904 [2197/13/5099] Copying clang's ppc_wrappers/bmiintrin.h...
94.904 [2197/12/5100] Copying clang's ppc_wrappers/bmi2intrin.h...
94.905 [2197/11/5101] Copying clang's ppc_wrappers/immintrin.h...
94.905 [2197/10/5102] Copying clang's ppc_wrappers/x86intrin.h...
94.905 [2197/9/5103] Copying clang's ppc_wrappers/x86gprintrin.h...
94.906 [2197/8/5104] Copying clang's openmp_wrappers/math.h...
95.290 [2197/7/5105] Building CXX object tools/mlir/tools/mlir-lsp-server/CMakeFiles/mlir-lsp-server.dir/mlir-lsp-server.cpp.o
95.313 [2197/6/5106] Linking CXX shared library lib/libMLIRNVVMDialect.so.20.0git
FAILED: lib/libMLIRNVVMDialect.so.20.0git 
: && /usr/local/bin/c++ -fPIC -stdlib=libc++ -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wc++98-compat-extra-semi -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wsuggest-override -Wstring-conversion -Wmisleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -ffunction-sections -fdata-sections -Wundef -Werror=mismatched-tags -Werror=global-constructors -O3 -DNDEBUG  -stdlib=libc++ -Wl,-z,defs -Wl,-z,nodelete   -Wl,-rpath-link,/home/tcwg-buildbot/worker/flang-aarch64-libcxx/build/./lib  -Wl,--gc-sections -shared -Wl,-soname,libMLIRNVVMDialect.so.20.0git -o lib/libMLIRNVVMDialect.so.20.0git tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/BasicPtxBuilderInterface.cpp.o  -Wl,-rpath,"\$ORIGIN/../lib:/home/tcwg-buildbot/worker/flang-aarch64-libcxx/build/lib:"  lib/libMLIRLLVMDialect.so.20.0git  lib/libMLIRSideEffectInterfaces.so.20.0git  lib/libMLIRCallInterfaces.so.20.0git  lib/libMLIRControlFlowInterfaces.so.20.0git  lib/libMLIRDataLayoutInterfaces.so.20.0git  lib/libMLIRFunctionInterfaces.so.20.0git  lib/libMLIRInferTypeOpInterface.so.20.0git  lib/libMLIRMemorySlotInterfaces.so.20.0git  lib/libMLIRIR.so.20.0git  lib/libMLIRSupport.so.20.0git  lib/libLLVMBitWriter.so.20.0git  lib/libLLVMAsmParser.so.20.0git  lib/libLLVMBitReader.so.20.0git  lib/libLLVMCore.so.20.0git  lib/libLLVMBinaryFormat.so.20.0git  lib/libLLVMSupport.so.20.0git  -Wl,-rpath-link,/home/tcwg-buildbot/worker/flang-aarch64-libcxx/build/lib && :
/usr/bin/ld: tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o: in function `mlir::detail::InferIntRangeInterfaceInterfaceTraits::Model<mlir::NVVM::BlockDimXOp>::inferResultRangesFromOptional(mlir::detail::InferIntRangeInterfaceInterfaceTraits::Concept const*, mlir::Operation*, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)':
NVVMDialect.cpp:(.text._ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM11BlockDimXOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE[_ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM11BlockDimXOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE]+0x8c): undefined reference to `mlir::intrange::detail::defaultInferResultRanges(mlir::InferIntRangeInterface, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)'
/usr/bin/ld: tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o: in function `mlir::detail::InferIntRangeInterfaceInterfaceTraits::Model<mlir::NVVM::BlockDimYOp>::inferResultRangesFromOptional(mlir::detail::InferIntRangeInterfaceInterfaceTraits::Concept const*, mlir::Operation*, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)':
NVVMDialect.cpp:(.text._ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM11BlockDimYOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE[_ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM11BlockDimYOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE]+0x8c): undefined reference to `mlir::intrange::detail::defaultInferResultRanges(mlir::InferIntRangeInterface, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)'
/usr/bin/ld: tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o: in function `mlir::detail::InferIntRangeInterfaceInterfaceTraits::Model<mlir::NVVM::BlockDimZOp>::inferResultRangesFromOptional(mlir::detail::InferIntRangeInterfaceInterfaceTraits::Concept const*, mlir::Operation*, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)':
NVVMDialect.cpp:(.text._ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM11BlockDimZOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE[_ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM11BlockDimZOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE]+0x8c): undefined reference to `mlir::intrange::detail::defaultInferResultRanges(mlir::InferIntRangeInterface, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)'
/usr/bin/ld: tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o: in function `mlir::detail::InferIntRangeInterfaceInterfaceTraits::Model<mlir::NVVM::BlockIdXOp>::inferResultRangesFromOptional(mlir::detail::InferIntRangeInterfaceInterfaceTraits::Concept const*, mlir::Operation*, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)':
NVVMDialect.cpp:(.text._ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM10BlockIdXOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE[_ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM10BlockIdXOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE]+0x8c): undefined reference to `mlir::intrange::detail::defaultInferResultRanges(mlir::InferIntRangeInterface, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)'
/usr/bin/ld: tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o: in function `mlir::detail::InferIntRangeInterfaceInterfaceTraits::Model<mlir::NVVM::BlockIdYOp>::inferResultRangesFromOptional(mlir::detail::InferIntRangeInterfaceInterfaceTraits::Concept const*, mlir::Operation*, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)':
NVVMDialect.cpp:(.text._ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM10BlockIdYOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE[_ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM10BlockIdYOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE]+0x8c): undefined reference to `mlir::intrange::detail::defaultInferResultRanges(mlir::InferIntRangeInterface, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)'
/usr/bin/ld: tools/mlir/lib/Dialect/LLVMIR/CMakeFiles/obj.MLIRNVVMDialect.dir/IR/NVVMDialect.cpp.o:NVVMDialect.cpp:(.text._ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM10BlockIdZOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE[_ZN4mlir6detail37InferIntRangeInterfaceInterfaceTraits5ModelINS_4NVVM10BlockIdZOpEE29inferResultRangesFromOptionalEPKNS1_7ConceptEPNS_9OperationEN4llvm8ArrayRefINS_17IntegerValueRangeEEENSB_12function_refIFvNS_5ValueERKSD_EEE]+0x8c): more undefined references to `mlir::intrange::detail::defaultInferResultRanges(mlir::InferIntRangeInterface, llvm::ArrayRef<mlir::IntegerValueRange>, llvm::function_ref<void (mlir::Value, mlir::IntegerValueRange const&)>)' follow
clang++: error: linker command failed with exit code 1 (use -v to see invocation)
126.957 [2197/5/5107] Building CXX object tools/mlir/tools/mlir-reduce/CMakeFiles/mlir-reduce.dir/mlir-reduce.cpp.o
136.759 [2197/4/5108] Building CXX object tools/mlir/lib/CAPI/RegisterEverything/CMakeFiles/obj.MLIRCAPIRegisterEverything.dir/RegisterEverything.cpp.o
147.288 [2197/3/5109] Building CXX object tools/mlir/tools/mlir-opt/CMakeFiles/mlir-opt.dir/mlir-opt.cpp.o
147.396 [2197/2/5110] Building CXX object tools/mlir/tools/mlir-opt/CMakeFiles/MLIRMlirOptMain.dir/mlir-opt.cpp.o
170.894 [2197/1/5111] Building CXX object tools/mlir/examples/transform-opt/CMakeFiles/mlir-transform-opt.dir/mlir-transform-opt.cpp.o
ninja: build stopped: subcommand failed.

@jplehr
Copy link
Contributor

jplehr commented Jan 10, 2025

I also see a breakage in https://lab.llvm.org/staging/#/builders/105/builds/13077

@jplehr
Copy link
Contributor

jplehr commented Jan 10, 2025

1ef2580 did not fix the build issue

@jplehr
Copy link
Contributor

jplehr commented Jan 10, 2025

Thank you.
2e6030e fixes the issue for me.

@grypp
Copy link
Member Author

grypp commented Jan 10, 2025

Thank you. 2e6030e fixes the issue for me.

Yes, this one should fix it.
Sorry for disturbance. I could not reproduce the bug locally in the first place.

BaiXilin pushed a commit to BaiXilin/llvm-fix-vnni-instr-types that referenced this pull request Jan 12, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

6 participants