NFC: Remove `Module::getFunctions` in favor of a general `getOps<T>`.

Modules can now contain more than just Functions, this just updates the iteration API to reflect that. The 'begin'/'end' methods have also been updated to iterate over opaque Operations.

PiperOrigin-RevId: 257099084

3 files changed, 3 insertions, 3 deletions

diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp
index da896afb090..1dbedf9fcee 100644
--- a/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp
@@ -64,7 +64,7 @@ public:
     LLVMInitializeNVPTXTargetMC();
     LLVMInitializeNVPTXAsmPrinter();
 
-    for (auto function : getModule()) {
+    for (auto function : getModule().getOps<FuncOp>()) {
       if (!gpu::GPUDialect::isKernel(function) || function.isExternal()) {
         continue;
       }
diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
index 0759324c77c..dafc5fa5730 100644
--- a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
@@ -130,7 +130,7 @@ public:
     // Cache the used LLVM types.
     initializeCachedTypes();
 
-    for (auto func : getModule()) {
+    for (auto func : getModule().getOps<FuncOp>()) {
       func.walk<mlir::gpu::LaunchFuncOp>(
           [this](mlir::gpu::LaunchFuncOp op) { translateGpuLaunchCalls(op); });
     }
diff --git a/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp b/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp
index 6de30433389..6306c567907 100644
--- a/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp
@@ -123,7 +123,7 @@ public:
     auto module = getModule();
     Builder builder(&getContext());
 
-    auto functions = module.getFunctions();
+    auto functions = module.getOps<FuncOp>();
     for (auto it = functions.begin(); it != functions.end();) {
       // Move iterator to after the current function so that potential insertion
       // of the accessor is after the kernel with cubin iself.