commit 887c276a6132832ea8330039992ff6ec0190542d
author Benoit Jacob <jacob.benoit.1@gmail.com> Thu Aug 01 10:04:57 2024 -0400
committer GitHub <noreply@github.com> Thu Aug 01 10:04:57 2024 -0400
tree 107dc48cb610ee8a39197364d76c7d8427f2c31a
parent 55e822841d532504908cefe2e8cd31162ad67da9

Fix `ConvertToNVVM` and drop local revert in `third_party/llvm-project` (#18064)

After debugging why the same testcase succeeds with
`iree-convert-to-llvm` but fails with `iree-convert-to-nvvm` unless a
local revert of https://github.com/llvm/llvm-project/commit/f6431f0c is
applied in `third_party/llvm-project`, comparing `-debug` logs, I found
that the difference was that `VectorToSCF` patterns were part of
`ConvertToLLVM` but not `ConvertToNVVM`.

Adding `VectorToSCF` to `ConvertToNVVM` fixes that and allows us to drop
the revert.

This requires adding the `AffineDialect` dependency similar to
https://github.com/iree-org/iree/pull/18062 because `VectorToSCF`
creates `affine.apply` ops.

---------

Signed-off-by: Benoit Jacob <jacob.benoit.1@gmail.com>
Signed-off-by: MaheshRavishankar <mahesh.ravishankar@gmail.com>
Co-authored-by: MaheshRavishankar <mahesh.ravishankar@gmail.com>

compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp

diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
index 05804ee..edcb2bb 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
@@ -44,8 +44,9 @@
 /// code.
 struct ConvertToNVVMPass : public ConvertToNVVMBase<ConvertToNVVMPass> {
   void getDependentDialects(DialectRegistry &registry) const override {
-    registry.insert<gpu::GPUDialect, IREE::GPU::IREEGPUDialect,
-                    LLVM::LLVMDialect, NVVM::NVVMDialect>();
+    registry
+        .insert<gpu::GPUDialect, IREE::GPU::IREEGPUDialect, LLVM::LLVMDialect,
+                NVVM::NVVMDialect, affine::AffineDialect>();
   }
   void runOnOperation() override {
     ModuleOp m = getOperation();
@@ -83,6 +84,8 @@
     // Run Vector -> Vector transformations ahead of conversion to LLVM.
     {
       RewritePatternSet patterns(&getContext());
+      populateVectorToSCFConversionPatterns(
+          patterns, VectorTransferToSCFOptions().enableFullUnroll());
       populateDropSharedMemoryDeallocOpPatterns(patterns);
       populateScalarizeMathOps(patterns);
       populateConvertSharedMemoryAllocOps(patterns);
@@ -144,8 +147,22 @@
       populateGpuToNVVMConversionPatterns(converter, llvmPatterns);
       populateNVGPUToNVVMConversionPatterns(converter, llvmPatterns);
       populateGpuWMMAToNVVMConversionPatterns(converter, llvmPatterns);
+
+      /// Target specification.
       LLVMConversionTarget target(getContext());
-      configureGpuToNVVMConversionLegality(target);
+      target.addIllegalOp<func::FuncOp>();
+      target.addLegalDialect<::mlir::LLVM::LLVMDialect>();
+      target.addLegalDialect<::mlir::NVVM::NVVMDialect>();
+      target.addIllegalDialect<gpu::GPUDialect>();
+      target.addIllegalOp<
+          LLVM::CopySignOp, LLVM::CosOp, LLVM::ExpOp, LLVM::Exp2Op,
+          LLVM::FAbsOp, LLVM::FCeilOp, LLVM::FFloorOp, LLVM::FRemOp,
+          LLVM::LogOp, LLVM::Log10Op, LLVM::Log2Op, LLVM::PowOp,
+          LLVM::RoundEvenOp, LLVM::RoundOp, LLVM::SinOp, LLVM::SqrtOp>();
+
+      // TODO: Remove once we support replacing non-root ops.
+      target.addLegalOp<gpu::YieldOp, gpu::GPUModuleOp, gpu::ModuleEndOp>();
+
       if (failed(applyPartialConversion(m, target, std::move(llvmPatterns)))) {
         signalPassFailure();
       }

third_party/llvm-project

diff --git a/third_party/llvm-project b/third_party/llvm-project
index 5ccb8d8..d8b985c 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 5ccb8d8a17285cacbb9ce11b0a8765194f243ccc
+Subproject commit d8b985c9490664f7ec923e59cf6c603d998179ae