[flang-commits] [flang] 3d66265 - [flang][cuda] Diagnose host reads of device data (#228287)
via flang-commits
flang-commits at lists.llvm.org
Sun Oct 4 19:24:06 PDT 2026
Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-10-05T02:23:57Z
New Revision: 3d66265c8a8d32d97cc9462401feda08870c8de7
URL: https://github.com/llvm/llvm-project/commit/3d66265c8a8d32d97cc9462401feda08870c8de7
DIFF: https://github.com/llvm/llvm-project/commit/3d66265c8a8d32d97cc9462401feda08870c8de7.diff
LOG: [flang][cuda] Diagnose host reads of device data (#228287)
Host code may only use device data in a data transfer assignment or as
an
actual argument. When device data appeared anywhere else, such as in an
IF
condition, flang accepted it silently and lowered it to a plain host
load of
device memory:
subroutine s(h, a)
double precision :: h, a(10)
attributes(device) :: a
if (a(3) > 2.5d0) h = 1.0d0
end subroutine
The CUDA checker now emits an error when host code reads data with the
DEVICE or CONSTANT attribute in a scalar expression. These expressions
include IF and ELSE IF conditions, DO bounds, DO WHILE, SELECT CASE
selectors, STOP codes and allocate bounds.
Device data that is only designated is not read and is still accepted.
This
covers inquiry intrinsics (ALLOCATED, SIZE, LBOUND, ...), address
intrinsics
(LOC, C_LOC, C_DEVLOC), and variables passed as actual arguments.
Subscripts
of these designators are still checked, so f(d(n_dev)) is an error.
Managed and unified data are host accessible and are not diagnosed.
Device subprograms, CUF kernels and OpenACC compute constructs are
skipped.
Assignments keep their existing data transfer checks.
Added:
flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
Modified:
flang/lib/Semantics/check-cuda.cpp
flang/lib/Semantics/check-cuda.h
Removed:
################################################################################
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index e152cca10c50f..42b8af80a29cb 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -207,6 +207,103 @@ 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"};
+
+// 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"};
+
+// 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 {};
+ }
+ return {symbol};
+ }
+ Result operator()(const evaluate::Component &x) const {
+ const Symbol &component{x.GetLastSymbol()};
+ if (evaluate::HasCUDADataAttr(component)) {
+ // The attribute of the component hides the one of the base.
+ return Combine((*this)(component),
+ CollectDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
+ x.base()));
+ }
+ return (*this)(x.base());
+ }
+ Result operator()(const evaluate::ArrayRef &x) const {
+ return Combine((*this)(x.base()),
+ CollectDeviceDataReadOnHost{context_}(x.subscript()));
+ }
+ Result operator()(const evaluate::Substring &x) const {
+ 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 CollectDeviceDataReadOnHost{context_, /*onlyDesignated=*/true}(
+ x.base());
+ }
+ Result operator()(const evaluate::TypeParamInquiry &) const { return {}; }
+ 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);
+ }
+ 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)};
+ result = Combine(std::move(result),
+ CollectDeviceDataReadOnHost{
+ context_, onlyDesignated_ || designated}(*expr));
+ }
+ }
+ 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) {
@@ -857,6 +954,61 @@ 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)}) {
+ 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());
+ }
+ }
+}
+
+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..564fea59fc1bc
--- /dev/null
+++ b/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
@@ -0,0 +1,166 @@
+! 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 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(:)
+ 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
+ !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
+
+subroutine host(ad)
+ use m
+ implicit none
+ real, device, allocatable :: ad(:)
+ real, device, allocatable :: ad2(:,:)
+ 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 '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 '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 '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 '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 '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 '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 '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 '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
+ 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
+ if (mc > 0) x = 14.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)
+ 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
More information about the flang-commits
mailing list