[flang-commits] [flang] [flang][cuda] Add option to flag descriptor I/O as error (PR #210170)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Thu Jul 16 23:07:38 PDT 2026
https://github.com/clementval updated https://github.com/llvm/llvm-project/pull/210170
>From 799a1f58939252223cad545de73b8edbe60b525d 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/3] [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 6353cdcf5cf86c202207778980744d9f97dd920f 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/3] 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
>From 5fdc5b7878e2e52328d8ff33672f4a864c4569b6 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 16 Jul 2026 23:07:19 -0700
Subject: [PATCH 3/3] Fix test option
---
flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir b/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir
index fcc92919bbc93..11b283bb5f1c7 100644
--- a/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir
+++ b/flang/test/Fir/CUDA/cuda-device-func-transform-aio.mlir
@@ -3,7 +3,7 @@
// 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: fir-opt --split-input-file --cuf-transform-device-func="check-io-output-descriptor=true" \
// RUN: --verify-diagnostics %s
func.func private @_FortranAioOutputDescriptor(!fir.ref<i8>, !fir.box<none>) -> i1
@@ -25,7 +25,7 @@ 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}}
+ // 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
"fir.end"() : () -> ()
}
More information about the flang-commits
mailing list