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

[via flang-commits]

Author: Valentin Clement (バレンタイン クレメン)
Date: 2025-01-06T14:30:08-08:00
New Revision: 3874c64418d2a7e36eab9af9253d905b48b36078

URL: https://github.com/llvm/llvm-project/commit/3874c64418d2a7e36eab9af9253d905b48b36078
DIFF: https://github.com/llvm/llvm-project/commit/3874c64418d2a7e36eab9af9253d905b48b36078.diff

LOG: [flang][cuda] Allow constant actual argument for device dummy (#121845)

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

Added: 
    

Modified: 
    flang/lib/Common/Fortran.cpp
    flang/test/Semantics/cuf10.cuf

Removed: 
    


################################################################################
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