[flang-commits] [flang] 2a68f82 - Revert "[flang][cuda] Add cuf.register_kernel operation" (#112306)

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2024-10-14T21:06:59-07:00
New Revision: 2a68f829895cf849c5cd666dedd5d508f0052649

URL: https://github.com/llvm/llvm-project/commit/2a68f829895cf849c5cd666dedd5d508f0052649
DIFF: https://github.com/llvm/llvm-project/commit/2a68f829895cf849c5cd666dedd5d508f0052649.diff

LOG: Revert "[flang][cuda] Add cuf.register_kernel operation" (#112306)

Reverts llvm/llvm-project#112268

Added: 
    

Modified: 
    flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
    flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
    flang/test/Fir/cuf-invalid.fir
    flang/tools/fir-opt/fir-opt.cpp

Removed: 
    flang/test/Fir/CUDA/cuda-register-func.fir


################################################################################
diff  --git a/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td b/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
index 98d1ef529738c7..f643674f1d5d6b 100644
--- a/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
+++ b/flang/include/flang/Optimizer/Dialect/CUF/CUFOps.td
@@ -288,23 +288,4 @@ def cuf_KernelOp : cuf_Op<"kernel", [AttrSizedOperandSegments,
   let hasVerifier = 1;
 }
 
-def cuf_RegisterKernelOp : cuf_Op<"register_kernel", []> {
-  let summary = "Register a CUDA kernel";
-
-  let arguments = (ins
-    SymbolRefAttr:$name
-  );
-
-  let assemblyFormat = [{
-    $name attr-dict
-  }];
-
-  let hasVerifier = 1;
-
-  let extraClassDeclaration = [{
-    mlir::StringAttr getKernelName();
-    mlir::StringAttr getKernelModuleName();
-  }];
-}
-
 #endif // FORTRAN_DIALECT_CUF_CUF_OPS

diff  --git a/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp b/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
index 9e3bbd1f9cbee9..7fb2dcf4af115c 100644
--- a/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
+++ b/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
@@ -15,7 +15,6 @@
 #include "flang/Optimizer/Dialect/CUF/CUFDialect.h"
 #include "flang/Optimizer/Dialect/FIRAttr.h"
 #include "flang/Optimizer/Dialect/FIRType.h"
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/IR/BuiltinOps.h"
@@ -254,42 +253,6 @@ llvm::LogicalResult cuf::KernelOp::verify() {
   return mlir::success();
 }
 
-//===----------------------------------------------------------------------===//
-// RegisterKernelOp
-//===----------------------------------------------------------------------===//
-
-mlir::StringAttr cuf::RegisterKernelOp::getKernelModuleName() {
-  return getName().getRootReference();
-}
-
-mlir::StringAttr cuf::RegisterKernelOp::getKernelName() {
-  return getName().getLeafReference();
-}
-
-mlir::LogicalResult cuf::RegisterKernelOp::verify() {
-  if (getKernelName() == getKernelModuleName())
-    return emitOpError("expect a module and a kernel name");
-
-  auto mod = getOperation()->getParentOfType<mlir::ModuleOp>();
-  if (!mod)
-    return emitOpError("expect to be in a module");
-
-  mlir::SymbolTable symTab(mod);
-  auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(getKernelModuleName());
-  if (!gpuMod)
-    return emitOpError("gpu module not found");
-
-  mlir::SymbolTable gpuSymTab(gpuMod);
-  auto func = gpuSymTab.lookup<mlir::gpu::GPUFuncOp>(getKernelName());
-  if (!func)
-    return emitOpError("device function not found");
-
-  if (!func.isKernel())
-    return emitOpError("only kernel gpu.func can be registered");
-
-  return mlir::success();
-}
-
 // Tablegen operators
 
 #define GET_OP_CLASSES

diff  --git a/flang/test/Fir/CUDA/cuda-register-func.fir b/flang/test/Fir/CUDA/cuda-register-func.fir
deleted file mode 100644
index a428f68eb3bf42..00000000000000
--- a/flang/test/Fir/CUDA/cuda-register-func.fir
+++ /dev/null
@@ -1,20 +0,0 @@
-// RUN: fir-opt %s | FileCheck %s
-
-module attributes {gpu.container_module} {
-  gpu.module @cuda_device_mod {
-    gpu.func @_QPsub_device1() kernel {
-      gpu.return
-    }
-    gpu.func @_QPsub_device2(%arg0: !fir.ref<f32>) kernel {
-      gpu.return
-    }
-  }
-  llvm.func internal @__cudaFortranConstructor() {
-    cuf.register_kernel @cuda_device_mod::@_QPsub_device1
-    cuf.register_kernel @cuda_device_mod::@_QPsub_device2
-    llvm.return
-  }
-}
-
-// CHECK: cuf.register_kernel @cuda_device_mod::@_QPsub_device1
-// CHECK: cuf.register_kernel @cuda_device_mod::@_QPsub_device2

diff  --git a/flang/test/Fir/cuf-invalid.fir b/flang/test/Fir/cuf-invalid.fir
index a5747b8ee4a3b3..e9aeaa281e2a85 100644
--- a/flang/test/Fir/cuf-invalid.fir
+++ b/flang/test/Fir/cuf-invalid.fir
@@ -125,53 +125,3 @@ func.func @_QPsub1(%arg0: !fir.ref<!fir.array<?xf32>> {cuf.data_attr = #cuf.cuda
   cuf.data_transfer %20#0 to %11#0, %19 : !fir.shape<1> {transfer_kind = #cuf.cuda_transfer<host_device>} : !fir.box<!fir.array<?xf32>>, !fir.box<!fir.array<?xf32>>
   return
 }
-
-// -----
-
-module attributes {gpu.container_module} {
-  gpu.module @cuda_device_mod {
-    gpu.func @_QPsub_device1() {
-      gpu.return
-    }
-  }
-  llvm.func internal @__cudaFortranConstructor() {
-    // expected-error at +1{{'cuf.register_kernel' op only kernel gpu.func can be registered}}
-    cuf.register_kernel @cuda_device_mod::@_QPsub_device1
-    llvm.return
-  }
-}
-
-// -----
-
-module attributes {gpu.container_module} {
-  gpu.module @cuda_device_mod {
-    gpu.func @_QPsub_device1() {
-      gpu.return
-    }
-  }
-  llvm.func internal @__cudaFortranConstructor() {
-    // expected-error at +1{{'cuf.register_kernel' op device function not found}}
-    cuf.register_kernel @cuda_device_mod::@_QPsub_device2
-    llvm.return
-  }
-}
-
-// -----
-
-module attributes {gpu.container_module} {
-  llvm.func internal @__cudaFortranConstructor() {
-    // expected-error at +1{{'cuf.register_kernel' op gpu module not found}}
-    cuf.register_kernel @cuda_device_mod::@_QPsub_device1
-    llvm.return
-  }
-}
-
-// -----
-
-module attributes {gpu.container_module} {
-  llvm.func internal @__cudaFortranConstructor() {
-    // expected-error at +1{{'cuf.register_kernel' op expect a module and a kernel name}}
-    cuf.register_kernel @_QPsub_device1
-    llvm.return
-  }
-}

diff  --git a/flang/tools/fir-opt/fir-opt.cpp b/flang/tools/fir-opt/fir-opt.cpp
index 84a74770cf0303..f75fba27c68f08 100644
--- a/flang/tools/fir-opt/fir-opt.cpp
+++ b/flang/tools/fir-opt/fir-opt.cpp
@@ -42,7 +42,6 @@ int main(int argc, char **argv) {
 #endif
   DialectRegistry registry;
   fir::support::registerDialects(registry);
-  registry.insert<mlir::gpu::GPUDialect>();
   fir::support::addFIRExtensions(registry);
   return failed(MlirOptMain(argc, argv, "FIR modular optimizer driver\n",
       registry));