[flang-commits] [flang] [flang][cuda] Add option to flag descriptor I/O as error (PR #210170)

[Valentin Clement バレンタイン クレメン via flang-commits]

https://github.com/clementval updated https://github.com/llvm/llvm-project/pull/210170

>From a24af455d9ad6c5f690cfa9e2aa6f39b9425e2e4 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 16 Jul 2026 13:27:33 -0700
Subject: [PATCH 1/2] [flang][cuda] Add option to flag descriptor I/O as error

---
 .../flang/Optimizer/Transforms/Passes.td      |  9 +++++++--
 .../CUDA/CUFDeviceFuncTransform.cpp           | 19 +++++++++++++++++++
 2 files changed, 26 insertions(+), 2 deletions(-)

diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td
index ae7977d0670c7..d6cbb26882df5 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.td
+++ b/flang/include/flang/Optimizer/Transforms/Passes.td
@@ -584,8 +584,13 @@ def CUFDeviceFuncTransform
   let dependentDialects = ["mlir::gpu::GPUDialect",
                            "mlir::cf::ControlFlowDialect",
                            "mlir::NVVM::NVVMDialect"];
-  let options = [Option<"computeCap", "compute-capability", "int",
-                        /*default=*/"0", "CUDA compute capability version">];
+  let options =
+      [Option<"computeCap", "compute-capability", "int",
+              /*default=*/"0", "CUDA compute capability version">,
+       Option<"checkioOutputDescriptor", "check-io-output-descriptor", "bool",
+              /*default=*/"false",
+              "Report an error if device code contains a call to "
+              "descriptor I/O">];
 }
 
 def CUFFunctionRewrite : Pass<"cuf-function-rewrite", ""> {
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
index 24c0f3151970b..2576dee47d079 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceFuncTransform.cpp
@@ -254,6 +254,25 @@ class CUFDeviceFuncTransform
       });
     });
 
+    // Optionally report an error when device code calls the runtime function
+    // _FortranAioOutputDescriptor, which is not supported on the device.
+    if (checkioOutputDescriptor) {
+      constexpr llvm::StringRef aioOutputDescriptor =
+          "_FortranAioOutputDescriptor";
+      auto checkForAioOutputDescriptor = [&](fir::CallOp op) {
+        if (op.getCallee() && op.getCallee()->getLeafReference().getValue() ==
+                                  aioOutputDescriptor) {
+          op.emitError("descriptor I/O is not supported in device code");
+          signalPassFailure();
+        }
+      };
+      for (auto funcOp : deviceFuncs)
+        funcOp.walk(checkForAioOutputDescriptor);
+      mod.walk([&](cuf::KernelOp kernelOp) {
+        kernelOp.walk(checkForAioOutputDescriptor);
+      });
+    }
+
     for (auto funcOp : funcsToClone)
       gpuModSymTab.insert(funcOp->clone());
 

>From 04dffd213ef56906e46ec4de3f420c3f338c3824 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 16 Jul 2026 19:23:22 -0700
Subject: [PATCH 2/2] Add test

---
 .../CUDA/cuda-device-func-transform-aio.mlir  | 35 +++++++++++++++++++
 1 file changed, 35 insertions(+)
 create mode 100644 flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir

diff --git a/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir b/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir
new file mode 100644
index 0000000000000..fcc92919bbc93
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir
@@ -0,0 +1,35 @@
+// Without the option, calls to _FortranAioOutputDescriptor in device code are
+// not reported and the pass succeeds.
+// RUN: fir-opt --split-input-file --cuf-transform-device-func %s | FileCheck %s
+
+// With the option enabled, such calls trigger an error and pass failure.
+// RUN: fir-opt --split-input-file --cuf-transform-device-func="check-aio-output-descriptor=true" \
+// RUN:   --verify-diagnostics %s
+
+func.func private @_FortranAioOutputDescriptor(!fir.ref<i8>, !fir.box<none>) -> i1
+
+func.func @_QPsub_aio_device(%arg0: !fir.ref<i8>, %arg1: !fir.box<none>) attributes {cuf.proc_attr = #cuf.cuda_proc<device>} {
+  // expected-error at +1 {{descriptor I/O is not supported in device code}}
+  %0 = fir.call @_FortranAioOutputDescriptor(%arg0, %arg1) : (!fir.ref<i8>, !fir.box<none>) -> i1
+  return
+}
+
+// CHECK-LABEL: gpu.module @cuda_device_mod
+// CHECK: gpu.func @_QPsub_aio_device
+
+// -----
+
+func.func private @_FortranAioOutputDescriptor(!fir.ref<i8>, !fir.box<none>) -> i1
+
+func.func @_QPsub_aio_host(%arg0: !fir.ref<i8>, %arg1: !fir.box<none>) {
+  %c1 = arith.constant 1 : index
+  %c1_i32 = arith.constant 1 : i32
+  cuf.kernel<<<%c1_i32, %c1_i32>>> (%iv : index) = (%c1 : index) to (%c1 : index) step (%c1 : index) {
+    // expected-error at +1 {{call to _FortranAioOutputDescriptor is not supported in device code}}
+    %0 = fir.call @_FortranAioOutputDescriptor(%arg0, %arg1) : (!fir.ref<i8>, !fir.box<none>) -> i1
+    "fir.end"() : () -> ()
+  }
+  return
+}
+
+// CHECK-LABEL: func.func @_QPsub_aio_host