[flang-commits] [flang] [flang][cuda] Relax the verifier for cuf.register_kernel op (PR #112585)

[via flang-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

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

Author: Valentin Clement (バレンタイン クレメン) (clementval)

<details>
<summary>Changes</summary>

Relax the verifier since the `gpu.func` might be converted to `llvm.func` before `cuf.register_kernel` is converted.

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


2 Files Affected:

- (modified) flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp (+6-6) 
- (modified) flang/test/Fir/cuf-invalid.fir (-15) 


``````````diff
diff --git a/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp b/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
index 9e3bbd1f9cbee9..5f7232e3dd4b09 100644
--- a/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
+++ b/flang/lib/Optimizer/Dialect/CUF/CUFOps.cpp
@@ -281,12 +281,12 @@ mlir::LogicalResult cuf::RegisterKernelOp::verify() {
 
   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");
-
+  if (func) {
+    // Only check if the gpu.func is found. It might be converted to LLVMFuncOp
+    // already.
+    if (!func.isKernel())
+      return emitOpError("only kernel gpu.func can be registered");
+  }
   return mlir::success();
 }
 
diff --git a/flang/test/Fir/cuf-invalid.fir b/flang/test/Fir/cuf-invalid.fir
index a5747b8ee4a3b3..ac488ab3e5abf5 100644
--- a/flang/test/Fir/cuf-invalid.fir
+++ b/flang/test/Fir/cuf-invalid.fir
@@ -143,21 +143,6 @@ module attributes {gpu.container_module} {
 
 // -----
 
-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}}

``````````

</details>


https://github.com/llvm/llvm-project/pull/112585