[flang-commits] [flang] [flang][cuda] Allow constant actual argument for device dummy (PR #121845)

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

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

The reference compiler allows this use case. Note that writing to this variable would result in CUDA error. 

>From 4e49b092a13e7786f394a5314076ff3a157a6c30 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Mon, 6 Jan 2025 13:13:52 -0800
Subject: [PATCH] [flang][cuda] Allow constant actual argument for device dummy

---
 flang/lib/Common/Fortran.cpp   | 3 ++-
 flang/test/Semantics/cuf10.cuf | 7 +++++++
 2 files changed, 9 insertions(+), 1 deletion(-)

diff --git a/flang/lib/Common/Fortran.cpp b/flang/lib/Common/Fortran.cpp
index 367a51f884e8a5..eec83419f9eb2b 100644
--- a/flang/lib/Common/Fortran.cpp
+++ b/flang/lib/Common/Fortran.cpp
@@ -136,7 +136,8 @@ bool AreCompatibleCUDADataAttrs(std::optional<CUDADataAttr> x,
       if (*x == CUDADataAttr::Device) {
         if ((y &&
                 (*y == CUDADataAttr::Managed || *y == CUDADataAttr::Unified ||
-                    *y == CUDADataAttr::Shared)) ||
+                    *y == CUDADataAttr::Shared ||
+                    *y == CUDADataAttr::Constant)) ||
             (!y && (isCudaUnified || isCudaManaged))) {
           if (y && *y == CUDADataAttr::Shared && warning) {
             *warning = "SHARED attribute ignored"s;
diff --git a/flang/test/Semantics/cuf10.cuf b/flang/test/Semantics/cuf10.cuf
index 047503b3cca4ea..24b596b1fa55db 100644
--- a/flang/test/Semantics/cuf10.cuf
+++ b/flang/test/Semantics/cuf10.cuf
@@ -2,6 +2,7 @@
 module m
   real, device :: a(4,8)
   real, managed, allocatable :: b(:,:)
+  integer, constant :: x = 1
  contains
   attributes(global) subroutine kernel(a,b,c,n,m)
     integer, value :: n
@@ -23,4 +24,10 @@ module m
       call devsub(c,4) ! not checked in OpenACC construct
     end do
   end
+  attributes(global) subroutine sub1(x)
+    integer :: x
+  end
+  subroutine sub2()
+    call sub1<<<1,1>>>(x) ! actual constant to device dummy
+  end
 end