[flang-commits] [flang] Revert "[flang][cuda] Add cuf.register_kernel operation" (PR #112306)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Mon Oct 14 21:06:55 PDT 2024
https://github.com/clementval created https://github.com/llvm/llvm-project/pull/112306
Reverts llvm/llvm-project#112268
>From b106c5edf89c5306de07766078f8449a28a39fc3 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Valentin=20Clement=20=28=E3=83=90=E3=83=AC=E3=83=B3?=
=?UTF-8?q?=E3=82=BF=E3=82=A4=E3=83=B3=20=E3=82=AF=E3=83=AC=E3=83=A1?=
=?UTF-8?q?=E3=83=B3=29?= <clementval at gmail.com>
Date: Mon, 14 Oct 2024 21:06:42 -0700
Subject: [PATCH] Revert "[flang][cuda] Add cuf.register_kernel operation
(#112268)"
This reverts commit cbe76a2ac3547258076cc93e8cbc42cdc6219d06.
---
.../flang/Optimizer/Dialect/CUF/CUFOps.td | 19 -------
flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp | 37 --------------
flang/test/Fir/CUDA/cuda-register-func.fir | 20 --------
flang/test/Fir/cuf-invalid.fir | 50 -------------------
flang/tools/fir-opt/fir-opt.cpp | 1 -
5 files changed, 127 deletions(-)
delete mode 100644 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));
More information about the flang-commits
mailing list