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

Reverts #112268

Author: clementval

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

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

flang/test/Fir/CUDA/cuda-register-func.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@+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@+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@+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@+1{{'cuf.register_kernel' op expect a module and a kernel name}}
-    cuf.register_kernel @_QPsub_device1
-    llvm.return
-  }
-}

flang/test/Fir/cuf-invalid.fir

@@ -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));