[flang-commits] [flang] [flang][cuda] Relax check on chevron syntax for bind(c) callee (PR #190861)

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

https://github.com/clementval created https://github.com/llvm/llvm-project/pull/190861

None

>From daa863f65aa5894790189ccc3ab9bbd51b895c33 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Tue, 7 Apr 2026 15:10:18 -0700
Subject: [PATCH] [flang][cuda] Relax check on chevron syntax for bind(c)
 callee

---
 flang/lib/Semantics/expression.cpp | 4 +++-
 flang/test/Semantics/cuf04.cuf     | 3 +++
 2 files changed, 6 insertions(+), 1 deletion(-)

diff --git a/flang/lib/Semantics/expression.cpp b/flang/lib/Semantics/expression.cpp
index 2d7c32fd5fb8f..97e9bd39f0630 100644
--- a/flang/lib/Semantics/expression.cpp
+++ b/flang/lib/Semantics/expression.cpp
@@ -3541,8 +3541,10 @@ void ExpressionAnalyzer::Analyze(const parser::CallStmt &callStmt) {
       ProcedureDesignator *proc{std::get_if<ProcedureDesignator>(&callee->u)};
       CHECK(proc);
       bool isKernel{false};
+      bool isBindC{false};
       if (const Symbol * procSym{proc->GetSymbol()}) {
         const Symbol &ultimate{procSym->GetUltimate()};
+        isBindC = ultimate.attrs().test(semantics::Attr::BIND_C);
         if (const auto *subpDetails{
                 ultimate.detailsIf<semantics::SubprogramDetails>()}) {
           if (auto attrs{subpDetails->cudaSubprogramAttrs()}) {
@@ -3558,7 +3560,7 @@ void ExpressionAnalyzer::Analyze(const parser::CallStmt &callStmt) {
               procSym->name());
         }
       }
-      if (!isKernel && !chevrons->empty()) {
+      if (!isKernel && !isBindC && !chevrons->empty()) {
         Say("Kernel launch parameters in chevrons may not be used unless calling a kernel subroutine"_err_en_US);
       }
       if (CheckCall(callStmt.source, *proc, callee->arguments)) {
diff --git a/flang/test/Semantics/cuf04.cuf b/flang/test/Semantics/cuf04.cuf
index 3c9ee7388ab2d..5350c7869f867 100644
--- a/flang/test/Semantics/cuf04.cuf
+++ b/flang/test/Semantics/cuf04.cuf
@@ -15,10 +15,13 @@ module m
   end subroutine
   subroutine boring
   end subroutine
+  subroutine c_kernel() bind(c,name='c_kernel')
+  end subroutine
   subroutine test
     !ERROR: 'globsubr' is a kernel subroutine and must be called with kernel launch parameters in chevrons
     call globsubr
     !ERROR: Kernel launch parameters in chevrons may not be used unless calling a kernel subroutine
     call boring<<<1,2>>>
+    call c_kernel<<<1,1>>>() ! ok because bind(c) subroutine
   end subroutine
 end module