[flang-commits] [flang] [flang][cuda] Add NYI message for CUDA dynamic parallelism (PR #205628)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Wed Jun 24 21:05:51 PDT 2026
Valentin Clement =?utf-8?b?KOODkOODrOODsw=?=,
Valentin Clement =?utf-8?b?KOODkOODrOODsw=?=,Valentin Clement
<clementval at gmail.com>,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/5] [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/5] 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/5] 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/5] 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
>From 52515b35aff521977dd219b10797aa1e5ae845f3 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Wed, 24 Jun 2026 21:05:36 -0700
Subject: [PATCH 5/5] Fix test
---
flang/test/Semantics/cuf30.cuf | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/flang/test/Semantics/cuf30.cuf b/flang/test/Semantics/cuf30.cuf
index cfe1af329ff24..04d063abc33d4 100644
--- a/flang/test/Semantics/cuf30.cuf
+++ b/flang/test/Semantics/cuf30.cuf
@@ -1,5 +1,9 @@
! 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
@@ -11,3 +15,5 @@ attributes(global) subroutine g2(a, m)
!ERROR: not yet implemented: CUDA dynamic parallelism
call g1<<<1, 1>>>(a, m)
end subroutine
+
+end module
More information about the flang-commits
mailing list