From 3874c64418d2a7e36eab9af9253d905b48b36078 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?= Date: Mon, 6 Jan 2025 14:30:08 -0800 Subject: [PATCH] [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. --- 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 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