[flang-commits] [flang] 79f4151 - Revert "[flang][cuda] Diagnose host reads of device data" (#228286)
via flang-commits
flang-commits at lists.llvm.org
Thu Oct 1 17:53:49 PDT 2026
Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-10-02T00:53:40Z
New Revision: 79f415178fb162f390e8c4628001fb4b8ce3998b
URL: https://github.com/llvm/llvm-project/commit/79f415178fb162f390e8c4628001fb4b8ce3998b
DIFF: https://github.com/llvm/llvm-project/commit/79f415178fb162f390e8c4628001fb4b8ce3998b.diff
LOG: Revert "[flang][cuda] Diagnose host reads of device data" (#228286)
Reverts llvm/llvm-project#228271
auto-merging was on by mistake
Added:
Modified:
flang/lib/Semantics/check-cuda.cpp
flang/lib/Semantics/check-cuda.h
Removed:
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 25e47091f4072..e152cca10c50f 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -207,98 +207,6 @@ 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"};
-
-// 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)) {
- 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());
- }
- 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 {
- 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 (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;
- }
- }
- }
- return nullptr;
- }
-
- SemanticsContext &context_;
- bool onlyDesignated_{false};
-};
-
template <typename A>
static MaybeMsg CheckUnwrappedExpr(
SemanticsContext &context, const A &x, bool allowHostCallees = false) {
@@ -949,47 +857,6 @@ 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 0623e2c09aeec..ef5e57ab41b81 100644
--- a/flang/lib/Semantics/check-cuda.h
+++ b/flang/lib/Semantics/check-cuda.h
@@ -50,21 +50,10 @@ 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
deleted file mode 100644
index 4c89571706ad7..0000000000000
--- a/flang/test/Semantics/CUDA/cuf-device-data-host-read.cuf
+++ /dev/null
@@ -1,126 +0,0 @@
-! 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, 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
- 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
More information about the flang-commits
mailing list