[Flang] Add missing dependent dialects to MLIR passes

[skatrak]

This patch updates several passes to include the DLTI dialect, since their use of the fir::support::getOrSetMLIRDataLayout() utility function could, in some cases, require this dialect to be loaded in advance.

Also, the CUFComputeSharedMemoryOffsetsAndSize pass has been updated with a dependency to the GPU dialect, as its invocation to cuf::getOrCreateGPUModule() would result in the same kind of error if no other operations or attributes from that dialect were present in the input MLIR module.

[llvmbot]

@llvm/pr-subscribers-flang-fir-hlfir

Author: Sergio Afonso (skatrak)

Changes

This patch updates several passes to include the DLTI dialect, since their use of the fir::support::getOrSetMLIRDataLayout() utility function could, in some cases, require this dialect to be loaded in advance.

Also, the CUFComputeSharedMemoryOffsetsAndSize pass has been updated with a dependency to the GPU dialect, as its invocation to cuf::getOrCreateGPUModule() would result in the same kind of error if no other operations or attributes from that dialect were present in the input MLIR module.

---

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

7 Files Affected:

• (modified) flang/include/flang/Optimizer/Transforms/Passes.td (+8-5)
• (modified) flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp (+1)
• (modified) flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp (+1)
• (modified) flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp (+1)
• (modified) flang/lib/Optimizer/Transforms/CUFOpConversion.cpp (+1)
• (modified) flang/lib/Optimizer/Transforms/LoopVersioning.cpp (+1)
• (added) flang/test/Transforms/dlti-dependency.fir (+21)

diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td
index 3243b44df9c7a..c0d88a8e19f80 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.td
+++ b/flang/include/flang/Optimizer/Transforms/Passes.td
@@ -356,7 +356,7 @@ def LoopVersioning : Pass<"loop-versioning", "mlir::func::FuncOp"> {
     an array has element sized stride. The element sizes stride allows some
     loops to be vectorized as well as other loop optimizations.
   }];
-  let dependentDialects = [ "fir::FIROpsDialect" ];
+  let dependentDialects = [ "fir::FIROpsDialect", "mlir::DLTIDialect" ];
 }
 
 def VScaleAttr : Pass<"vscale-attr", "mlir::func::FuncOp"> {
@@ -436,7 +436,7 @@ def AssumedRankOpConversion : Pass<"fir-assumed-rank-op", "mlir::ModuleOp"> {
 def CUFOpConversion : Pass<"cuf-convert", "mlir::ModuleOp"> {
   let summary = "Convert some CUF operations to runtime calls";
   let dependentDialects = [
-    "fir::FIROpsDialect", "mlir::gpu::GPUDialect"
+    "fir::FIROpsDialect", "mlir::gpu::GPUDialect", "mlir::DLTIDialect"
   ];
 }
 
@@ -451,14 +451,14 @@ def CUFDeviceGlobal :
 def CUFAddConstructor : Pass<"cuf-add-constructor", "mlir::ModuleOp"> {
   let summary = "Add constructor to register CUDA Fortran allocators";
   let dependentDialects = [
-    "cuf::CUFDialect", "mlir::func::FuncDialect"
+    "cuf::CUFDialect", "mlir::func::FuncDialect", "mlir::DLTIDialect"
   ];
 }
 
 def CUFGPUToLLVMConversion : Pass<"cuf-gpu-convert-to-llvm", "mlir::ModuleOp"> {
   let summary = "Convert some GPU operations lowered from CUF to runtime calls";
   let dependentDialects = [
-    "mlir::LLVM::LLVMDialect"
+    "mlir::LLVM::LLVMDialect", "mlir::DLTIDialect"
   ];
 }
 
@@ -472,7 +472,10 @@ def CUFComputeSharedMemoryOffsetsAndSize
     the global and set it.
   }];
 
-  let dependentDialects = ["cuf::CUFDialect", "fir::FIROpsDialect"];
+  let dependentDialects = [
+    "cuf::CUFDialect", "fir::FIROpsDialect", "mlir::gpu::GPUDialect",
+    "mlir::DLTIDialect"
+  ];
 }
 
 def SetRuntimeCallAttributes
diff --git a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp
index 064f0f363f699..2dd6950b34897 100644
--- a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp
+++ b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp
@@ -22,6 +22,7 @@
 #include "flang/Optimizer/Support/DataLayout.h"
 #include "flang/Runtime/CUDA/registration.h"
 #include "flang/Runtime/entry-names.h"
+#include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMAttrs.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
diff --git a/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp b/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp
index 8009522a82e27..f6381ef8a8a21 100644
--- a/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp
+++ b/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp
@@ -22,6 +22,7 @@
 #include "flang/Optimizer/Support/DataLayout.h"
 #include "flang/Runtime/CUDA/registration.h"
 #include "flang/Runtime/entry-names.h"
+#include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Value.h"
diff --git a/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp b/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp
index 2549fdcb8baee..fe69ffa8350af 100644
--- a/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp
+++ b/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp
@@ -14,6 +14,7 @@
 #include "flang/Runtime/CUDA/common.h"
 #include "flang/Support/Fortran.h"
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
+#include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Pass/Pass.h"
diff --git a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp
index e70ceb3a67d98..7477a3c53c3ef 100644
--- a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp
+++ b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp
@@ -24,6 +24,7 @@
 #include "flang/Runtime/allocatable.h"
 #include "flang/Support/Fortran.h"
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
+#include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/Pass/Pass.h"
diff --git a/flang/lib/Optimizer/Transforms/LoopVersioning.cpp b/flang/lib/Optimizer/Transforms/LoopVersioning.cpp
index 42e149bb3dba2..94dd8db2bafcb 100644
--- a/flang/lib/Optimizer/Transforms/LoopVersioning.cpp
+++ b/flang/lib/Optimizer/Transforms/LoopVersioning.cpp
@@ -51,6 +51,7 @@
 #include "flang/Optimizer/Dialect/Support/KindMapping.h"
 #include "flang/Optimizer/Support/DataLayout.h"
 #include "flang/Optimizer/Transforms/Passes.h"
+#include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Dominance.h"
 #include "mlir/IR/Matchers.h"
diff --git a/flang/test/Transforms/dlti-dependency.fir b/flang/test/Transforms/dlti-dependency.fir
new file mode 100644
index 0000000000000..c1c3da19fb8d6
--- /dev/null
+++ b/flang/test/Transforms/dlti-dependency.fir
@@ -0,0 +1,21 @@
+// This test only makes sure that passes with a DLTI dialect dependency are able
+// to obtain the dlti.dl_spec module attribute from an llvm.data_layout string.
+//
+// If dependencies for the pass are not properly set, this test causes a
+// compiler error due to the DLTI dialect not being loaded.
+
+// RUN: fir-opt --add-debug-info %s
+// RUN: fir-opt --cuf-add-constructor %s
+// RUN: fir-opt --cuf-compute-shared-memory %s
+// RUN: fir-opt --cuf-gpu-convert-to-llvm %s
+// RUN: fir-opt --cuf-convert %s
+// RUN: fir-opt --loop-versioning %s
+
+module attributes {llvm.data_layout = "e-m:e-p:64:64-i64:64-i128:128-n32:64-S128"} {
+  llvm.func @foo(%arg0 : i32) {
+    llvm.return
+  }
+}
+
+// CHECK: module attributes {
+// CHECK-SAME: dlti.dl_spec = #dlti.dl_spec<

[tblah]

LGTM but please wait for @clementval