[flang-commits] [flang] a67566b - Allow do concurrent inside cuf kernel directive (#127693)
via flang-commits
flang-commits at lists.llvm.org
Thu Feb 20 14:05:48 PST 2025
Author: Zhen Wang
Date: 2025-02-20T14:05:44-08:00
New Revision: a67566b185c56ce84f6b858e431e4d412b40fdaa
URL: https://github.com/llvm/llvm-project/commit/a67566b185c56ce84f6b858e431e4d412b40fdaa
DIFF: https://github.com/llvm/llvm-project/commit/a67566b185c56ce84f6b858e431e4d412b40fdaa.diff
LOG: Allow do concurrent inside cuf kernel directive (#127693)
Allow do concurrent inside cuf kernel directive to avoid the following
Lowering error:
```
void {anonymous}::FirConverter::genFIR(const Fortran::parser::CUFKernelDoConstruct&): Assertion `bounds && "Expected bounds on the loop construct"' failed.
```
---------
Co-authored-by: Valentin Clement (バレンタイン クレメン) <clementval at gmail.com>
Added:
flang/test/Lower/CUDA/cuda-doconc.cuf
Modified:
flang/lib/Lower/Bridge.cpp
flang/lib/Semantics/check-cuda.cpp
flang/test/Semantics/cuf09.cuf
Removed:
################################################################################
diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp
index 1b24ed12e04f1..f824e4c621c8e 100644
--- a/flang/lib/Lower/Bridge.cpp
+++ b/flang/lib/Lower/Bridge.cpp
@@ -3114,50 +3114,127 @@ class FirConverter : public Fortran::lower::AbstractConverter {
llvm::SmallVector<mlir::Value> ivValues;
Fortran::lower::pft::Evaluation *loopEval =
&getEval().getFirstNestedEvaluation();
- for (unsigned i = 0; i < nestedLoops; ++i) {
- const Fortran::parser::LoopControl *loopControl;
- mlir::Location crtLoc = loc;
- if (i == 0) {
- loopControl = &*outerDoConstruct->GetLoopControl();
- crtLoc =
- genLocation(Fortran::parser::FindSourceLocation(outerDoConstruct));
- } else {
- auto *doCons = loopEval->getIf<Fortran::parser::DoConstruct>();
- assert(doCons && "expect do construct");
- loopControl = &*doCons->GetLoopControl();
- crtLoc = genLocation(Fortran::parser::FindSourceLocation(*doCons));
+ if (outerDoConstruct->IsDoConcurrent()) {
+ // Handle DO CONCURRENT
+ locs.push_back(
+ genLocation(Fortran::parser::FindSourceLocation(outerDoConstruct)));
+ const Fortran::parser::LoopControl *loopControl =
+ &*outerDoConstruct->GetLoopControl();
+ const auto &concurrent =
+ std::get<Fortran::parser::LoopControl::Concurrent>(loopControl->u);
+
+ if (!std::get<std::list<Fortran::parser::LocalitySpec>>(concurrent.t)
+ .empty())
+ TODO(loc, "DO CONCURRENT with locality spec");
+
+ const auto &concurrentHeader =
+ std::get<Fortran::parser::ConcurrentHeader>(concurrent.t);
+ const auto &controls =
+ std::get<std::list<Fortran::parser::ConcurrentControl>>(
+ concurrentHeader.t);
+
+ for (const auto &control : controls) {
+ mlir::Value lb = fir::getBase(genExprValue(
+ *Fortran::semantics::GetExpr(std::get<1>(control.t)), stmtCtx));
+ mlir::Value ub = fir::getBase(genExprValue(
+ *Fortran::semantics::GetExpr(std::get<2>(control.t)), stmtCtx));
+ mlir::Value step;
+
+ if (const auto &expr =
+ std::get<std::optional<Fortran::parser::ScalarIntExpr>>(
+ control.t))
+ step = fir::getBase(
+ genExprValue(*Fortran::semantics::GetExpr(*expr), stmtCtx));
+ else
+ step = builder->create<mlir::arith::ConstantIndexOp>(
+ loc, 1); // Use index type directly
+
+ // Ensure lb, ub, and step are of index type using fir.convert
+ mlir::Type indexType = builder->getIndexType();
+ lb = builder->create<fir::ConvertOp>(loc, indexType, lb);
+ ub = builder->create<fir::ConvertOp>(loc, indexType, ub);
+ step = builder->create<fir::ConvertOp>(loc, indexType, step);
+
+ lbs.push_back(lb);
+ ubs.push_back(ub);
+ steps.push_back(step);
+
+ const auto &name = std::get<Fortran::parser::Name>(control.t);
+
+ // Handle induction variable
+ mlir::Value ivValue = getSymbolAddress(*name.symbol);
+ std::size_t ivTypeSize = name.symbol->size();
+ if (ivTypeSize == 0)
+ llvm::report_fatal_error("unexpected induction variable size");
+ mlir::Type ivTy = builder->getIntegerType(ivTypeSize * 8);
+
+ if (!ivValue) {
+ // DO CONCURRENT induction variables are not mapped yet since they are
+ // local to the DO CONCURRENT scope.
+ mlir::OpBuilder::InsertPoint insPt = builder->saveInsertionPoint();
+ builder->setInsertionPointToStart(builder->getAllocaBlock());
+ ivValue = builder->createTemporaryAlloc(
+ loc, ivTy, toStringRef(name.symbol->name()));
+ builder->restoreInsertionPoint(insPt);
+ }
+
+ // Create the hlfir.declare operation using the symbol's name
+ auto declareOp = builder->create<hlfir::DeclareOp>(
+ loc, ivValue, toStringRef(name.symbol->name()));
+ ivValue = declareOp.getResult(0);
+
+ // Bind the symbol to the declared variable
+ bindSymbol(*name.symbol, ivValue);
+ ivValues.push_back(ivValue);
+ ivTypes.push_back(ivTy);
+ ivLocs.push_back(loc);
}
+ } else {
+ for (unsigned i = 0; i < nestedLoops; ++i) {
+ const Fortran::parser::LoopControl *loopControl;
+ mlir::Location crtLoc = loc;
+ if (i == 0) {
+ loopControl = &*outerDoConstruct->GetLoopControl();
+ crtLoc = genLocation(
+ Fortran::parser::FindSourceLocation(outerDoConstruct));
+ } else {
+ auto *doCons = loopEval->getIf<Fortran::parser::DoConstruct>();
+ assert(doCons && "expect do construct");
+ loopControl = &*doCons->GetLoopControl();
+ crtLoc = genLocation(Fortran::parser::FindSourceLocation(*doCons));
+ }
- locs.push_back(crtLoc);
-
- const Fortran::parser::LoopControl::Bounds *bounds =
- std::get_if<Fortran::parser::LoopControl::Bounds>(&loopControl->u);
- assert(bounds && "Expected bounds on the loop construct");
-
- Fortran::semantics::Symbol &ivSym =
- bounds->name.thing.symbol->GetUltimate();
- ivValues.push_back(getSymbolAddress(ivSym));
-
- lbs.push_back(builder->createConvert(
- crtLoc, idxTy,
- fir::getBase(genExprValue(*Fortran::semantics::GetExpr(bounds->lower),
- stmtCtx))));
- ubs.push_back(builder->createConvert(
- crtLoc, idxTy,
- fir::getBase(genExprValue(*Fortran::semantics::GetExpr(bounds->upper),
- stmtCtx))));
- if (bounds->step)
- steps.push_back(builder->createConvert(
+ locs.push_back(crtLoc);
+
+ const Fortran::parser::LoopControl::Bounds *bounds =
+ std::get_if<Fortran::parser::LoopControl::Bounds>(&loopControl->u);
+ assert(bounds && "Expected bounds on the loop construct");
+
+ Fortran::semantics::Symbol &ivSym =
+ bounds->name.thing.symbol->GetUltimate();
+ ivValues.push_back(getSymbolAddress(ivSym));
+
+ lbs.push_back(builder->createConvert(
crtLoc, idxTy,
fir::getBase(genExprValue(
- *Fortran::semantics::GetExpr(bounds->step), stmtCtx))));
- else // If `step` is not present, assume it is `1`.
- steps.push_back(builder->createIntegerConstant(loc, idxTy, 1));
-
- ivTypes.push_back(idxTy);
- ivLocs.push_back(crtLoc);
- if (i < nestedLoops - 1)
- loopEval = &*std::next(loopEval->getNestedEvaluations().begin());
+ *Fortran::semantics::GetExpr(bounds->lower), stmtCtx))));
+ ubs.push_back(builder->createConvert(
+ crtLoc, idxTy,
+ fir::getBase(genExprValue(
+ *Fortran::semantics::GetExpr(bounds->upper), stmtCtx))));
+ if (bounds->step)
+ steps.push_back(builder->createConvert(
+ crtLoc, idxTy,
+ fir::getBase(genExprValue(
+ *Fortran::semantics::GetExpr(bounds->step), stmtCtx))));
+ else // If `step` is not present, assume it is `1`.
+ steps.push_back(builder->createIntegerConstant(loc, idxTy, 1));
+
+ ivTypes.push_back(idxTy);
+ ivLocs.push_back(crtLoc);
+ if (i < nestedLoops - 1)
+ loopEval = &*std::next(loopEval->getNestedEvaluations().begin());
+ }
}
auto op = builder->create<cuf::KernelOp>(
diff --git a/flang/lib/Semantics/check-cuda.cpp b/flang/lib/Semantics/check-cuda.cpp
index c85a84ea5527f..cb7a383284e63 100644
--- a/flang/lib/Semantics/check-cuda.cpp
+++ b/flang/lib/Semantics/check-cuda.cpp
@@ -525,6 +525,21 @@ static int DoConstructTightNesting(
return 0;
}
innerBlock = &std::get<parser::Block>(doConstruct->t);
+ if (doConstruct->IsDoConcurrent()) {
+ const auto &loopControl = doConstruct->GetLoopControl();
+ if (loopControl) {
+ if (const auto *concurrentControl{
+ std::get_if<parser::LoopControl::Concurrent>(&loopControl->u)}) {
+ const auto &concurrentHeader =
+ std::get<Fortran::parser::ConcurrentHeader>(concurrentControl->t);
+ const auto &controls =
+ std::get<std::list<Fortran::parser::ConcurrentControl>>(
+ concurrentHeader.t);
+ return controls.size();
+ }
+ }
+ return 0;
+ }
if (innerBlock->size() == 1) {
if (const auto *execConstruct{
std::get_if<parser::ExecutableConstruct>(&innerBlock->front().u)}) {
@@ -598,9 +613,14 @@ void CUDAChecker::Enter(const parser::CUFKernelDoConstruct &x) {
std::get<std::optional<parser::DoConstruct>>(x.t))};
const parser::Block *innerBlock{nullptr};
if (DoConstructTightNesting(doConstruct, innerBlock) < depth) {
- context_.Say(source,
- "!$CUF KERNEL DO (%jd) must be followed by a DO construct with tightly nested outer levels of counted DO loops"_err_en_US,
- std::intmax_t{depth});
+ if (doConstruct && doConstruct->IsDoConcurrent())
+ context_.Say(source,
+ "!$CUF KERNEL DO (%jd) must be followed by a DO CONCURRENT construct with at least %jd indices"_err_en_US,
+ std::intmax_t{depth}, std::intmax_t{depth});
+ else
+ context_.Say(source,
+ "!$CUF KERNEL DO (%jd) must be followed by a DO construct with tightly nested outer levels of counted DO loops"_err_en_US,
+ std::intmax_t{depth});
}
if (innerBlock) {
DeviceContextChecker<true>{context_}.Check(*innerBlock);
diff --git a/flang/test/Lower/CUDA/cuda-doconc.cuf b/flang/test/Lower/CUDA/cuda-doconc.cuf
new file mode 100644
index 0000000000000..32cd1676b22f4
--- /dev/null
+++ b/flang/test/Lower/CUDA/cuda-doconc.cuf
@@ -0,0 +1,39 @@
+! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s
+
+! Check if do concurrent works inside cuf kernel directive
+
+subroutine doconc1
+ integer :: i, n
+ integer, managed :: a(3)
+ a(:) = -1
+ n = 3
+ n = n - 1
+ !$cuf kernel do
+ do concurrent(i=1:n)
+ a(i) = 1
+ end do
+end
+
+! CHECK: func.func @_QPdoconc1() {
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc1Ei"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: cuf.kernel<<<*, *>>>
+! CHECK: %{{.*}} = fir.load %[[DECL]]#0 : !fir.ref<i32>
+
+subroutine doconc2
+ integer :: i, j, m, n
+ integer, managed :: a(2, 4)
+ m = 2
+ n = 4
+ a(:,:) = -1
+ !$cuf kernel do
+ do concurrent(i=1:m,j=1:n)
+ a(i,j) = i+j
+ end do
+end
+
+! CHECK: func.func @_QPdoconc2() {
+! CHECK: %[[DECLI:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc2Ei"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: %[[DECLJ:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc2Ej"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: cuf.kernel<<<*, *>>> (%arg0 : i32, %arg1 : i32) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index) {
+! CHECK: %{{.*}} = fir.load %[[DECLI]]#0 : !fir.ref<i32>
+! CHECK: %{{.*}} = fir.load %[[DECLJ]]#0 : !fir.ref<i32>
diff --git a/flang/test/Semantics/cuf09.cuf b/flang/test/Semantics/cuf09.cuf
index 7d32e0d70ba36..a8c62db65c6d5 100644
--- a/flang/test/Semantics/cuf09.cuf
+++ b/flang/test/Semantics/cuf09.cuf
@@ -133,6 +133,10 @@ program main
!$cuf kernel do <<< 1, 2 >>>
do concurrent (j=1:10)
end do
+ !ERROR: !$CUF KERNEL DO (2) must be followed by a DO CONCURRENT construct with at least 2 indices
+ !$cuf kernel do(2) <<< 1, 2 >>>
+ do concurrent (j=1:10)
+ end do
!$cuf kernel do <<< 1, 2 >>>
do 1 j=1,10
1 continue ! ok
More information about the flang-commits
mailing list