Skip to content

Commit 3874c64

Browse files
authored
[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.
1 parent d0c00cf commit 3874c64

File tree

2 files changed

+9
-1
lines changed

2 files changed

+9
-1
lines changed

flang/lib/Common/Fortran.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -136,7 +136,8 @@ bool AreCompatibleCUDADataAttrs(std::optional<CUDADataAttr> x,
136136
if (*x == CUDADataAttr::Device) {
137137
if ((y &&
138138
(*y == CUDADataAttr::Managed || *y == CUDADataAttr::Unified ||
139-
*y == CUDADataAttr::Shared)) ||
139+
*y == CUDADataAttr::Shared ||
140+
*y == CUDADataAttr::Constant)) ||
140141
(!y && (isCudaUnified || isCudaManaged))) {
141142
if (y && *y == CUDADataAttr::Shared && warning) {
142143
*warning = "SHARED attribute ignored"s;

flang/test/Semantics/cuf10.cuf

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22
module m
33
real, device :: a(4,8)
44
real, managed, allocatable :: b(:,:)
5+
integer, constant :: x = 1
56
contains
67
attributes(global) subroutine kernel(a,b,c,n,m)
78
integer, value :: n
@@ -23,4 +24,10 @@ module m
2324
call devsub(c,4) ! not checked in OpenACC construct
2425
end do
2526
end
27+
attributes(global) subroutine sub1(x)
28+
integer :: x
29+
end
30+
subroutine sub2()
31+
call sub1<<<1,1>>>(x) ! actual constant to device dummy
32+
end
2633
end

0 commit comments

Comments
 (0)