[flang][cuda] Do not produce data transfer in offloaded do concurrent

Author: clementval

flang/include/flang/Optimizer/Builder/CUFCommon.h

@@ -27,7 +27,7 @@ mlir::gpu::GPUModuleOp getOrCreateGPUModule(mlir::ModuleOp mod,
                                             mlir::SymbolTable &symTab);
 
 bool isCUDADeviceContext(mlir::Operation *op);
-bool isCUDADeviceContext(mlir::Region &);
+bool isCUDADeviceContext(mlir::Region &, bool isStdParEnabled = false);
 bool isRegisteredDeviceGlobal(fir::GlobalOp op);
 bool isRegisteredDeviceAttr(std::optional<cuf::DataAttribute> attr);
 

flang/include/flang/Support/Fortran-features.h

@@ -55,7 +55,7 @@ ENUM_CLASS(LanguageFeature, BackslashEscapes, OldDebugLines,
     SavedLocalInSpecExpr, PrintNamelist, AssumedRankPassedToNonAssumedRank,
     IgnoreIrrelevantAttributes, Unsigned, AmbiguousStructureConstructor,
     ContiguousOkForSeqAssociation, ForwardRefExplicitTypeDummy,
-    InaccessibleDeferredOverride, CudaWarpMatchFunction)
+    InaccessibleDeferredOverride, CudaWarpMatchFunction, StdPar)
 
 // Portability and suspicious usage warnings
 ENUM_CLASS(UsageWarning, Portability, PointerToUndefinable,

flang/lib/Lower/Bridge.cpp

@@ -4886,7 +4886,9 @@ class FirConverter : public Fortran::lower::AbstractConverter {
     mlir::Location loc = getCurrentLocation();
     fir::FirOpBuilder &builder = getFirOpBuilder();
 
-    bool isInDeviceContext = cuf::isCUDADeviceContext(builder.getRegion());
+    bool isInDeviceContext = cuf::isCUDADeviceContext(
+        builder.getRegion(), getFoldingContext().languageFeatures().IsEnabled(
+                                 Fortran::common::LanguageFeature::StdPar));
 
     bool isCUDATransfer =
         IsCUDADataTransfer(assign.lhs, assign.rhs) && !isInDeviceContext;

flang/lib/Optimizer/Builder/CUFCommon.cpp

@@ -43,7 +43,7 @@ bool cuf::isCUDADeviceContext(mlir::Operation *op) {
 // for it.
 // If the insertion point is inside an OpenACC region op, it is considered
 // device context.
-bool cuf::isCUDADeviceContext(mlir::Region &region) {
+bool cuf::isCUDADeviceContext(mlir::Region &region, bool isStdParEnabled) {
   if (region.getParentOfType<cuf::KernelOp>())
     return true;
   if (region.getParentOfType<mlir::acc::ComputeRegionOpInterface>())
@@ -56,6 +56,8 @@ bool cuf::isCUDADeviceContext(mlir::Region &region) {
              cudaProcAttr.getValue() != cuf::ProcAttribute::HostDevice;
     }
   }
+  if (isStdParEnabled && region.getParentOfType<fir::DoConcurrentLoopOp>())
+    return true;
   return false;
 }
 

flang/lib/Semantics/assignment.cpp

@@ -260,6 +260,20 @@ void AssignmentChecker::Enter(const parser::CUFKernelDoConstruct &x) {
 void AssignmentChecker::Leave(const parser::CUFKernelDoConstruct &) {
   --context_.value().deviceConstructDepth_;
 }
+void AssignmentChecker::Enter(const parser::DoConstruct &x) {
+  if (x.IsDoConcurrent() &&
+      context().foldingContext().languageFeatures().IsEnabled(
+          common::LanguageFeature::StdPar)) {
+    ++context_.value().deviceConstructDepth_;
+  }
+}
+void AssignmentChecker::Leave(const parser::DoConstruct &x) {
+  if (x.IsDoConcurrent() &&
+      context().foldingContext().languageFeatures().IsEnabled(
+          common::LanguageFeature::StdPar)) {
+    --context_.value().deviceConstructDepth_;
+  }
+}
 static bool IsOpenACCComputeConstruct(const parser::OpenACCBlockConstruct &x) {
   const auto &beginBlockDirective =
       std::get<Fortran::parser::AccBeginBlockDirective>(x.t);

flang/lib/Semantics/assignment.h

@@ -21,6 +21,7 @@ struct MaskedElsewhereStmt;
 struct PointerAssignmentStmt;
 struct WhereConstructStmt;
 struct WhereStmt;
+struct DoConstruct;
 } // namespace Fortran::parser
 
 namespace Fortran::semantics {
@@ -54,6 +55,8 @@ class AssignmentChecker : public virtual BaseChecker {
   void Leave(const parser::OpenACCCombinedConstruct &);
   void Enter(const parser::OpenACCLoopConstruct &);
   void Leave(const parser::OpenACCLoopConstruct &);
+  void Enter(const parser::DoConstruct &);
+  void Leave(const parser::DoConstruct &);
 
   SemanticsContext &context();
 

flang/lib/Semantics/semantics.cpp

@@ -197,11 +197,11 @@ static void WarnUndefinedFunctionResult(
 
 using StatementSemanticsPass1 = ExprChecker;
 using StatementSemanticsPass2 = SemanticsVisitor<AllocateChecker,
-    ArithmeticIfStmtChecker, AssignmentChecker, CaseChecker, CoarrayChecker,
-    DataChecker, DeallocateChecker, DoForallChecker, IfStmtChecker, IoChecker,
-    MiscChecker, NamelistChecker, NullifyChecker, PurityChecker,
-    ReturnStmtChecker, SelectRankConstructChecker, SelectTypeChecker,
-    StopChecker>;
+    ArithmeticIfStmtChecker, CaseChecker, CoarrayChecker, DataChecker,
+    DeallocateChecker, DoForallChecker, IfStmtChecker, IoChecker, MiscChecker,
+    NamelistChecker, NullifyChecker, PurityChecker, ReturnStmtChecker,
+    SelectRankConstructChecker, SelectTypeChecker, StopChecker>;
+using StatementSemanticsPass3 = SemanticsVisitor<AssignmentChecker>;
 
 static bool PerformStatementSemantics(
     SemanticsContext &context, parser::Program &program) {
@@ -212,6 +212,7 @@ static bool PerformStatementSemantics(
   StatementSemanticsPass1{context}.Walk(program);
   StatementSemanticsPass2 pass2{context};
   pass2.Walk(program);
+  StatementSemanticsPass3{context}.Walk(program);
   if (context.languageFeatures().IsEnabled(common::LanguageFeature::OpenACC)) {
     SemanticsVisitor<AccStructureChecker>{context}.Walk(program);
   }

flang/test/Lower/CUDA/cuda-data-transfer.cuf

@@ -403,3 +403,19 @@ end subroutine
 ! CHECK-LABEL: func.func @_QPsub20()
 ! CHECK-NOT: cuf.data_transfer
 ! CHECK: hlfir.assign
+
+subroutine sub21()
+  real, allocatable,device:: a(:,:), b(:,:)
+  real:: s
+  integer:: i,j,N=16
+  allocate(a(N,N),b(N,N))
+  do concurrent(i=1:N, j=1:N) reduce(+:s)
+    b(i,j)=a(i,j)**2
+    s=s+b(i,j)
+  end do
+end subroutine
+
+! CHECK-LABEL: func.func @_QPsub21()
+! CHECK: fir.do_concurrent.loop
+! CHECK-NOT: cuf.data_transfer
+! CHECK: hlfir.assign

flang/tools/bbc/bbc.cpp

@@ -223,6 +223,10 @@ static llvm::cl::opt<bool> enableCUDA("fcuda",
                                       llvm::cl::desc("enable CUDA Fortran"),
                                       llvm::cl::init(false));
 
+static llvm::cl::opt<bool> enableStdPar("stdpar",
+                                        llvm::cl::desc("enable stdpar"),
+                                        llvm::cl::init(false));
+
 static llvm::cl::opt<bool>
     disableCUDAWarpFunction("fcuda-disable-warp-function",
                             llvm::cl::desc("Disable CUDA Warp Function"),
@@ -608,6 +612,10 @@ int main(int argc, char **argv) {
     options.features.Enable(Fortran::common::LanguageFeature::CUDA);
   }
 
+  if (enableStdPar) {
+    options.features.Enable(Fortran::common::LanguageFeature::StdPar);
+  }
+
   if (disableCUDAWarpFunction) {
     options.features.Enable(
         Fortran::common::LanguageFeature::CudaWarpMatchFunction, false);