[flang-commits] [flang] [flang][cuda] Add NYI message for CUDA dynamic parallelism (PR #205628)

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

Valentin Clement =?utf-8?b?KOODkOODrOODsw=?=,
Valentin Clement =?utf-8?b?KOODkOODrOODsw=?=,Valentin Clement
 <clementval at gmail.com>
Message-ID:
In-Reply-To: <llvm.org/llvm/llvm-project/pull/205628 at github.com>


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

>From 589d7d8b463404ec54d80760472bd6948317d602 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Wed, 24 Jun 2026 11:07:07 -0700
Subject: [PATCH 1/4] [flang][cuda] Add NYI message for CUDA dynamic
 parallelism

---
 flang/lib/Semantics/check-cuda.cpp |  4 ++++
 flang/lib/Semantics/expression.cpp | 12 ++++++++++--
 flang/test/Semantics/cuf29.cuf     |  7 +++++++
 3 files changed, 21 insertions(+), 2 deletions(-)
 create mode 100644 flang/test/Semantics/cuf29.cuf

diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 7d2c012d3a4dc..32ba35e83a120 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -92,6 +92,10 @@ struct DeviceExprChecker
             }
             return {};
           }
+          if (*attrs == common::CUDASubprogramAttrs::Global) {
+            return parser::MessageFormattedText(
+                "not yet implemented: CUDA dynamic parallelism"_err_en_US);
+          }
         }
       }
 
diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index 6c0a21cc769c1..06bcb7134f127 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3429,8 +3429,16 @@ void ExpressionAnalyzer::CheckForBadRecursion(
             "Assumed-length CHARACTER(*) function '%s' cannot call itself"_err_en_US,
             callSite);
       } else if (FindCUDADeviceContext(scope)) {
-        msg = Say(
-            "Device subprogram '%s' cannot call itself"_err_en_US, callSite);
+        const auto *subp{
+            proc.GetUltimate().detailsIf<semantics::SubprogramDetails>()};
+        bool isGlobalCUDA{subp && subp->cudaSubprogramAttrs() &&
+            *subp->cudaSubprogramAttrs() ==
+                common::CUDASubprogramAttrs::Global};
+        // CUDA global call diagnostics are handled by CUDA checks.
+        if (!isGlobalCUDA) {
+          msg = Say(
+              "Device subprogram '%s' cannot call itself"_err_en_US, callSite);
+        }
       }
       AttachDeclaration(msg, proc);
     }
diff --git a/flang/test/Semantics/cuf29.cuf b/flang/test/Semantics/cuf29.cuf
new file mode 100644
index 0000000000000..9eea2c309fa68
--- /dev/null
+++ b/flang/test/Semantics/cuf29.cuf
@@ -0,0 +1,7 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1 -Werror
+
+attributes(global) subroutine g1(a, m)
+  integer :: a(*), m
+  !ERROR: not yet implemented: CUDA dynamic parallelism
+  call g1<<<1, 1>>>(a, m)
+end subroutine

>From 9419a53331580fde2f467715766df7e0b34760e6 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: Wed, 24 Jun 2026 14:07:06 -0700
Subject: [PATCH 2/4] Add non recursive test case

---
 flang/test/Semantics/cuf29.cuf | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/flang/test/Semantics/cuf29.cuf b/flang/test/Semantics/cuf29.cuf
index 9eea2c309fa68..2dca5de501644 100644
--- a/flang/test/Semantics/cuf29.cuf
+++ b/flang/test/Semantics/cuf29.cuf
@@ -5,3 +5,9 @@ attributes(global) subroutine g1(a, m)
   !ERROR: not yet implemented: CUDA dynamic parallelism
   call g1<<<1, 1>>>(a, m)
 end subroutine
+
+attributes(global) subroutine g1(a, m)
+  integer :: a(*), m
+  !ERROR: not yet implemented: CUDA dynamic parallelism
+  call g1<<<1, 1>>>(a, m)
+end subroutine

>From 3e26fc0c5f48d7e6eee0b3a0a9b141f9298c070d 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: Wed, 24 Jun 2026 14:15:28 -0700
Subject: [PATCH 3/4] Fix name

---
 flang/test/Semantics/cuf29.cuf | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/flang/test/Semantics/cuf29.cuf b/flang/test/Semantics/cuf29.cuf
index 2dca5de501644..cfe1af329ff24 100644
--- a/flang/test/Semantics/cuf29.cuf
+++ b/flang/test/Semantics/cuf29.cuf
@@ -6,7 +6,7 @@ attributes(global) subroutine g1(a, m)
   call g1<<<1, 1>>>(a, m)
 end subroutine
 
-attributes(global) subroutine g1(a, m)
+attributes(global) subroutine g2(a, m)
   integer :: a(*), m
   !ERROR: not yet implemented: CUDA dynamic parallelism
   call g1<<<1, 1>>>(a, m)

>From 60c7f73bb63760add12a6596c232648bb042bf3e Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Wed, 24 Jun 2026 18:57:16 -0700
Subject: [PATCH 4/4] move cuf29 to cuf30

---
 flang/test/Semantics/{cuf29.cuf => cuf30.cuf} | 0
 1 file changed, 0 insertions(+), 0 deletions(-)
 rename flang/test/Semantics/{cuf29.cuf => cuf30.cuf} (100%)

diff --git a/flang/test/Semantics/cuf29.cuf b/flang/test/Semantics/cuf30.cuf
similarity index 100%
rename from flang/test/Semantics/cuf29.cuf
rename to flang/test/Semantics/cuf30.cuf