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

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-06-25T04:28:45Z
New Revision: 2fe65489978870631729ccb5c242c77f80f01a99

URL: https://github.com/llvm/llvm-project/commit/2fe65489978870631729ccb5c242c77f80f01a99
DIFF: https://github.com/llvm/llvm-project/commit/2fe65489978870631729ccb5c242c77f80f01a99.diff

LOG: [flang][cuda] Add NYI message for CUDA dynamic parallelism (#205628)

Added: 
    flang/test/Semantics/cuf30.cuf

Modified: 
    flang/lib/Semantics/check-cuda.cpp
    flang/lib/Semantics/expression.cpp

Removed: 
    


################################################################################
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 aea685e575754..02263fae0e55d 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3468,8 +3468,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/cuf30.cuf b/flang/test/Semantics/cuf30.cuf
new file mode 100644
index 0000000000000..04d063abc33d4
--- /dev/null
+++ b/flang/test/Semantics/cuf30.cuf
@@ -0,0 +1,19 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1 -Werror
+
+module m
+
+contains
+
+attributes(global) subroutine g1(a, m)
+  integer :: a(*), m
+  !ERROR: not yet implemented: CUDA dynamic parallelism
+  call g1<<<1, 1>>>(a, m)
+end subroutine
+
+attributes(global) subroutine g2(a, m)
+  integer :: a(*), m
+  !ERROR: not yet implemented: CUDA dynamic parallelism
+  call g1<<<1, 1>>>(a, m)
+end subroutine
+
+end module