[flang-commits] [flang] [flang][cuda] Diagnose host reads of device data (PR #228287)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Fri Oct 2 13:55:24 PDT 2026
https://github.com/clementval updated https://github.com/llvm/llvm-project/pull/228287
>From 8171fd2b0da9d3e9b1b6136b48d2872740f31a43 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 1 Oct 2026 15:34:59 -0700
Subject: [PATCH 1/6] [flang][cuda] Diagnose host reads of device data
---
flang/lib/Semantics/check-cuda.cpp | 116 ++++++++++++++++++
flang/lib/Semantics/check-cuda.h | 11 ++
.../CUDA/cuf-device-data-host-read.cuf | 112 +++++++++++++++++
3 files changed, 239 insertions(+)
create mode 100644 flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index e152cca10c50f..bab1d4819fb0a 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -207,6 +207,81 @@ struct FindHostArray
}
};
+// Intrinsics that only use the address or the descriptor of their arguments.
+static const llvm::StringSet<> hostAddressIntrinsics_ = {
+ "__builtin_c_devloc", "__builtin_c_loc", "c_sizeof", "loc", "sizeof"};
+
+// Traverses an expression evaluated by host code in search of device data
+// whose value would have to be read from the host. Device data that is only
+// designated (actual argument to a procedure, argument of an inquiry
+// intrinsic) is not read; the subscripts of its designator still are.
+struct FindDeviceDataReadOnHost
+ : public evaluate::AnyTraverse<FindDeviceDataReadOnHost, const Symbol *> {
+ using Result = const Symbol *;
+ using Base = evaluate::AnyTraverse<FindDeviceDataReadOnHost, Result>;
+ explicit FindDeviceDataReadOnHost(
+ SemanticsContext &c, bool onlyDesignated = false)
+ : Base(*this), context_{c}, onlyDesignated_{onlyDesignated} {}
+ using Base::operator();
+ Result operator()(const Symbol &symbol) const {
+ if (!onlyDesignated_ &&
+ evaluate::IsCUDADeviceOnlySymbol(GetAssociationRoot(symbol))) {
+ return &symbol;
+ }
+ return nullptr;
+ }
+ Result operator()(const evaluate::Component &x) const {
+ const Symbol &component{x.GetLastSymbol()};
+ if (evaluate::HasCUDADataAttr(component)) {
+ return (*this)(component);
+ }
+ return (*this)(x.base());
+ }
+ Result operator()(const evaluate::ArrayRef &x) const {
+ if (Result result{(*this)(x.base())}) {
+ return result;
+ }
+ return FindDeviceDataReadOnHost{context_}(x.subscript());
+ }
+ Result operator()(const evaluate::Substring &x) const {
+ if (Result result{(*this)(x.parent())}) {
+ return result;
+ }
+ FindDeviceDataReadOnHost readChecker{context_};
+ if (Result result{readChecker(x.lower())}) {
+ return result;
+ }
+ return readChecker(x.upper());
+ }
+ Result operator()(const evaluate::DescriptorInquiry &) const {
+ return nullptr;
+ }
+ Result operator()(const evaluate::TypeParamInquiry &) const {
+ return nullptr;
+ }
+ Result operator()(const evaluate::ProcedureRef &x) const {
+ if (const auto *intrinsic{x.proc().GetSpecificIntrinsic()}) {
+ if (context_.intrinsics().GetIntrinsicClass(intrinsic->name) !=
+ evaluate::IntrinsicClass::inquiryFunction &&
+ !hostAddressIntrinsics_.contains(intrinsic->name)) {
+ return (*this)(x.arguments());
+ }
+ }
+ for (const auto &arg : x.arguments()) {
+ if (const auto *expr{arg ? arg->UnwrapExpr() : nullptr}) {
+ if (Result result{FindDeviceDataReadOnHost{context_,
+ onlyDesignated_ || evaluate::IsVariable(*expr)}(*expr)}) {
+ return result;
+ }
+ }
+ }
+ return nullptr;
+ }
+
+ SemanticsContext &context_;
+ bool onlyDesignated_{false};
+};
+
template <typename A>
static MaybeMsg CheckUnwrappedExpr(
SemanticsContext &context, const A &x, bool allowHostCallees = false) {
@@ -857,6 +932,47 @@ void CUDAChecker::Enter(const parser::AssignmentStmt &x) {
}
}
+template <typename A> void CUDAChecker::EnterHostScalarExpr(const A &x) {
+ if (hostScalarExprDepth_++ > 0) {
+ return; // Checked with the enclosing expression.
+ }
+ const auto &expr{DEREF(parser::Unwrap<parser::Expr>(x))};
+ const Scope &progUnit{
+ GetProgramUnitContaining(context_.FindScope(expr.source))};
+ if (IsCUDADeviceContext(&progUnit) || deviceConstructDepth_ > 0) {
+ return;
+ }
+ if (const auto *typedExpr{GetExpr(context_, expr)}) {
+ if (const Symbol *deviceData{
+ FindDeviceDataReadOnHost{context_}(*typedExpr)}) {
+ context_.Say(expr.source,
+ "Device data '%s' may not be referenced in host code outside of a data transfer or an actual argument"_err_en_US,
+ deviceData->name());
+ }
+ }
+}
+
+void CUDAChecker::Enter(const parser::Scalar<parser::Expr> &x) {
+ EnterHostScalarExpr(x);
+}
+void CUDAChecker::Leave(const parser::Scalar<parser::Expr> &) {
+ --hostScalarExprDepth_;
+}
+void CUDAChecker::Enter(const parser::ScalarExpr &x) { EnterHostScalarExpr(x); }
+void CUDAChecker::Leave(const parser::ScalarExpr &) { --hostScalarExprDepth_; }
+void CUDAChecker::Enter(const parser::ScalarIntExpr &x) {
+ EnterHostScalarExpr(x);
+}
+void CUDAChecker::Leave(const parser::ScalarIntExpr &) {
+ --hostScalarExprDepth_;
+}
+void CUDAChecker::Enter(const parser::ScalarLogicalExpr &x) {
+ EnterHostScalarExpr(x);
+}
+void CUDAChecker::Leave(const parser::ScalarLogicalExpr &) {
+ --hostScalarExprDepth_;
+}
+
void CUDAChecker::Enter(const parser::PrintStmt &x) {
CHECK(context_.location());
const Scope &scope{context_.FindScope(*context_.location())};
diff --git a/flang/lib/Semantics/check-cuda.h b/flang/lib/Semantics/check-cuda.h
index ef5e57ab41b81..0623e2c09aeec 100644
--- a/flang/lib/Semantics/check-cuda.h
+++ b/flang/lib/Semantics/check-cuda.h
@@ -50,10 +50,21 @@ class CUDAChecker : public virtual BaseChecker {
void Enter(const parser::DoConstruct &);
void Leave(const parser::DoConstruct &);
void Enter(const parser::PrintStmt &);
+ void Enter(const parser::Scalar<parser::Expr> &);
+ void Leave(const parser::Scalar<parser::Expr> &);
+ void Enter(const parser::ScalarExpr &);
+ void Leave(const parser::ScalarExpr &);
+ void Enter(const parser::ScalarIntExpr &);
+ void Leave(const parser::ScalarIntExpr &);
+ void Enter(const parser::ScalarLogicalExpr &);
+ void Leave(const parser::ScalarLogicalExpr &);
private:
+ template <typename A> void EnterHostScalarExpr(const A &);
+
SemanticsContext &context_;
int deviceConstructDepth_{0};
+ int hostScalarExprDepth_{0};
};
bool CanonicalizeCUDA(parser::Program &);
diff --git a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
new file mode 100644
index 0000000000000..efa6db17f14e8
--- /dev/null
+++ b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
@@ -0,0 +1,112 @@
+! RUN: %python %S/../test_errors.py %s %flang_fc1 -fopenacc
+
+! Device data may not be read by host code outside of a data transfer
+! assignment or an actual argument.
+
+module m
+ real, device :: md(10)
+ integer, device :: mn
+ integer, constant :: mc = 10
+ real, managed :: mm(10)
+ type t
+ real, device, allocatable :: c(:)
+ real :: h(10)
+ end type
+contains
+ attributes(global) subroutine k(x)
+ real :: x(:)
+ end subroutine
+ real function f(x)
+ real, device :: x(*)
+ f = 0.0
+ end function
+end module
+
+subroutine dummy_expl_ifcond(h, a)
+ implicit none
+ double precision :: h
+ double precision :: a(10)
+ integer, device :: i
+ attributes(device) :: a
+ !ERROR: Device data 'a' may not be referenced in host code outside of a data transfer or an actual argument
+ if (a(3) > 2.5d0) h = 1.0d0
+ !ERROR: Device data 'i' may not be referenced in host code outside of a data transfer or an actual argument
+ if (i > 0) h = 2.0d0
+end subroutine
+
+subroutine host(ad)
+ use m
+ implicit none
+ real, device, allocatable :: ad(:)
+ real :: h(10), x
+ integer :: i
+ type(t) :: tt
+
+ !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
+ if (mn > 0) x = 1.0
+ if (x > 0.0) then
+ !ERROR: Device data 'md' may not be referenced in host code outside of a data transfer or an actual argument
+ else if (md(1) > 0.0) then
+ end if
+ !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
+ do i = 1, mn
+ end do
+ !ERROR: Device data 'md' may not be referenced in host code outside of a data transfer or an actual argument
+ do while (md(1) > 0.0)
+ end do
+ !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
+ select case (mn)
+ case (1)
+ !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
+ stop mn
+ end select
+ !ERROR: Device data 'md' may not be referenced in host code outside of a data transfer or an actual argument
+ if (md(1) + 1.0 > 0.0) x = 2.0
+ !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
+ if (h(mn) > 0.0) x = 3.0
+ !ERROR: Device data 'c' may not be referenced in host code outside of a data transfer or an actual argument
+ if (tt%c(1) > 0.0) x = 4.0
+ !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
+ if (f(md(mn)) > 0.0) x = 5.0
+
+ if (allocated(ad)) x = 6.0
+ if (size(ad) > 3) x = 7.0
+ if (lbound(ad, 1) == 1) x = 8.0
+ if (allocated(tt%c)) x = 9.0
+ if (f(md) > 0.0) x = 10.0
+ if (f(md(2)) > 0.0) x = 11.0
+ if (mm(1) > 0.0) x = 12.0
+ if (tt%h(1) > 0.0) x = 13.0
+ !ERROR: Device data 'mc' may not be referenced in host code outside of a data transfer or an actual argument
+ if (mc > 0) x = 14.0
+ mc = 10
+ x = md(1)
+ h = md
+ md = h
+ call k<<<1, 1>>>(md)
+ !$cuf kernel do <<<*,*>>>
+ do i = 1, 10
+ if (md(i) > 0.0) md(i) = 0.0
+ end do
+
+ !$acc parallel
+ if (mn > 0) md(1) = 0.0
+ !$acc end parallel
+ !$acc serial
+ if (md(1) > 0.0) md(1) = 0.0
+ !$acc end serial
+ !$acc kernels
+ do i = 1, mn
+ if (md(i) > 0.0) md(i) = 0.0
+ end do
+ !$acc end kernels
+ !$acc parallel loop
+ do i = 1, mn
+ if (md(i) > 0.0) md(i) = 0.0
+ end do
+end subroutine
+
+attributes(global) subroutine kernel(a)
+ real, device :: a(10)
+ if (a(1) > 0.0) a(1) = 1.0
+end subroutine
>From e511ad78207965438ca4a9ec578e33db7c830b8f Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 1 Oct 2026 15:58:09 -0700
Subject: [PATCH 2/6] Add more use cases
---
flang/lib/Semantics/check-cuda.cpp | 25 ++++++++++++++++---
.../CUDA/cuf-device-data-host-read.cuf | 14 +++++++++++
2 files changed, 35 insertions(+), 4 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index bab1d4819fb0a..d5e7da0424c98 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -211,6 +211,12 @@ struct FindHostArray
static const llvm::StringSet<> hostAddressIntrinsics_ = {
"__builtin_c_devloc", "__builtin_c_loc", "c_sizeof", "loc", "sizeof"};
+// Inquiry intrinsics whose arguments are all inquired objects. Other inquiry
+// intrinsics only inquire about their first argument and read the values of
+// the others (DIM=, KIND=).
+static const llvm::StringSet<> allArgsInquiryIntrinsics_ = {
+ "associated", "extends_type_of", "same_type_as"};
+
// Traverses an expression evaluated by host code in search of device data
// whose value would have to be read from the host. Device data that is only
// designated (actual argument to a procedure, argument of an inquiry
@@ -233,7 +239,12 @@ struct FindDeviceDataReadOnHost
Result operator()(const evaluate::Component &x) const {
const Symbol &component{x.GetLastSymbol()};
if (evaluate::HasCUDADataAttr(component)) {
- return (*this)(component);
+ if (Result result{(*this)(component)}) {
+ return result;
+ }
+ // The attribute of the component hides the one of the base.
+ return FindDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
+ x.base());
}
return (*this)(x.base());
}
@@ -260,17 +271,23 @@ struct FindDeviceDataReadOnHost
return nullptr;
}
Result operator()(const evaluate::ProcedureRef &x) const {
+ bool onlyFirstArgDesignated{false};
if (const auto *intrinsic{x.proc().GetSpecificIntrinsic()}) {
if (context_.intrinsics().GetIntrinsicClass(intrinsic->name) !=
evaluate::IntrinsicClass::inquiryFunction &&
!hostAddressIntrinsics_.contains(intrinsic->name)) {
return (*this)(x.arguments());
}
+ onlyFirstArgDesignated =
+ !allArgsInquiryIntrinsics_.contains(intrinsic->name);
}
- for (const auto &arg : x.arguments()) {
+ for (std::size_t j{0}; j < x.arguments().size(); ++j) {
+ const auto &arg{x.arguments()[j]};
if (const auto *expr{arg ? arg->UnwrapExpr() : nullptr}) {
- if (Result result{FindDeviceDataReadOnHost{context_,
- onlyDesignated_ || evaluate::IsVariable(*expr)}(*expr)}) {
+ bool designated{evaluate::IsVariable(*expr) &&
+ (j == 0 || !onlyFirstArgDesignated)};
+ if (Result result{FindDeviceDataReadOnHost{
+ context_, onlyDesignated_ || designated}(*expr)}) {
return result;
}
}
diff --git a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
index efa6db17f14e8..4c89571706ad7 100644
--- a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
+++ b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
@@ -38,9 +38,23 @@ subroutine host(ad)
use m
implicit none
real, device, allocatable :: ad(:)
+ real, device, allocatable :: ad2(:,:)
+ integer, device :: nd
real :: h(10), x
integer :: i
type(t) :: tt
+ type(t) :: tta(10)
+
+ !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
+ if (allocated(tta(nd)%c)) continue
+ !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
+ if (f(tta(nd)%c) > 0.0) continue
+ if (allocated(tta(2)%c)) continue
+ !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
+ if (size(ad2, dim=nd) > 0) continue
+ !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
+ if (lbound(ad2, nd) > 0) continue
+ if (size(ad2, dim=2) > 0) continue
!ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
if (mn > 0) x = 1.0
>From 31bbbe4cb1de39cc27b1bf1fdf5f6191efba9c64 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 1 Oct 2026 16:02:44 -0700
Subject: [PATCH 3/6] format
---
flang/lib/Semantics/check-cuda.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index d5e7da0424c98..25e47091f4072 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -284,8 +284,8 @@ struct FindDeviceDataReadOnHost
for (std::size_t j{0}; j < x.arguments().size(); ++j) {
const auto &arg{x.arguments()[j]};
if (const auto *expr{arg ? arg->UnwrapExpr() : nullptr}) {
- bool designated{evaluate::IsVariable(*expr) &&
- (j == 0 || !onlyFirstArgDesignated)};
+ bool designated{
+ evaluate::IsVariable(*expr) && (j == 0 || !onlyFirstArgDesignated)};
if (Result result{FindDeviceDataReadOnHost{
context_, onlyDesignated_ || designated}(*expr)}) {
return result;
>From 9774d8583a9ae7ec8f2a009d43d1212877a11072 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 1 Oct 2026 17:41:15 -0700
Subject: [PATCH 4/6] Add more use case
---
flang/lib/Semantics/check-cuda.cpp | 6 ++++--
flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf | 6 ++++++
2 files changed, 10 insertions(+), 2 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 25e47091f4072..0c9c169fd0b6c 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -264,8 +264,10 @@ struct FindDeviceDataReadOnHost
}
return readChecker(x.upper());
}
- Result operator()(const evaluate::DescriptorInquiry &) const {
- return nullptr;
+ Result operator()(const evaluate::DescriptorInquiry &x) const {
+ // Accessing descriptor metadata is allowed, but selecting the descriptor
+ // may require reading device data in subscripts.
+ return FindDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(x.base());
}
Result operator()(const evaluate::TypeParamInquiry &) const {
return nullptr;
diff --git a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
index 4c89571706ad7..869bf8ff16cd0 100644
--- a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
+++ b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
@@ -93,6 +93,9 @@ subroutine host(ad)
if (tt%h(1) > 0.0) x = 13.0
!ERROR: Device data 'mc' may not be referenced in host code outside of a data transfer or an actual argument
if (mc > 0) x = 14.0
+ !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
+ if (size(tta(nd)%c, dim=1) > 0) x = 15.0
+ if (size(tta(2)%c, dim=1) > 0) x = 14.0
mc = 10
x = md(1)
h = md
@@ -118,6 +121,9 @@ subroutine host(ad)
do i = 1, mn
if (md(i) > 0.0) md(i) = 0.0
end do
+
+
+
end subroutine
attributes(global) subroutine kernel(a)
>From 58c9e5d9fb0926c8bf77f40a1afbc924d0358e43 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 1 Oct 2026 17:41:36 -0700
Subject: [PATCH 5/6] format
---
flang/lib/Semantics/check-cuda.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 0c9c169fd0b6c..88ebed35ae89b 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -267,7 +267,8 @@ struct FindDeviceDataReadOnHost
Result operator()(const evaluate::DescriptorInquiry &x) const {
// Accessing descriptor metadata is allowed, but selecting the descriptor
// may require reading device data in subscripts.
- return FindDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(x.base());
+ return FindDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
+ x.base());
}
Result operator()(const evaluate::TypeParamInquiry &) const {
return nullptr;
>From ed0f8417951d9fe09a6babd42ca9f932975154f9 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Fri, 2 Oct 2026 13:54:33 -0700
Subject: [PATCH 6/6] Relax some usage
---
flang/lib/Semantics/check-cuda.cpp | 104 ++++++++++--------
.../CUDA/cuf-device-data-host-read.cuf | 86 ++++++++++-----
2 files changed, 120 insertions(+), 70 deletions(-)
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index 88ebed35ae89b..42b8af80a29cb 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -217,62 +217,57 @@ static const llvm::StringSet<> hostAddressIntrinsics_ = {
static const llvm::StringSet<> allArgsInquiryIntrinsics_ = {
"associated", "extends_type_of", "same_type_as"};
-// Traverses an expression evaluated by host code in search of device data
-// whose value would have to be read from the host. Device data that is only
-// designated (actual argument to a procedure, argument of an inquiry
-// intrinsic) is not read; the subscripts of its designator still are.
-struct FindDeviceDataReadOnHost
- : public evaluate::AnyTraverse<FindDeviceDataReadOnHost, const Symbol *> {
- using Result = const Symbol *;
- using Base = evaluate::AnyTraverse<FindDeviceDataReadOnHost, Result>;
- explicit FindDeviceDataReadOnHost(
+// Collects, in order of appearance, the device data whose value host code
+// reads when it evaluates an expression. Device data that is only designated
+// (actual argument to a procedure, argument of an inquiry intrinsic) is not
+// read; the subscripts of its designator still are. A component with a CUDA
+// data attribute is collected instead of its base.
+struct CollectDeviceDataReadOnHost
+ : public evaluate::Traverse<CollectDeviceDataReadOnHost, SymbolVector> {
+ using Result = SymbolVector;
+ using Base = evaluate::Traverse<CollectDeviceDataReadOnHost, Result>;
+ explicit CollectDeviceDataReadOnHost(
SemanticsContext &c, bool onlyDesignated = false)
: Base(*this), context_{c}, onlyDesignated_{onlyDesignated} {}
using Base::operator();
+ static Result Default() { return {}; }
+ static Result Combine(Result &&x, Result &&y) {
+ x.insert(x.end(), y.begin(), y.end());
+ return std::move(x);
+ }
Result operator()(const Symbol &symbol) const {
- if (!onlyDesignated_ &&
- evaluate::IsCUDADeviceOnlySymbol(GetAssociationRoot(symbol))) {
- return &symbol;
+ if (onlyDesignated_ ||
+ !evaluate::IsCUDADeviceOnlySymbol(GetAssociationRoot(symbol))) {
+ return {};
}
- return nullptr;
+ return {symbol};
}
Result operator()(const evaluate::Component &x) const {
const Symbol &component{x.GetLastSymbol()};
if (evaluate::HasCUDADataAttr(component)) {
- if (Result result{(*this)(component)}) {
- return result;
- }
// The attribute of the component hides the one of the base.
- return FindDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
- x.base());
+ return Combine((*this)(component),
+ CollectDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
+ x.base()));
}
return (*this)(x.base());
}
Result operator()(const evaluate::ArrayRef &x) const {
- if (Result result{(*this)(x.base())}) {
- return result;
- }
- return FindDeviceDataReadOnHost{context_}(x.subscript());
+ return Combine((*this)(x.base()),
+ CollectDeviceDataReadOnHost{context_}(x.subscript()));
}
Result operator()(const evaluate::Substring &x) const {
- if (Result result{(*this)(x.parent())}) {
- return result;
- }
- FindDeviceDataReadOnHost readChecker{context_};
- if (Result result{readChecker(x.lower())}) {
- return result;
- }
- return readChecker(x.upper());
+ CollectDeviceDataReadOnHost readCollector{context_};
+ return Combine((*this)(x.parent()),
+ Combine(readCollector(x.lower()), readCollector(x.upper())));
}
Result operator()(const evaluate::DescriptorInquiry &x) const {
// Accessing descriptor metadata is allowed, but selecting the descriptor
// may require reading device data in subscripts.
- return FindDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
+ return CollectDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
x.base());
}
- Result operator()(const evaluate::TypeParamInquiry &) const {
- return nullptr;
- }
+ Result operator()(const evaluate::TypeParamInquiry &) const { return {}; }
Result operator()(const evaluate::ProcedureRef &x) const {
bool onlyFirstArgDesignated{false};
if (const auto *intrinsic{x.proc().GetSpecificIntrinsic()}) {
@@ -284,24 +279,31 @@ struct FindDeviceDataReadOnHost
onlyFirstArgDesignated =
!allArgsInquiryIntrinsics_.contains(intrinsic->name);
}
+ Result result;
for (std::size_t j{0}; j < x.arguments().size(); ++j) {
const auto &arg{x.arguments()[j]};
if (const auto *expr{arg ? arg->UnwrapExpr() : nullptr}) {
bool designated{
evaluate::IsVariable(*expr) && (j == 0 || !onlyFirstArgDesignated)};
- if (Result result{FindDeviceDataReadOnHost{
- context_, onlyDesignated_ || designated}(*expr)}) {
- return result;
- }
+ result = Combine(std::move(result),
+ CollectDeviceDataReadOnHost{
+ context_, onlyDesignated_ || designated}(*expr));
}
}
- return nullptr;
+ return result;
}
SemanticsContext &context_;
bool onlyDesignated_{false};
};
+// A scalar variable other than a component, an allocatable or a pointer.
+static bool IsPlainScalar(const Symbol &symbol) {
+ const Symbol &ultimate{symbol.GetUltimate()};
+ return ultimate.has<ObjectEntityDetails>() && ultimate.Rank() == 0 &&
+ !ultimate.owner().IsDerivedType() && !IsAllocatableOrPointer(ultimate);
+}
+
template <typename A>
static MaybeMsg CheckUnwrappedExpr(
SemanticsContext &context, const A &x, bool allowHostCallees = false) {
@@ -963,11 +965,25 @@ template <typename A> void CUDAChecker::EnterHostScalarExpr(const A &x) {
return;
}
if (const auto *typedExpr{GetExpr(context_, expr)}) {
- if (const Symbol *deviceData{
- FindDeviceDataReadOnHost{context_}(*typedExpr)}) {
- context_.Say(expr.source,
- "Device data '%s' may not be referenced in host code outside of a data transfer or an actual argument"_err_en_US,
- deviceData->name());
+ const Symbol *deviceScalar{nullptr};
+ for (const Symbol &deviceData :
+ CollectDeviceDataReadOnHost{context_}(*typedExpr)) {
+ if (!IsPlainScalar(deviceData)) {
+ context_.Say(expr.source,
+ "Device data '%s' may not be referenced in host code outside of a data transfer or an actual argument"_err_en_US,
+ deviceData.name());
+ return;
+ }
+ // Host code reads a module scalar from its host copy.
+ if (!deviceScalar &&
+ deviceData.GetUltimate().owner().kind() != Scope::Kind::Module) {
+ deviceScalar = &deviceData;
+ }
+ }
+ if (deviceScalar) {
+ context_.Warn(common::UsageWarning::CUDAUsage, expr.source,
+ "Device data '%s' is read in host code outside of a data transfer or an actual argument"_warn_en_US,
+ deviceScalar->name());
}
}
}
diff --git a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
index 869bf8ff16cd0..564fea59fc1bc 100644
--- a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
+++ b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
@@ -1,12 +1,14 @@
! RUN: %python %S/../test_errors.py %s %flang_fc1 -fopenacc
! Device data may not be read by host code outside of a data transfer
-! assignment or an actual argument.
+! assignment or an actual argument. Module scalars are read from their host
+! copy, and other device scalars are only warned about.
module m
real, device :: md(10)
integer, device :: mn
integer, constant :: mc = 10
+ integer, constant :: mca(3)
real, managed :: mm(10)
type t
real, device, allocatable :: c(:)
@@ -30,7 +32,7 @@ subroutine dummy_expl_ifcond(h, a)
attributes(device) :: a
!ERROR: Device data 'a' may not be referenced in host code outside of a data transfer or an actual argument
if (a(3) > 2.5d0) h = 1.0d0
- !ERROR: Device data 'i' may not be referenced in host code outside of a data transfer or an actual argument
+ !WARNING: Device data 'i' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
if (i > 0) h = 2.0d0
end subroutine
@@ -39,49 +41,80 @@ subroutine host(ad)
implicit none
real, device, allocatable :: ad(:)
real, device, allocatable :: ad2(:,:)
- integer, device :: nd
+ integer, device :: nd, nda(2)
+ integer, device, allocatable :: ndl
real :: h(10), x
integer :: i
type(t) :: tt
type(t) :: tta(10)
- !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
- if (allocated(tta(nd)%c)) continue
- !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
- if (f(tta(nd)%c) > 0.0) continue
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (allocated(tta(nda(1))%c)) continue
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (f(tta(nda(1))%c) > 0.0) continue
if (allocated(tta(2)%c)) continue
- !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
- if (size(ad2, dim=nd) > 0) continue
- !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
- if (lbound(ad2, nd) > 0) continue
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (size(ad2, dim=nda(1)) > 0) continue
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (lbound(ad2, nda(1)) > 0) continue
if (size(ad2, dim=2) > 0) continue
- !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
- if (mn > 0) x = 1.0
+ !ERROR: Device data 'ndl' may not be referenced in host code outside of a data transfer or an actual argument
+ if (ndl > 0) x = 1.0
if (x > 0.0) then
!ERROR: Device data 'md' may not be referenced in host code outside of a data transfer or an actual argument
else if (md(1) > 0.0) then
end if
- !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
- do i = 1, mn
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ do i = 1, nda(2)
end do
!ERROR: Device data 'md' may not be referenced in host code outside of a data transfer or an actual argument
do while (md(1) > 0.0)
end do
- !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
- select case (mn)
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ select case (nda(1))
case (1)
- !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
- stop mn
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ stop nda(1)
end select
!ERROR: Device data 'md' may not be referenced in host code outside of a data transfer or an actual argument
if (md(1) + 1.0 > 0.0) x = 2.0
- !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
- if (h(mn) > 0.0) x = 3.0
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (h(nda(1)) > 0.0) x = 3.0
!ERROR: Device data 'c' may not be referenced in host code outside of a data transfer or an actual argument
if (tt%c(1) > 0.0) x = 4.0
- !ERROR: Device data 'mn' may not be referenced in host code outside of a data transfer or an actual argument
- if (f(md(mn)) > 0.0) x = 5.0
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (f(md(nda(1))) > 0.0) x = 5.0
+
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ if (nd > 0) x = 1.0
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ if (h(nd) > 0.0) x = 3.0
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ if (f(md(nd)) > 0.0) x = 5.0
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ if (allocated(tta(nd)%c)) continue
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ if (size(ad2, dim=nd) > 0) continue
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ do i = 1, nd
+ end do
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ select case (nd)
+ case (1)
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ stop nd
+ end select
+ !WARNING: Device data 'nd' is read in host code outside of a data transfer or an actual argument [-Wcuda-usage]
+ allocate(ad(nd))
+
+ ! Module scalars are read from their host copy.
+ if (mn > 0) x = 1.0
+ if (h(mn) > 0.0) x = 3.0
+ if (mn + mc > 0) x = 1.0
+ do i = 1, mn
+ end do
+ allocate(ad2(mn, mc))
if (allocated(ad)) x = 6.0
if (size(ad) > 3) x = 7.0
@@ -91,10 +124,11 @@ subroutine host(ad)
if (f(md(2)) > 0.0) x = 11.0
if (mm(1) > 0.0) x = 12.0
if (tt%h(1) > 0.0) x = 13.0
- !ERROR: Device data 'mc' may not be referenced in host code outside of a data transfer or an actual argument
if (mc > 0) x = 14.0
- !ERROR: Device data 'nd' may not be referenced in host code outside of a data transfer or an actual argument
- if (size(tta(nd)%c, dim=1) > 0) x = 15.0
+ !ERROR: Device data 'mca' may not be referenced in host code outside of a data transfer or an actual argument
+ if (mca(1) > 0) x = 14.0
+ !ERROR: Device data 'nda' may not be referenced in host code outside of a data transfer or an actual argument
+ if (size(tta(nda(1))%c, dim=1) > 0) x = 15.0
if (size(tta(2)%c, dim=1) > 0) x = 14.0
mc = 10
x = md(1)
More information about the flang-commits
mailing list