[flang][cuda] Relax check on chevron syntax for bind(c) callee

[clementval]

No description provided.

[llvmbot]

@llvm/pr-subscribers-flang-semantics

Author: Valentin Clement (バレンタイン クレメン) (clementval)

Changes

---

Full diff: https://github.com/llvm/llvm-project/pull/190861.diff

2 Files Affected:

• (modified) flang/lib/Semantics/expression.cpp (+3-1)
• (modified) flang/test/Semantics/cuf04.cuf (+3)

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