[Mlir-commits] [flang] [llvm] [mlir] [Flang][OpenMP] Fix non-contiguous array section target updates (PR #235006)
Akash Banerjee
llvmlistbot at llvm.org
Fri Oct 23 10:47:21 PDT 2026
https://github.com/TIFitis updated https://github.com/llvm/llvm-project/pull/235006
>From e85ab4dc6bda20a4188fad522bb636a3fc211302 Mon Sep 17 00:00:00 2001
From: Akash Banerjee <Akash.Banerjee at amd.com>
Date: Thu, 22 Oct 2026 15:29:21 +0100
Subject: [PATCH 1/2] [Flang][OpenMP] Fix non-contiguous array section target
updates
Multidimensional sections such as A(:,NB+1:,:,:) were transferred as one contiguous block, overwriting elements outside the section and leaving later slices unchanged. Emit non-contiguous runtime descriptors that preserve the bounds and physical strides so TO/FROM updates transfer only the selected elements.
Also fix runtime dimension merging, empty-transfer handling, large-count traversal, and character bounds and stride normalization. Add lowering, LLVM translation, and offload regressions covering these cases.
Co-authored-by: Codex <codex at openai.com>
---
.../Optimizer/Builder/DirectivesCommon.h | 13 +-
flang/lib/Lower/OpenMP/ClauseProcessor.cpp | 4 +-
flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp | 43 +++-
.../bounds-generation-for-char-arrays.f90 | 127 +++++++++++-
.../OpenMP/dynamic-len-char-bounds-gen.f90 | 34 +++
.../OpenMP/target-update-non-contiguous.f90 | 59 ++++++
mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td | 12 ++
.../OpenMP/OpenMPToLLVMIRTranslation.cpp | 123 +++++++++--
.../omptarget-update-non-contiguous.mlir | 194 ++++++++++++++++++
offload/libomptarget/omptarget.cpp | 39 ++--
.../fortran/dtype-char-array-map-2.f90 | 46 ++++-
.../fortran/target-update-empty.f90 | 52 +++++
.../fortran/target-update-noncontiguous.f90 | 137 +++++++++++++
.../fortran/target-update-wide-character.f90 | 74 +++++++
.../test/offloading/non_contiguous_update.cpp | 24 +++
.../non_contiguous_update_empty.cpp | 107 ++++++++++
16 files changed, 1040 insertions(+), 48 deletions(-)
create mode 100644 flang/test/Lower/OpenMP/target-update-non-contiguous.f90
create mode 100644 mlir/test/Target/LLVMIR/omptarget-update-non-contiguous.mlir
create mode 100644 offload/test/offloading/fortran/target-update-empty.f90
create mode 100644 offload/test/offloading/fortran/target-update-noncontiguous.f90
create mode 100644 offload/test/offloading/fortran/target-update-wide-character.f90
create mode 100644 offload/test/offloading/non_contiguous_update_empty.cpp
diff --git a/flang/include/flang/Optimizer/Builder/DirectivesCommon.h b/flang/include/flang/Optimizer/Builder/DirectivesCommon.h
index 1315eca5622491..795c952aa462aa 100644
--- a/flang/include/flang/Optimizer/Builder/DirectivesCommon.h
+++ b/flang/include/flang/Optimizer/Builder/DirectivesCommon.h
@@ -151,6 +151,15 @@ genBoundsOpFromBoxChar(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Type lenType = builder.getCharacterLengthType();
mlir::Value zero = builder.createIntegerConstant(loc, idxTy, 0);
mlir::Value one = builder.createIntegerConstant(loc, idxTy, 1);
+ auto boxCharType =
+ mlir::cast<fir::BoxCharType>(fir::unwrapRefType(info.addr.getType()));
+ unsigned codeUnitBytes =
+ builder.getKindMap().getCharacterBitsize(boxCharType.getKind()) / 8;
+ // The extent counts character code units, but the stride is in bytes.
+ mlir::Value byteStride =
+ codeUnitBytes == 1
+ ? one
+ : builder.createIntegerConstant(loc, idxTy, codeUnitBytes);
using ExtentAndStride = std::tuple<mlir::Value, mlir::Value>;
auto [extent, stride] = [&]() -> ExtentAndStride {
if (info.isPresent) {
@@ -169,7 +178,7 @@ genBoundsOpFromBoxChar(fir::FirOpBuilder &builder, mlir::Location loc,
auto unboxed = fir::UnboxCharOp::create(builder, loc, refType,
lenType, boxChar);
mlir::SmallVector<mlir::Value> results = {unboxed.getResult(1),
- one};
+ byteStride};
fir::ResultOp::create(builder, loc, results);
})
.genElse([&]() {
@@ -189,7 +198,7 @@ genBoundsOpFromBoxChar(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::Type refType = builder.getRefType(boxCharType.getEleTy());
auto unboxed =
fir::UnboxCharOp::create(builder, loc, refType, lenType, boxChar);
- return {unboxed.getResult(1), one};
+ return {unboxed.getResult(1), byteStride};
}();
mlir::Value ub = mlir::arith::SubIOp::create(builder, loc, extent, one);
diff --git a/flang/lib/Lower/OpenMP/ClauseProcessor.cpp b/flang/lib/Lower/OpenMP/ClauseProcessor.cpp
index 34961d3d44ba2f..3494f275935f4a 100644
--- a/flang/lib/Lower/OpenMP/ClauseProcessor.cpp
+++ b/flang/lib/Lower/OpenMP/ClauseProcessor.cpp
@@ -1868,7 +1868,9 @@ void ClauseProcessor::processMapObjects(
mlir::omp::MapBoundsType>(
converter, firOpBuilder, semaCtx, stmtCtx, *object.sym(),
object.ref(), clauseLocation, asFortran, bounds,
- treatIndexAsSection);
+ treatIndexAsSection, /*unwrapFirBox=*/true,
+ /*genDefaultBounds=*/true,
+ /*strideIncludeLowerExtent=*/isMotionModifier);
mlir::Value baseOp = info.rawInput;
if (object.sym()->owner().IsDerivedType() && !isMotionModifier) {
diff --git a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
index 35819444387e03..a6cfca049ff859 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
@@ -57,16 +57,16 @@ struct MapInfoOpConversion
mlir::omp::MapBoundsOp
createBoundsForCharString(mlir::ConversionPatternRewriter &rewriter,
- unsigned int len, mlir::Location loc) const {
+ int64_t len, mlir::Location loc) const {
mlir::Type i64Ty = rewriter.getIntegerType(64);
auto lBound = mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, 0);
- auto uBoundAndExt =
- mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, len - 1);
+ auto uBound = mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, len - 1);
+ auto extent = mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, len);
auto stride = mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, 1);
auto baseLb = mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, 1);
auto mapBoundType = rewriter.getType<mlir::omp::MapBoundsType>();
return mlir::omp::MapBoundsOp::create(rewriter, loc, mapBoundType, lBound,
- uBoundAndExt, uBoundAndExt, stride,
+ uBound, extent, stride,
/*strideInBytes*/ false, baseLb);
}
@@ -164,8 +164,41 @@ struct MapInfoOpConversion
curOp, resTypes, adaptor.getOperands(), newProperties,
curOp->getDiscardableAttrDictionary().getValue());
if (mapBoundsOp) {
+ rewriter.setInsertionPoint(newOp);
+ // Bounds are ordered from fastest to slowest varying dimension. Each
+ // string contributes an innermost dimension of character code units.
+ // Element strides still count whole strings here, whereas descriptor
+ // byte strides already include the character length and kind.
+ llvm::SmallVector<mlir::Value> bounds{mapBoundsOp};
+ for (mlir::Value bound : newOp.getBounds()) {
+ auto original = bound.getDefiningOp<mlir::omp::MapBoundsOp>();
+ if (original.getStrideInBytes()) {
+ bounds.push_back(bound);
+ continue;
+ }
+ mlir::Value stride = original.getStride();
+ if (stride) {
+ // Bounds operands may use different integer widths. Match the i64
+ // character extent before scaling the physical element stride.
+ unsigned strideWidth =
+ mlir::cast<mlir::IntegerType>(stride.getType()).getWidth();
+ if (strideWidth < 64)
+ stride = mlir::LLVM::SExtOp::create(rewriter, newOp.getLoc(),
+ rewriter.getI64Type(), stride);
+ else if (strideWidth > 64)
+ stride = mlir::LLVM::TruncOp::create(rewriter, newOp.getLoc(),
+ rewriter.getI64Type(), stride);
+ stride = mlir::LLVM::MulOp::create(rewriter, newOp.getLoc(), stride,
+ mapBoundsOp.getExtent());
+ }
+ bounds.push_back(mlir::omp::MapBoundsOp::create(
+ rewriter, newOp.getLoc(), original.getType(),
+ original.getLowerBound(), original.getUpperBound(),
+ original.getExtent(), stride, /*strideInBytes=*/false,
+ original.getStartIdx()));
+ }
rewriter.startOpModification(newOp);
- newOp.getBoundsMutable().append(mlir::ValueRange{mapBoundsOp});
+ newOp.getBoundsMutable().assign(bounds);
rewriter.finalizeOpModification(newOp);
}
diff --git a/flang/test/Fir/OpenMP/bounds-generation-for-char-arrays.f90 b/flang/test/Fir/OpenMP/bounds-generation-for-char-arrays.f90
index ef5f5f63f13ca7..4f1df036f5beab 100644
--- a/flang/test/Fir/OpenMP/bounds-generation-for-char-arrays.f90
+++ b/flang/test/Fir/OpenMP/bounds-generation-for-char-arrays.f90
@@ -25,11 +25,16 @@ module attributes {omp.is_target_device = false} {
// CHECK: %[[VAL_5:.*]] = omp.map.bounds lower_bound(%[[VAL_1]] : i64) upper_bound(%[[VAL_0]] : i64) extent(%[[VAL_3]] : i64) stride(%[[VAL_2]] : i64) start_idx(%[[VAL_2]] : i64)
// CHECK: %[[VAL_6:.*]] = llvm.mlir.constant(0 : i64) : i64
// CHECK: %[[VAL_7:.*]] = llvm.mlir.constant(15 : i64) : i64
-// CHECK: %[[VAL_8:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[VAL_8:.*]] = llvm.mlir.constant(16 : i64) : i64
// CHECK: %[[VAL_9:.*]] = llvm.mlir.constant(1 : i64) : i64
-// CHECK: %[[VAL_10:.*]] = omp.map.bounds lower_bound(%[[VAL_6]] : i64) upper_bound(%[[VAL_7]] : i64) extent(%[[VAL_7]] : i64) stride(%[[VAL_8]] : i64) start_idx(%[[VAL_9]] : i64)
-// CHECK: %[[VAL_11:.*]] = omp.map.info var_ptr(%[[ARG0]] : !llvm.ptr, i8) map_clauses(tofrom) capture(ByRef) bounds(%[[VAL_4]], %[[VAL_5]], %[[VAL_10]]) name("") -> !llvm.ptr
-// CHECK: omp.target kernel_type(generic) map_entries(%[[VAL_11]] -> %[[VAL_12:.*]] : !llvm.ptr) {
+// CHECK: %[[VAL_10:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[LEN_BOUND:.*]] = omp.map.bounds lower_bound(%[[VAL_6]] : i64) upper_bound(%[[VAL_7]] : i64) extent(%[[VAL_8]] : i64) stride(%[[VAL_9]] : i64) start_idx(%[[VAL_10]] : i64)
+// CHECK: %[[STRIDE0:.*]] = llvm.mul %[[VAL_2]], %[[VAL_8]] : i64
+// CHECK: %[[BOUND0:.*]] = omp.map.bounds lower_bound(%[[VAL_1]] : i64) upper_bound(%[[VAL_0]] : i64) extent(%[[VAL_3]] : i64) stride(%[[STRIDE0]] : i64)
+// CHECK: %[[STRIDE1:.*]] = llvm.mul %[[VAL_2]], %[[VAL_8]] : i64
+// CHECK: %[[BOUND1:.*]] = omp.map.bounds lower_bound(%[[VAL_1]] : i64) upper_bound(%[[VAL_0]] : i64) extent(%[[VAL_3]] : i64) stride(%[[STRIDE1]] : i64)
+// CHECK: %[[MAP:.*]] = omp.map.info var_ptr(%[[ARG0]] : !llvm.ptr, i8) map_clauses(tofrom) capture(ByRef) bounds(%[[LEN_BOUND]], %[[BOUND0]], %[[BOUND1]]) name("") -> !llvm.ptr
+// CHECK: omp.target kernel_type(generic) map_entries(%[[MAP]] -> %{{.*}} : !llvm.ptr) {
// CHECK: omp.terminator
// CHECK: }
// CHECK: llvm.return
@@ -83,14 +88,122 @@ module attributes {omp.is_target_device = false} {
// CHECK: %[[VAL_22:.*]] = llvm.getelementptr %[[ARG0]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<2 x array<3 x i64>>)>
// CHECK: %[[VAL_23:.*]] = llvm.mlir.constant(0 : i64) : i64
// CHECK: %[[VAL_24:.*]] = llvm.mlir.constant(15 : i64) : i64
-// CHECK: %[[VAL_25:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[VAL_25:.*]] = llvm.mlir.constant(16 : i64) : i64
// CHECK: %[[VAL_26:.*]] = llvm.mlir.constant(1 : i64) : i64
-// CHECK: %[[VAL_27:.*]] = omp.map.bounds lower_bound(%[[VAL_23]] : i64) upper_bound(%[[VAL_24]] : i64) extent(%[[VAL_24]] : i64) stride(%[[VAL_25]] : i64) start_idx(%[[VAL_26]] : i64)
-// CHECK: %[[VAL_28:.*]] = omp.map.info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<2 x array<3 x i64>>)>) map_clauses(tofrom) capture(ByRef) var_ptr_ptr(%[[VAL_22]] : !llvm.ptr, i8) bounds(%[[VAL_12]], %[[VAL_21]], %[[VAL_27]]) name("") -> !llvm.ptr
+// CHECK: %[[VAL_27:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[LEN_BOUND:.*]] = omp.map.bounds lower_bound(%[[VAL_23]] : i64) upper_bound(%[[VAL_24]] : i64) extent(%[[VAL_25]] : i64) stride(%[[VAL_26]] : i64) start_idx(%[[VAL_27]] : i64)
+// CHECK-NOT: llvm.mul
+// CHECK: %[[VAL_28:.*]] = omp.map.info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<2 x array<3 x i64>>)>) map_clauses(tofrom) capture(ByRef) var_ptr_ptr(%[[VAL_22]] : !llvm.ptr, i8) bounds(%[[LEN_BOUND]], %[[VAL_12]], %[[VAL_21]]) name("") -> !llvm.ptr
// CHECK: %[[VAL_29:.*]] = omp.map.info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<2 x array<3 x i64>>)>) map_clauses(to) capture(ByRef) members(%[[VAL_28]] : [0] : !llvm.ptr) name("csv_chem_list_a") -> !llvm.ptr
// CHECK: omp.target kernel_type(generic) map_entries(%[[VAL_29]] -> %[[VAL_30:.*]], %[[VAL_28]] -> %[[VAL_31:.*]] : !llvm.ptr, !llvm.ptr) {
// CHECK: omp.terminator
// CHECK: }
// CHECK: llvm.return
// CHECK: }
+
+ // Motion-clause element strides are measured in whole strings. Normalizing
+ // the element type to i8 adds LEN as the first (fastest varying) dimension
+ // and scales the original physical strides by LEN.
+ func.func @char_array_section(%arg0 : !fir.ref<!fir.array<4x5x!fir.char<1,7>>>) {
+ %c1 = arith.constant 1 : index
+ %c2 = arith.constant 2 : index
+ %c4 = arith.constant 4 : index
+ %c5 = arith.constant 5 : index
+ %b0 = omp.map.bounds lower_bound(%c1 : index) upper_bound(%c2 : index) extent(%c4 : index) stride(%c1 : index) start_idx(%c1 : index)
+ %b1 = omp.map.bounds lower_bound(%c2 : index) upper_bound(%c4 : index) extent(%c5 : index) stride(%c4 : index) start_idx(%c1 : index)
+ %map = omp.map.info var_ptr(%arg0 : !fir.ref<!fir.array<4x5x!fir.char<1,7>>>, !fir.array<4x5x!fir.char<1,7>>) map_clauses(to) capture(ByRef) bounds(%b0, %b1) -> !fir.ref<!fir.array<4x5x!fir.char<1,7>>>
+ omp.target_update map_entries(%map : !fir.ref<!fir.array<4x5x!fir.char<1,7>>>)
+ return
+ }
+
+// CHECK-LABEL: llvm.func @char_array_section(
+// CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[TWO:.*]] = llvm.mlir.constant(2 : i64) : i64
+// CHECK: %[[FOUR:.*]] = llvm.mlir.constant(4 : i64) : i64
+// CHECK: %[[FIVE:.*]] = llvm.mlir.constant(5 : i64) : i64
+// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i64) : i64
+// CHECK: %[[SIX:.*]] = llvm.mlir.constant(6 : i64) : i64
+// CHECK: %[[LEN:.*]] = llvm.mlir.constant(7 : i64) : i64
+// CHECK: %[[LEN_BOUND:.*]] = omp.map.bounds lower_bound(%[[ZERO]] : i64) upper_bound(%[[SIX]] : i64) extent(%[[LEN]] : i64) stride(%{{.*}} : i64)
+// CHECK: %[[STRIDE0:.*]] = llvm.mul %[[ONE]], %[[LEN]] : i64
+// CHECK: %[[BOUND0:.*]] = omp.map.bounds lower_bound(%[[ONE]] : i64) upper_bound(%[[TWO]] : i64) extent(%[[FOUR]] : i64) stride(%[[STRIDE0]] : i64)
+// CHECK: %[[STRIDE1:.*]] = llvm.mul %[[FOUR]], %[[LEN]] : i64
+// CHECK: %[[BOUND1:.*]] = omp.map.bounds lower_bound(%[[TWO]] : i64) upper_bound(%[[FOUR]] : i64) extent(%[[FIVE]] : i64) stride(%[[STRIDE1]] : i64)
+// CHECK: %[[MAP:.*]] = omp.map.info var_ptr(%{{.*}} : !llvm.ptr, i8) map_clauses(to) capture(ByRef) bounds(%[[LEN_BOUND]], %[[BOUND0]], %[[BOUND1]])
+// CHECK: omp.target_update map_entries(%[[MAP]] : !llvm.ptr)
+
+ // Wide characters use i32 code units. LEN scales element strides by three,
+ // while the subsequent LLVM translation accounts for four bytes per i32.
+ func.func @wide_char_array(%arg0 : !fir.ref<!fir.array<2x3x!fir.char<4,3>>>) {
+ %c0 = arith.constant 0 : index
+ %c1 = arith.constant 1 : index
+ %c2 = arith.constant 2 : index
+ %c3 = arith.constant 3 : index
+ %b0 = omp.map.bounds lower_bound(%c0 : index) upper_bound(%c1 : index) extent(%c2 : index) stride(%c1 : index) start_idx(%c1 : index)
+ %b1 = omp.map.bounds lower_bound(%c1 : index) upper_bound(%c2 : index) extent(%c3 : index) stride(%c2 : index) start_idx(%c1 : index)
+ %map = omp.map.info var_ptr(%arg0 : !fir.ref<!fir.array<2x3x!fir.char<4,3>>>, !fir.array<2x3x!fir.char<4,3>>) map_clauses(from) capture(ByRef) bounds(%b0, %b1) -> !fir.ref<!fir.array<2x3x!fir.char<4,3>>>
+ omp.target_update map_entries(%map : !fir.ref<!fir.array<2x3x!fir.char<4,3>>>)
+ return
+ }
+
+// CHECK-LABEL: llvm.func @wide_char_array(
+// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i64) : i64
+// CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[TWO:.*]] = llvm.mlir.constant(2 : i64) : i64
+// CHECK: %[[THREE:.*]] = llvm.mlir.constant(3 : i64) : i64
+// CHECK: %[[LEN:.*]] = llvm.mlir.constant(3 : i64) : i64
+// CHECK: %[[LEN_BOUND:.*]] = omp.map.bounds {{.*}} extent(%[[LEN]] : i64) stride(%{{.*}} : i64)
+// CHECK: %[[STRIDE0:.*]] = llvm.mul %[[ONE]], %[[LEN]] : i64
+// CHECK: %[[BOUND0:.*]] = omp.map.bounds lower_bound(%[[ZERO]] : i64) upper_bound(%[[ONE]] : i64) extent(%[[TWO]] : i64) stride(%[[STRIDE0]] : i64)
+// CHECK: %[[STRIDE1:.*]] = llvm.mul %[[TWO]], %[[LEN]] : i64
+// CHECK: %[[BOUND1:.*]] = omp.map.bounds lower_bound(%[[ONE]] : i64) upper_bound(%[[TWO]] : i64) extent(%[[THREE]] : i64) stride(%[[STRIDE1]] : i64)
+// CHECK: %[[MAP:.*]] = omp.map.info var_ptr(%{{.*}} : !llvm.ptr, i32) map_clauses(from) capture(ByRef) bounds(%[[LEN_BOUND]], %[[BOUND0]], %[[BOUND1]])
+// CHECK: omp.target_update map_entries(%[[MAP]] : !llvm.ptr)
+
+ // Fixed-length scalars need only the character-code-unit bound and have no
+ // array stride to scale.
+ func.func @scalar_char(%arg0 : !fir.ref<!fir.char<1,3>>) {
+ %map = omp.map.info var_ptr(%arg0 : !fir.ref<!fir.char<1,3>>, !fir.char<1,3>) map_clauses(to) capture(ByRef) -> !fir.ref<!fir.char<1,3>>
+ omp.target_update map_entries(%map : !fir.ref<!fir.char<1,3>>)
+ return
+ }
+
+// CHECK-LABEL: llvm.func @scalar_char(
+// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i64) : i64
+// CHECK: %[[TWO:.*]] = llvm.mlir.constant(2 : i64) : i64
+// CHECK: %[[LEN:.*]] = llvm.mlir.constant(3 : i64) : i64
+// CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[BASE_LB:.*]] = llvm.mlir.constant(1 : i64) : i64
+// CHECK: %[[LEN_BOUND:.*]] = omp.map.bounds lower_bound(%[[ZERO]] : i64) upper_bound(%[[TWO]] : i64) extent(%[[LEN]] : i64) stride(%[[ONE]] : i64) start_idx(%[[BASE_LB]] : i64)
+// CHECK-NOT: llvm.mul
+// CHECK: %[[MAP:.*]] = omp.map.info var_ptr(%{{.*}} : !llvm.ptr, i8) map_clauses(to) capture(ByRef) bounds(%[[LEN_BOUND]])
+// CHECK: omp.target_update map_entries(%[[MAP]] : !llvm.ptr)
+
+ // Each bounds operand has its own integer type. Normalize both narrower and
+ // wider physical strides to i64 before multiplying by the character length.
+ func.func @char_array_mixed_stride_types(%arg0 : !fir.ref<!fir.array<4x5x!fir.char<1,7>>>, %stride0 : i32, %stride1 : i128) {
+ %c0 = arith.constant 0 : i64
+ %c1 = arith.constant 1 : i64
+ %c3 = arith.constant 3 : i64
+ %c4 = arith.constant 4 : i64
+ %c5 = arith.constant 5 : i64
+ %b0 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c3 : i64) extent(%c4 : i64) stride(%stride0 : i32) start_idx(%c1 : i64)
+ %b1 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c4 : i64) extent(%c5 : i64) stride(%stride1 : i128) start_idx(%c1 : i64)
+ %map = omp.map.info var_ptr(%arg0 : !fir.ref<!fir.array<4x5x!fir.char<1,7>>>, !fir.array<4x5x!fir.char<1,7>>) map_clauses(to) capture(ByRef) bounds(%b0, %b1) -> !fir.ref<!fir.array<4x5x!fir.char<1,7>>>
+ omp.target_update map_entries(%map : !fir.ref<!fir.array<4x5x!fir.char<1,7>>>)
+ return
+ }
+
+// CHECK-LABEL: llvm.func @char_array_mixed_stride_types(
+// CHECK-SAME: %[[ARRAY:.*]]: !llvm.ptr, %[[NARROW:.*]]: i32, %[[WIDE:.*]]: i128)
+// CHECK: %[[LEN:.*]] = llvm.mlir.constant(7 : i64) : i64
+// CHECK: %[[LEN_BOUND:.*]] = omp.map.bounds {{.*}} extent(%[[LEN]] : i64)
+// CHECK: %[[NARROW64:.*]] = llvm.sext %[[NARROW]] : i32 to i64
+// CHECK: %[[STRIDE0:.*]] = llvm.mul %[[NARROW64]], %[[LEN]] : i64
+// CHECK: %[[BOUND0:.*]] = omp.map.bounds {{.*}} stride(%[[STRIDE0]] : i64)
+// CHECK: %[[WIDE64:.*]] = llvm.trunc %[[WIDE]] : i128 to i64
+// CHECK: %[[STRIDE1:.*]] = llvm.mul %[[WIDE64]], %[[LEN]] : i64
+// CHECK: %[[BOUND1:.*]] = omp.map.bounds {{.*}} stride(%[[STRIDE1]] : i64)
+// CHECK: %[[MAP:.*]] = omp.map.info var_ptr(%[[ARRAY]] : !llvm.ptr, i8) map_clauses(to) capture(ByRef) bounds(%[[LEN_BOUND]], %[[BOUND0]], %[[BOUND1]])
+// CHECK: omp.target_update map_entries(%[[MAP]] : !llvm.ptr)
}
diff --git a/flang/test/Lower/OpenMP/dynamic-len-char-bounds-gen.f90 b/flang/test/Lower/OpenMP/dynamic-len-char-bounds-gen.f90
index bcd1f606bd8232..435d2025f88e59 100644
--- a/flang/test/Lower/OpenMP/dynamic-len-char-bounds-gen.f90
+++ b/flang/test/Lower/OpenMP/dynamic-len-char-bounds-gen.f90
@@ -17,3 +17,37 @@ end subroutine TestCharLenBounds
!CHECK: %[[UB:.*]] = arith.subi %[[EXTENT]]#1, %[[STRIDE]] : index
!CHECK: %[[BOUNDS:.*]] = omp.map.bounds lower_bound(%[[LB_START_IDX]] : index) upper_bound(%[[UB]] : index) extent(%[[EXTENT]]#1 : index) stride(%[[STRIDE]] : index) start_idx(%[[LB_START_IDX]] : index) stride_in_bytes(true)
!CHECK: %{{.*}} = omp.map.info {{.*}} bounds(%[[BOUNDS]]) {{.*}}
+
+! A dynamic length is a count of code units. Wide-character byte strides
+! must include the kind's storage size, in both motion directions.
+! CHECK-LABEL: func.func @_QPupdate_wide_char(
+! CHECK: %[[ZERO:.*]] = arith.constant 0 : index
+! CHECK: %[[ONE:.*]] = arith.constant 1 : index
+! CHECK: %[[WIDTH:.*]] = arith.constant 4 : index
+! CHECK: %[[UNBOX:.*]]:2 = fir.unboxchar {{.*}} : (!fir.boxchar<4>) -> (!fir.ref<!fir.char<4,?>>, index)
+! CHECK: %[[UB:.*]] = arith.subi %[[UNBOX]]#1, %[[ONE]] : index
+! CHECK: %[[BOUNDS:.*]] = omp.map.bounds lower_bound(%[[ZERO]] : index) upper_bound(%[[UB]] : index) extent(%[[UNBOX]]#1 : index) stride(%[[WIDTH]] : index) start_idx(%[[ZERO]] : index) stride_in_bytes(true)
+! CHECK: omp.map.info {{.*}} map_clauses(to) {{.*}} bounds(%[[BOUNDS]])
+! CHECK: omp.target_update
+! CHECK: %[[WIDTH_FROM:.*]] = arith.constant 4 : index
+! CHECK: omp.map.bounds {{.*}} stride(%[[WIDTH_FROM]] : index) {{.*}} stride_in_bytes(true)
+! CHECK: omp.map.info {{.*}} map_clauses(from)
+! CHECK: omp.target_update
+subroutine update_wide_char(clen)
+ character(kind=4,len=*) :: clen
+ !$omp target update to(clen)
+ !$omp target update from(clen)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPupdate_optional_wide_char(
+! CHECK: %[[WIDTH:.*]] = arith.constant 2 : index
+! CHECK: %[[PRESENT_BOUNDS:.*]]:2 = fir.if {{.*}} -> (index, index) {
+! CHECK: fir.result {{.*}}, %[[WIDTH]] : index, index
+! CHECK: } else {
+! CHECK: fir.result %{{.*}}, %{{.*}} : index, index
+! CHECK: omp.map.bounds {{.*}} stride(%[[PRESENT_BOUNDS]]#1 : index) {{.*}} stride_in_bytes(true)
+! CHECK: omp.target_update
+subroutine update_optional_wide_char(clen)
+ character(kind=2,len=*), optional :: clen
+ !$omp target update to(clen)
+end subroutine
diff --git a/flang/test/Lower/OpenMP/target-update-non-contiguous.f90 b/flang/test/Lower/OpenMP/target-update-non-contiguous.f90
new file mode 100644
index 00000000000000..d9c3038ffb6b4f
--- /dev/null
+++ b/flang/test/Lower/OpenMP/target-update-non-contiguous.f90
@@ -0,0 +1,59 @@
+! RUN: %flang_fc1 -emit-hlfir -fopenmp %s -o - | FileCheck %s
+
+! A scalar subscript has extent one, but subsequent physical strides must
+! include the full source extent of that dimension. Both motion directions
+! need these strides to describe non-contiguous array sections.
+
+! CHECK-LABEL: func.func @_QPupdate_to(
+! CHECK: %[[ARRAY:.*]]:2 = hlfir.declare %arg0(
+! CHECK: %[[D0:.*]]:3 = fir.box_dims %[[ARRAY]]#0,
+! CHECK: %[[B0:.*]] = omp.map.bounds {{.*}} stride(%[[D0]]#2 : index) {{.*}} stride_in_bytes(true)
+! CHECK: %[[D1:.*]]:3 = fir.box_dims %[[ARRAY]]#0,
+! CHECK: %[[B1:.*]] = omp.map.bounds {{.*}} stride(%[[D1]]#2 : index) {{.*}} stride_in_bytes(true)
+! CHECK: %[[D2:.*]]:3 = fir.box_dims %[[ARRAY]]#0,
+! CHECK: %[[B2:.*]] = omp.map.bounds {{.*}} stride(%[[D2]]#2 : index) {{.*}} stride_in_bytes(true)
+! CHECK: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(to) {{.*}} bounds(%[[B0]], %[[B1]], %[[B2]])
+! CHECK: omp.target_update map_entries(%[[MAP]]
+subroutine update_to(a, n, m, k)
+ integer :: n, m, k
+ real :: a(n, m, k)
+ !$omp target update to(a(2, 3:, :))
+end subroutine
+
+! CHECK-LABEL: func.func @_QPupdate_from(
+! CHECK: %[[ARRAY:.*]]:2 = hlfir.declare %arg0(
+! CHECK: %[[D0:.*]]:3 = fir.box_dims %[[ARRAY]]#0,
+! CHECK: %[[B0:.*]] = omp.map.bounds {{.*}} stride(%[[D0]]#2 : index) {{.*}} stride_in_bytes(true)
+! CHECK: %[[D1:.*]]:3 = fir.box_dims %[[ARRAY]]#0,
+! CHECK: %[[B1:.*]] = omp.map.bounds {{.*}} stride(%[[D1]]#2 : index) {{.*}} stride_in_bytes(true)
+! CHECK: %[[D2:.*]]:3 = fir.box_dims %[[ARRAY]]#0,
+! CHECK: %[[B2:.*]] = omp.map.bounds {{.*}} stride(%[[D2]]#2 : index) {{.*}} stride_in_bytes(true)
+! CHECK: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(from) {{.*}} bounds(%[[B0]], %[[B1]], %[[B2]])
+! CHECK: omp.target_update map_entries(%[[MAP]]
+subroutine update_from(a, n, m, k)
+ integer :: n, m, k
+ real :: a(n, m, k)
+ !$omp target update from(a(2, 3:, :))
+end subroutine
+
+! CHECK-LABEL: func.func @_QPupdate_explicit(
+! CHECK: %[[ONE:.*]] = arith.constant 1 : index
+! CHECK: %[[FOUR:.*]] = arith.constant 4 : index
+! CHECK: %[[B0:.*]] = omp.map.bounds {{.*}} extent(%[[ONE]] : index) stride(%[[ONE]] : index)
+! CHECK: %[[TWENTY:.*]] = arith.constant 20 : index
+! CHECK: %[[B1:.*]] = omp.map.bounds {{.*}} stride(%[[FOUR]] : index)
+! CHECK: %[[B2:.*]] = omp.map.bounds {{.*}} stride(%[[TWENTY]] : index)
+! CHECK: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(to) {{.*}} bounds(%[[B0]], %[[B1]], %[[B2]])
+! CHECK: omp.target_update map_entries(%[[MAP]]
+! CHECK: %[[FROM_FOUR:.*]] = arith.constant 4 : index
+! CHECK: %[[FROM_B0:.*]] = omp.map.bounds {{.*}} extent(%{{.*}} : index) stride(%{{.*}} : index)
+! CHECK: %[[FROM_TWENTY:.*]] = arith.constant 20 : index
+! CHECK: %[[FROM_B1:.*]] = omp.map.bounds {{.*}} stride(%[[FROM_FOUR]] : index)
+! CHECK: %[[FROM_B2:.*]] = omp.map.bounds {{.*}} stride(%[[FROM_TWENTY]] : index)
+! CHECK: %[[FROM_MAP:.*]] = omp.map.info {{.*}} map_clauses(from) {{.*}} bounds(%[[FROM_B0]], %[[FROM_B1]], %[[FROM_B2]])
+! CHECK: omp.target_update map_entries(%[[FROM_MAP]]
+subroutine update_explicit(a)
+ real :: a(4, 5, 6)
+ !$omp target update to(a(2, 3:, :))
+ !$omp target update from(a(2, 3:, :))
+end subroutine
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
index 307b91ebe99af2..d6c71b2782b269 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
@@ -1406,6 +1406,18 @@ def MapBoundsOp : OpenMP_Op<"map.bounds",
requirement in data clauses - since a `lower_bound` of 0 means looking
at data at the zero offset from pointer.
+ For target-update translation, bounds are ordered from fastest to slowest
+ varying dimension. With `stride_in_bytes`, `stride` is the byte distance
+ between adjacent elements of that source dimension, including any faster
+ varying dimensions. Otherwise strides are measured in scalar elements.
+ When the mapped type is a nested LLVM array, translation reconstructs the
+ physical element strides from its complete shape; this also supports
+ existing static-array bounds that carry unit strides. When only the scalar
+ element type is available, each explicit element stride must already
+ include the complete source extents of faster varying dimensions. Selected
+ section lengths must not replace those source extents when computing a
+ physical stride. Explicit byte strides take precedence over the type shape.
+
This operation must have an `upper_bound` or `extent` (or both are allowed -
but not checked for consistency). When the source language's arrays are
not zero-based, the `start_idx` must specify the zero-position index.
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index 1f2024060d4025..7e400a536f6dfc 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -7953,6 +7953,98 @@ getOverlappedMembers(llvm::SmallVectorImpl<size_t> &overlapMapDataIdxs,
overlapMapDataIdxs.push_back(i);
}
+/// Describe a bounded update as a set of contiguous pieces. The runtime uses
+/// byte offsets and strides, with the slowest varying dimension first (the
+/// OpenMPIRBuilder reverses the vectors below when emitting the descriptor).
+static void addNonContiguousUpdate(LLVM::ModuleTranslation &moduleTranslation,
+ llvm::IRBuilderBase &builder,
+ MapInfoData &mapData, size_t mapDataIdx,
+ MapInfosTy &combinedInfo) {
+ auto mapOp = cast<omp::MapInfoOp>(mapData.MapClause[mapDataIdx]);
+ if (mapOp.getBounds().empty() || mapOp.getMapperId() ||
+ mapOp->getParentOfType<omp::DeclareMapperOp>() ||
+ bitEnumContainsAll(mapOp.getMapType(), omp::ClauseMapFlags::attach))
+ return;
+
+ llvm::Type *elementType = mapData.BaseType[mapDataIdx];
+ SmallVector<uint64_t> arrayExtents;
+ while (auto *arrayType = dyn_cast<llvm::ArrayType>(elementType)) {
+ arrayExtents.push_back(arrayType->getNumElements());
+ elementType = arrayType->getElementType();
+ }
+ llvm::Value *elementSize = builder.getInt64(
+ builder.GetInsertBlock()->getModule()->getDataLayout().getTypeAllocSize(
+ elementType));
+
+ // Include a byte dimension so that a non-unit descriptor stride in the
+ // innermost array dimension still transfers one element at a time.
+ llvm::OpenMPIRBuilder::MapValuesArrayTy offsets{builder.getInt64(0)};
+ llvm::OpenMPIRBuilder::MapValuesArrayTy counts{elementSize};
+ llvm::OpenMPIRBuilder::MapValuesArrayTy strides{builder.getInt64(1)};
+ llvm::Value *sourceStride = elementSize;
+ llvm::Value *contiguousSize = elementSize;
+ bool isContiguous = true;
+ for (auto [dim, bound] : llvm::enumerate(mapOp.getBounds())) {
+ auto bounds = cast<omp::MapBoundsOp>(bound.getDefiningOp());
+ auto getInt64 = [&](Value value) {
+ return builder.CreateSExtOrTrunc(moduleTranslation.lookupValue(value),
+ builder.getInt64Ty());
+ };
+ llvm::Value *lower = bounds.getLowerBound()
+ ? getInt64(bounds.getLowerBound())
+ : builder.getInt64(0);
+ llvm::Value *count = builder.CreateAdd(
+ builder.CreateSub(getInt64(bounds.getUpperBound()), lower),
+ builder.getInt64(1));
+ count =
+ builder.CreateSelect(builder.CreateICmpSGT(count, builder.getInt64(0)),
+ count, builder.getInt64(0));
+ llvm::Value *stride = sourceStride;
+ if (bounds.getStrideInBytes() && bounds.getStride())
+ stride = getInt64(bounds.getStride());
+ else if (arrayExtents.empty() && bounds.getStride())
+ stride = builder.CreateMul(getInt64(bounds.getStride()), elementSize);
+ offsets.push_back(builder.CreateMul(lower, stride));
+ counts.push_back(count);
+ strides.push_back(stride);
+ isContiguous &= stride == contiguousSize;
+ contiguousSize = builder.CreateMul(contiguousSize, count);
+
+ // Static array types carry the full source shape, including dimensions
+ // selected with scalar subscripts. Otherwise the frontend supplies the
+ // physical stride in elements (or descriptor bytes) for every dimension.
+ if (dim < arrayExtents.size())
+ sourceStride = builder.CreateMul(
+ sourceStride,
+ builder.getInt64(arrayExtents[arrayExtents.size() - dim - 1]));
+ else if (bounds.getExtent())
+ sourceStride =
+ builder.CreateMul(sourceStride, getInt64(bounds.getExtent()));
+ }
+ if (isContiguous)
+ return;
+
+ llvm::Value *base = mapData.OriginalValue[mapDataIdx];
+ if (mapOp.getVarPtrPtr())
+ base = builder.CreateLoad(builder.getPtrTy(), base);
+ // Absent optional arrays and unallocated descriptors must not issue
+ // transfers, including sections with explicit nonempty bounds.
+ counts[1] = builder.CreateSelect(builder.CreateIsNull(base),
+ builder.getInt64(0), counts[1]);
+ // The descriptor offsets already locate the first selected element. The
+ // runtime needs the pointee base, not the descriptor or a section address.
+ combinedInfo.BasePointers.back() = base;
+ combinedInfo.Types.back() |=
+ llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG;
+ auto &nonContig = combinedInfo.NonContigInfo;
+ nonContig.IsNonContiguous = true;
+ nonContig.Dims.resize(combinedInfo.Types.size(), 1);
+ nonContig.Dims.back() = counts.size();
+ nonContig.Offsets.push_back(std::move(offsets));
+ nonContig.Counts.push_back(std::move(counts));
+ nonContig.Strides.push_back(std::move(strides));
+}
+
/// This function handles the insertion of a single item of map data from
/// MapInfoData into the OMPIRBuilder's MapInfo list. Utilising this function
/// means the map being inserted can be treated as a non-parent map entity,
@@ -7962,14 +8054,13 @@ getOverlappedMembers(llvm::SmallVectorImpl<size_t> &overlapMapDataIdxs,
/// a ref_ptr or ref_ptee map, if it's a member of a record, what construct
/// the map belongs to and the various map type bit flags that are set for
/// the map.
-static void
-processIndividualMap(llvm::IRBuilderBase &builder,
- llvm::OpenMPIRBuilder &ompBuilder, MapInfoData &mapData,
- size_t mapDataIdx, MapInfosTy &combinedInfo,
- TargetDirectiveEnumTy targetDirective,
- llvm::omp::OpenMPOffloadMappingFlags memberOfFlag =
- llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE,
- bool isTargetParam = true, int mapDataParentIdx = -1) {
+static void processIndividualMap(
+ LLVM::ModuleTranslation &moduleTranslation, llvm::IRBuilderBase &builder,
+ llvm::OpenMPIRBuilder &ompBuilder, MapInfoData &mapData, size_t mapDataIdx,
+ MapInfosTy &combinedInfo, TargetDirectiveEnumTy targetDirective,
+ llvm::omp::OpenMPOffloadMappingFlags memberOfFlag =
+ llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE,
+ bool isTargetParam = true, int mapDataParentIdx = -1) {
auto mapFlag = mapData.Types[mapDataIdx];
auto mapInfoOp = llvm::cast<omp::MapInfoOp>(mapData.MapClause[mapDataIdx]);
@@ -8062,6 +8153,9 @@ processIndividualMap(llvm::IRBuilderBase &builder,
builder.CreateIsNull(mapData.Pointers[mapDataIdx]),
builder.getInt64(0), mapData.Sizes[mapDataIdx])
: mapData.Sizes[mapDataIdx]);
+ if (targetDirective == TargetDirectiveEnumTy::TargetUpdate)
+ addNonContiguousUpdate(moduleTranslation, builder, mapData, mapDataIdx,
+ combinedInfo);
}
// This creates two insertions into the MapInfosTy data structure for the
@@ -8351,8 +8445,8 @@ static void processMapWithMembersOf(LLVM::ModuleTranslation &moduleTranslation,
// currently do), whereas we treat them as arrays in all cases
// currently.
processIndividualMap(
- builder, ompBuilder, mapData, memberDataIdx, combinedInfo,
- targetDirective,
+ moduleTranslation, builder, ompBuilder, mapData, memberDataIdx,
+ combinedInfo, targetDirective,
/*MemberOfFlag=*/llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE,
/*isTargetParam=*/true, mapDataIndex);
return;
@@ -8392,7 +8486,8 @@ static void processMapWithMembersOf(LLVM::ModuleTranslation &moduleTranslation,
targetDirective);
} else {
processIndividualMap(
- builder, ompBuilder, mapData, idx, combinedInfo, targetDirective,
+ moduleTranslation, builder, ompBuilder, mapData, idx, combinedInfo,
+ targetDirective,
parentIsPrivatizeableAttach
? llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE
: memberOfFlag,
@@ -8519,9 +8614,11 @@ static void genMapInfos(llvm::IRBuilderBase &builder,
continue;
}
- processIndividualMap(builder, *ompBuilder, mapData, i, combinedInfo,
- targetDirective);
+ processIndividualMap(moduleTranslation, builder, *ompBuilder, mapData, i,
+ combinedInfo, targetDirective);
}
+ if (!combinedInfo.NonContigInfo.Dims.empty())
+ combinedInfo.NonContigInfo.Dims.resize(combinedInfo.Types.size(), 1);
}
static llvm::Expected<llvm::Function *>
diff --git a/mlir/test/Target/LLVMIR/omptarget-update-non-contiguous.mlir b/mlir/test/Target/LLVMIR/omptarget-update-non-contiguous.mlir
new file mode 100644
index 00000000000000..b8a97d7ac988d7
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/omptarget-update-non-contiguous.mlir
@@ -0,0 +1,194 @@
+// RUN: mlir-translate -mlir-to-llvmir -split-input-file %s | FileCheck %s
+
+// Descriptor strides are bytes and need not describe contiguous storage, even
+// for the first dimension. Preserve all four source strides and add a final
+// byte dimension so that each selected element transfers exactly four bytes.
+module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
+ llvm.func @descriptor_update(%descriptor: !llvm.ptr, %stride0: i64, %stride1: i64, %stride2: i64, %stride3: i64) {
+ %c0 = llvm.mlir.constant(0 : i64) : i64
+ %c1 = llvm.mlir.constant(1 : i64) : i64
+ %c2 = llvm.mlir.constant(2 : i64) : i64
+ %c3 = llvm.mlir.constant(3 : i64) : i64
+ %c4 = llvm.mlir.constant(4 : i64) : i64
+ %c5 = llvm.mlir.constant(5 : i64) : i64
+ %b0 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c2 : i64) extent(%c3 : i64) stride(%stride0 : i64) start_idx(%c1 : i64) stride_in_bytes(true)
+ %b1 = omp.map.bounds lower_bound(%c2 : i64) upper_bound(%c4 : i64) extent(%c5 : i64) stride(%stride1 : i64) start_idx(%c1 : i64) stride_in_bytes(true)
+ %b2 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c1 : i64) extent(%c2 : i64) stride(%stride2 : i64) start_idx(%c1 : i64) stride_in_bytes(true)
+ %b3 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c1 : i64) extent(%c2 : i64) stride(%stride3 : i64) start_idx(%c1 : i64) stride_in_bytes(true)
+ %to = omp.map.info var_ptr(%descriptor : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<4 x array<3 x i64>>)>) map_clauses(to) capture(ByRef) var_ptr_ptr(%descriptor : !llvm.ptr, i32) bounds(%b0, %b1, %b2, %b3) -> !llvm.ptr
+ omp.target_update map_entries(%to : !llvm.ptr)
+ %from = omp.map.info var_ptr(%descriptor : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<4 x array<3 x i64>>)>) map_clauses(from) capture(ByRef) var_ptr_ptr(%descriptor : !llvm.ptr, i32) bounds(%b0, %b1, %b2, %b3) -> !llvm.ptr
+ omp.target_update map_entries(%from : !llvm.ptr)
+ llvm.return
+ }
+}
+
+// CHECK: @.offload_sizes = {{.*}}[1 x i64] [i64 5]
+// CHECK: @.offload_maptypes = {{.*}}[1 x i64] [i64 17592186044417]
+// CHECK: @.offload_sizes.1 = {{.*}}[1 x i64] [i64 5]
+// CHECK: @.offload_maptypes.2 = {{.*}}[1 x i64] [i64 17592186044418]
+// CHECK-LABEL: define void @descriptor_update(
+// CHECK-SAME: ptr %[[DESC:.*]], i64 %[[S0:.*]], i64 %[[S1:.*]], i64 %[[S2:.*]], i64 %[[S3:.*]])
+// CHECK: %[[FROM_DIMS:.*]] = alloca [5 x %struct.descriptor_dim.0]
+// CHECK: %[[TO_DIMS:.*]] = alloca [5 x %struct.descriptor_dim]
+// CHECK: %[[BAND_OFFSET:.*]] = mul i64 2, %[[S1]]
+// CHECK: %[[BASE:.*]] = load ptr, ptr %[[DESC]]
+// CHECK: %[[IS_NULL:.*]] = icmp eq ptr %[[BASE]], null
+// CHECK: %[[INNER_COUNT:.*]] = select i1 %[[IS_NULL]], i64 0, i64 3
+// CHECK: %[[BASE_ADDR:.*]] = getelementptr inbounds [1 x ptr], ptr %{{.*}}, i32 0, i32 0
+// CHECK-NEXT: store ptr %[[BASE]], ptr %[[BASE_ADDR]]
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[TO_DIMS]], i64 0, i64 0
+// CHECK: store i64 2,
+// CHECK: store i64 %[[S3]],
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[TO_DIMS]], i64 0, i64 1
+// CHECK: store i64 2,
+// CHECK: store i64 %[[S2]],
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[TO_DIMS]], i64 0, i64 2
+// CHECK: store i64 %[[BAND_OFFSET]],
+// CHECK: store i64 3,
+// CHECK: store i64 %[[S1]],
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[TO_DIMS]], i64 0, i64 3
+// CHECK: store i64 %[[INNER_COUNT]],
+// CHECK: store i64 %[[S0]],
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[TO_DIMS]], i64 0, i64 4
+// CHECK: store i64 0,
+// CHECK: store i64 4,
+// CHECK: store i64 1,
+// CHECK: store ptr %[[TO_DIMS]],
+// CHECK: call void @__tgt_target_data_update_mapper({{.*}}ptr @.offload_sizes, ptr @.offload_maptypes,
+// CHECK: %[[FROM_BAND_OFFSET:.*]] = mul i64 2, %[[S1]]
+// CHECK: %[[FROM_BASE:.*]] = load ptr, ptr %[[DESC]]
+// CHECK: store ptr %[[FROM_BASE]],
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim.0], ptr %[[FROM_DIMS]], i64 0, i64 2
+// CHECK: store i64 %[[FROM_BAND_OFFSET]],
+// CHECK: store i64 3,
+// CHECK: store i64 %[[S1]],
+// CHECK: store ptr %[[FROM_DIMS]],
+// CHECK: call void @__tgt_target_data_update_mapper({{.*}}ptr @.offload_sizes.1, ptr @.offload_maptypes.2,
+
+// -----
+
+// Static array types provide complete source extents even when the bounds
+// carry legacy unit strides. Reconstruct slice strides from those extents,
+// rather than the selected section lengths. The scalars before and after the
+// array retain their byte sizes and acquire no non-contiguous descriptors.
+// A null raw array pointer represents an absent optional array and must give
+// a zero first-dimension count, even though the section bounds are nonempty.
+module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
+ llvm.func @explicit_array_update(%before: !llvm.ptr, %array: !llvm.ptr, %after: !llvm.ptr) {
+ %c0 = llvm.mlir.constant(0 : i64) : i64
+ %c1 = llvm.mlir.constant(1 : i64) : i64
+ %c2 = llvm.mlir.constant(2 : i64) : i64
+ %c4 = llvm.mlir.constant(4 : i64) : i64
+ %b0 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c2 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %b1 = omp.map.bounds lower_bound(%c2 : i64) upper_bound(%c4 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %b2 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c1 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %b3 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c1 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %first = omp.map.info var_ptr(%before : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr
+ %section = omp.map.info var_ptr(%array : !llvm.ptr, !llvm.array<2 x array<2 x array<5 x array<3 x i32>>>>) map_clauses(from) capture(ByRef) bounds(%b0, %b1, %b2, %b3) -> !llvm.ptr
+ %last = omp.map.info var_ptr(%after : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr
+ omp.target_update map_entries(%first, %section, %last : !llvm.ptr, !llvm.ptr, !llvm.ptr)
+ llvm.return
+ }
+}
+
+// CHECK: @.offload_sizes = {{.*}}[3 x i64] [i64 4, i64 5, i64 4]
+// CHECK: @.offload_maptypes = {{.*}}[3 x i64] [i64 2, i64 17592186044418, i64 2]
+// CHECK-LABEL: define void @explicit_array_update(
+// CHECK-SAME: ptr %[[BEFORE:.*]], ptr %[[ARRAY:.*]], ptr %[[AFTER:.*]])
+// CHECK: %[[DIMS:.*]] = alloca [5 x %struct.descriptor_dim]
+// CHECK-NOT: alloca {{.*}}%struct.descriptor_dim
+// CHECK: %[[IS_NULL:.*]] = icmp eq ptr %[[ARRAY]], null
+// CHECK-NEXT: %[[INNER_COUNT:.*]] = select i1 %[[IS_NULL]], i64 0, i64 3
+// CHECK: store ptr %[[BEFORE]],
+// CHECK: store ptr %[[ARRAY]],
+// CHECK: store ptr %[[AFTER]],
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 0
+// CHECK: store i64 0,
+// CHECK: store i64 2,
+// CHECK: store i64 120,
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 1
+// CHECK: store i64 0,
+// CHECK: store i64 2,
+// CHECK: store i64 60,
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 2
+// CHECK: store i64 24,
+// CHECK: store i64 3,
+// CHECK: store i64 12,
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 3
+// CHECK: store i64 0,
+// CHECK: store i64 %[[INNER_COUNT]],
+// CHECK: store i64 4,
+// CHECK: getelementptr inbounds [5 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 4
+// CHECK: store i64 0,
+// CHECK: store i64 4,
+// CHECK: store i64 1,
+// CHECK: %[[SECTION_PTR:.*]] = getelementptr inbounds [3 x ptr], ptr %{{.*}}, i32 0, i32 1
+// CHECK-NEXT: store ptr %[[DIMS]], ptr %[[SECTION_PTR]]
+// CHECK: call void @__tgt_target_data_update_mapper({{.*}}i64 -1, i32 3, {{.*}}ptr @.offload_sizes, ptr @.offload_maptypes,
+
+// -----
+
+// A scalar subscript has extent one, but still retains the physical stride of
+// the original array. The next stride includes the full source extent.
+module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
+ llvm.func @scalar_subscript_update(%array: !llvm.ptr, %extent0: i64, %extent1: i64) {
+ %c0 = llvm.mlir.constant(0 : i64) : i64
+ %c1 = llvm.mlir.constant(1 : i64) : i64
+ %c2 = llvm.mlir.constant(2 : i64) : i64
+ %b0 = omp.map.bounds lower_bound(%c1 : i64) upper_bound(%c1 : i64) extent(%c1 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %b1 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c1 : i64) extent(%extent1 : i64) stride(%extent0 : i64) start_idx(%c1 : i64)
+ %section = omp.map.info var_ptr(%array : !llvm.ptr, i32) map_clauses(to) capture(ByRef) bounds(%b0, %b1) -> !llvm.ptr
+ omp.target_update map_entries(%section : !llvm.ptr)
+ llvm.return
+ }
+}
+
+// CHECK: @.offload_sizes = {{.*}}[1 x i64] [i64 3]
+// CHECK: @.offload_maptypes = {{.*}}[1 x i64] [i64 17592186044417]
+// CHECK-LABEL: define void @scalar_subscript_update(
+// CHECK-SAME: ptr %[[ARRAY:.*]], i64 %[[EXTENT:.*]], i64 %{{.*}})
+// CHECK: %[[DIMS:.*]] = alloca [3 x %struct.descriptor_dim]
+// CHECK: %[[OUTER_STRIDE:.*]] = mul i64 %[[EXTENT]], 4
+// CHECK: %[[IS_NULL:.*]] = icmp eq ptr %[[ARRAY]], null
+// CHECK-NEXT: %[[INNER_COUNT:.*]] = select i1 %[[IS_NULL]], i64 0, i64 1
+// CHECK: getelementptr inbounds [3 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 0
+// CHECK: store i64 2,
+// CHECK: store i64 %[[OUTER_STRIDE]],
+// CHECK: getelementptr inbounds [3 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 1
+// CHECK: store i64 4,
+// CHECK: store i64 %[[INNER_COUNT]],
+// CHECK: store i64 4,
+// CHECK: getelementptr inbounds [3 x %struct.descriptor_dim], ptr %[[DIMS]], i64 0, i64 2
+// CHECK: store i64 0,
+// CHECK: store i64 4,
+// CHECK: store i64 1,
+// CHECK: call void @__tgt_target_data_update_mapper
+
+// -----
+
+// A full innermost dimension followed by a section of the outermost dimension
+// is contiguous. Keep the usual byte size and the selected starting address.
+module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
+ llvm.func @contiguous_update(%array: !llvm.ptr) {
+ %c0 = llvm.mlir.constant(0 : i64) : i64
+ %c1 = llvm.mlir.constant(1 : i64) : i64
+ %c2 = llvm.mlir.constant(2 : i64) : i64
+ %c3 = llvm.mlir.constant(3 : i64) : i64
+ %b0 = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c2 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %b1 = omp.map.bounds lower_bound(%c1 : i64) upper_bound(%c2 : i64) stride(%c3 : i64) start_idx(%c1 : i64)
+ %section = omp.map.info var_ptr(%array : !llvm.ptr, !llvm.array<5 x array<3 x i32>>) map_clauses(to) capture(ByRef) bounds(%b0, %b1) -> !llvm.ptr
+ omp.target_update map_entries(%section : !llvm.ptr)
+ llvm.return
+ }
+}
+
+// CHECK: @.offload_sizes = {{.*}}[1 x i64] [i64 24]
+// CHECK: @.offload_maptypes = {{.*}}[1 x i64] [i64 1]
+// CHECK-LABEL: define void @contiguous_update(
+// CHECK-SAME: ptr %[[ARRAY:.*]])
+// CHECK-NOT: %struct.descriptor_dim
+// CHECK: %[[START:.*]] = getelementptr inbounds [5 x [3 x i32]], ptr %[[ARRAY]], i64 0, i64 1, i64 0
+// CHECK: store ptr %[[START]],
+// CHECK: call void @__tgt_target_data_update_mapper({{.*}}ptr @.offload_sizes, ptr @.offload_maptypes,
+// CHECK: ret void
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 61f14a39ec8905..d694d6859bef07 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1343,6 +1343,9 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
void *HstPtrBegin, int64_t ArgSize,
int64_t ArgType, AsyncInfoTy &AsyncInfo) {
+ // An empty update requires neither a presence check nor a data transfer.
+ if (ArgSize == 0)
+ return OFFLOAD_SUCCESS;
TargetPointerResultTy TPR = Device.getMappingInfo().getTgtPtrBegin(
HstPtrBegin, ArgSize, /*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, /*MustContain=*/true);
@@ -1462,20 +1465,21 @@ static int targetDataNonContiguous(ident_t *Loc, DeviceTy &Device,
AsyncInfoTy &AsyncInfo) {
int Ret = OFFLOAD_SUCCESS;
if (CurrentDim < DimSize) {
- for (unsigned int I = 0; I < NonContig[CurrentDim].Count; ++I) {
+ // The last dimension describes an already contiguous piece. Visit its
+ // first element once instead of iterating over all of its elements.
+ uint64_t Count = NonContig[CurrentDim].Count;
+ if (CurrentDim == DimSize - 1)
+ Count = std::min(Count, uint64_t(1));
+ for (uint64_t I = 0; I < Count; ++I) {
uint64_t CurOffset =
NonContig[CurrentDim].Offset + I * NonContig[CurrentDim].Stride;
- // we only need to transfer the first element for the last dimension
- // since we've already got a contiguous piece.
- if (CurrentDim != DimSize - 1 || I == 0) {
- Ret = targetDataNonContiguous(Loc, Device, ArgsBase, NonContig, Size,
- ArgType, CurrentDim + 1, DimSize,
- Offset + CurOffset, AsyncInfo);
- // Stop the whole process if any contiguous piece returns anything
- // other than OFFLOAD_SUCCESS.
- if (Ret != OFFLOAD_SUCCESS)
- return Ret;
- }
+ Ret = targetDataNonContiguous(Loc, Device, ArgsBase, NonContig, Size,
+ ArgType, CurrentDim + 1, DimSize,
+ Offset + CurOffset, AsyncInfo);
+ // Stop the whole process if any contiguous piece returns anything
+ // other than OFFLOAD_SUCCESS.
+ if (Ret != OFFLOAD_SUCCESS)
+ return Ret;
}
} else {
void *Ptr = reinterpret_cast<void *>((char *)ArgsBase + Offset);
@@ -1490,9 +1494,12 @@ static int targetDataNonContiguous(ident_t *Loc, DeviceTy &Device,
static int getNonContigMergedDimension(__tgt_target_non_contig *NonContig,
int32_t DimSize) {
int RemovedDim = 0;
+ // Only consecutive innermost dimensions form a contiguous transfer.
for (int I = DimSize - 1; I > 0; --I) {
if (NonContig[I].Count * NonContig[I].Stride == NonContig[I - 1].Stride)
RemovedDim++;
+ else
+ break;
}
return RemovedDim;
}
@@ -1536,10 +1543,16 @@ int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
__tgt_target_non_contig *NonContig = (__tgt_target_non_contig *)Args[I];
int32_t DimSize = ArgSizes[I];
ODBG(ODT_DataTransfer) << "Non contig descriptor:";
- for (int I = 0; I < DimSize; I++)
+ bool IsEmpty = false;
+ for (int I = 0; I < DimSize; I++) {
ODBG(ODT_DataTransfer)
<< " Dim " << I << ": Offset " << NonContig[I].Offset << " Count "
<< NonContig[I].Count << " Stride " << NonContig[I].Stride;
+ IsEmpty |= NonContig[I].Count == 0;
+ }
+ // Empty sections require neither a presence check nor a data transfer.
+ if (IsEmpty)
+ continue;
int32_t MergedDim = getNonContigMergedDimension(NonContig, DimSize);
ODBG(ODT_DataTransfer) << "Merged " << MergedDim << " dimensions";
__tgt_target_non_contig &FirstMergedDim =
diff --git a/offload/test/offloading/fortran/dtype-char-array-map-2.f90 b/offload/test/offloading/fortran/dtype-char-array-map-2.f90
index f17ea9e53853b8..a0748dbe3f40d6 100644
--- a/offload/test/offloading/fortran/dtype-char-array-map-2.f90
+++ b/offload/test/offloading/fortran/dtype-char-array-map-2.f90
@@ -1,5 +1,4 @@
-! Offloading test that verifies certain type of character string arrays
-! map to and from device without problem.
+! Check whole-array and noncontiguous updates of a character-array component.
! REQUIRES: flang, amdgpu
! RUN: %libomptarget-compile-fortran-run-and-check-generic
@@ -9,17 +8,50 @@ program main
CHARACTER(LEN=16), dimension(10,10) :: char_arr
end type char_t
type(char_t) :: dtype_char
+ character(len=16) :: actual(10,10), expected(10,10), host_values(10,10)
+ character(len=16), parameter :: host_sentinel = 'host sentinel '
+ character(len=16), parameter :: device_sentinel = 'device sentinel '
+ integer :: i, j
+ dtype_char%char_arr = host_sentinel
!$omp target enter data map(alloc:dtype_char%char_arr)
-!$omp target
- dtype_char%char_arr(2,2) = 'c'
+ ! Check the entire string, including embedded and trailing blanks.
+!$omp target map(present,alloc:dtype_char%char_arr)
+ dtype_char%char_arr = 'whole from GPU '
!$omp end target
-
!$omp target update from(dtype_char%char_arr)
+ if (any(dtype_char%char_arr /= 'whole from GPU ')) stop 1
+
+ do j = 1, 10
+ do i = 1, 10
+ write(host_values(i,j), '(i2.2,a,i2.2,a)') i, '/', j, ': ab cd xy '
+ end do
+ end do
+ dtype_char%char_arr = host_sentinel
+ dtype_char%char_arr(3:6,:) = host_values(3:6,:)
+!$omp target map(present,alloc:dtype_char%char_arr)
+ dtype_char%char_arr = device_sentinel
+!$omp end target
+
+ ! Four complete strings per column, separated by six unselected strings.
+!$omp target update to(dtype_char%char_arr(3:6,:))
+!$omp target map(present,alloc:dtype_char%char_arr) map(from:actual)
+ actual = dtype_char%char_arr
+ dtype_char%char_arr = 'from GPU value '
+!$omp end target
+ expected = device_sentinel
+ expected(3:6,:) = host_values(3:6,:)
+ if (any(actual /= expected)) stop 2
+ dtype_char%char_arr = host_sentinel
+!$omp target update from(dtype_char%char_arr(3:6,:))
+ expected = host_sentinel
+ expected(3:6,:) = 'from GPU value '
+ if (any(dtype_char%char_arr /= expected)) stop 3
+!$omp target exit data map(delete:dtype_char%char_arr)
- print *, dtype_char%char_arr(2,2)
+ print *, 'PASS'
end program
-!CHECK: c
+! CHECK: PASS
diff --git a/offload/test/offloading/fortran/target-update-empty.f90 b/offload/test/offloading/fortran/target-update-empty.f90
new file mode 100644
index 00000000000000..41042ec09cad79
--- /dev/null
+++ b/offload/test/offloading/fortran/target-update-empty.f90
@@ -0,0 +1,52 @@
+! REQUIRES: flang, amdgpu
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+
+! Empty dimensions must remain empty when contiguous dimensions are merged.
+program target_update_empty
+ implicit none
+ integer, allocatable :: a(:,:)
+ integer :: storage(4,3), actual(4,3), n, lo, hi
+ integer, parameter :: host_sentinel = 7777, device_sentinel = -9999
+
+ n = 0
+ allocate(a(n,3))
+!$omp target enter data map(to:a)
+!$omp target update to(present:a)
+!$omp target update from(present:a)
+!$omp target update to(present:a(:,2:))
+!$omp target update from(present:a(:,2:))
+!$omp target exit data map(delete:a)
+ deallocate(a)
+
+ allocate(a(4,n))
+!$omp target enter data map(to:a)
+!$omp target update to(present:a)
+!$omp target update from(present:a)
+!$omp target update to(present:a(2:3,:))
+!$omp target update from(present:a(2:3,:))
+!$omp target exit data map(delete:a)
+ deallocate(a)
+
+ ! Empty sections of nonempty arrays must preserve every surrounding element.
+ storage = host_sentinel
+!$omp target enter data map(alloc:storage)
+!$omp target map(present,alloc:storage)
+ storage = device_sentinel
+!$omp end target
+ lo = 3
+ hi = 2
+!$omp target update to(present:storage(lo:hi,:))
+!$omp target update to(present:storage(:,lo:hi))
+!$omp target map(present,alloc:storage) map(from:actual)
+ actual = storage
+ storage = -42
+!$omp end target
+ if (any(actual /= device_sentinel)) stop 1
+!$omp target update from(present:storage(lo:hi,:))
+!$omp target update from(present:storage(:,lo:hi))
+ if (any(storage /= host_sentinel)) stop 2
+!$omp target exit data map(delete:storage)
+ print *, 'PASS'
+end program
+
+! CHECK: PASS
diff --git a/offload/test/offloading/fortran/target-update-noncontiguous.f90 b/offload/test/offloading/fortran/target-update-noncontiguous.f90
new file mode 100644
index 00000000000000..f9e2e134e8c0da
--- /dev/null
+++ b/offload/test/offloading/fortran/target-update-noncontiguous.f90
@@ -0,0 +1,137 @@
+! REQUIRES: flang, amdgpu
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+
+! A section of the second dimension has gaps between the (k, s) slices.
+! Check both the values transferred and the values outside the section.
+program target_update_noncontiguous
+ implicit none
+ type :: waves
+ integer, pointer :: a(:,:,:,:)
+ end type
+ type(waves), target :: w
+ integer :: actual(0:2, 5, -1:0, 2:3)
+ integer :: values(0:2, 5, -1:0, 2:3)
+ integer :: expected(0:2, 5, -1:0, 2:3)
+ integer :: plain(-2:0, 5, 2, 2), plain_actual(-2:0, 5, 2, 2)
+ integer :: i, b, k, s, nb
+ integer, parameter :: host_sentinel = 7777, device_sentinel = -9999
+
+ nb = 2
+ allocate(w%a(0:2, 5, -1:0, 2:3))
+ do s = 2, 3
+ do k = -1, 0
+ do b = 1, 5
+ do i = 0, 2
+ values(i,b,k,s) = i + 10*b + 100*k + 1000*s
+ end do
+ end do
+ end do
+ end do
+ w%a = values
+!$omp target enter data map(to:w)
+!$omp target enter data map(alloc:w%a)
+!$omp target map(present,to:w)
+ w%a = device_sentinel
+!$omp end target
+
+!$omp target update to(w%a(:,nb+1:,:,:))
+ ! Read through a kernel: a whole-array update would hide the original bug.
+!$omp target map(present,to:w) map(from:actual)
+ do s = 2, 3
+ do k = -1, 0
+ do b = 1, 5
+ do i = 0, 2
+ actual(i,b,k,s) = w%a(i,b,k,s)
+ w%a(i,b,k,s) = -(i + 10*b + 100*k + 1000*s)
+ end do
+ end do
+ end do
+ end do
+!$omp end target
+ expected = device_sentinel
+ expected(:,nb+1:,:,:) = values(:,nb+1:,:,:)
+ if (any(actual /= expected)) stop 1
+
+ w%a = host_sentinel
+!$omp target update from(w%a(:,nb+1:,:,:))
+ expected = host_sentinel
+ expected(:,nb+1:,:,:) = -values(:,nb+1:,:,:)
+ if (any(w%a /= expected)) stop 2
+!$omp target exit data map(delete:w%a)
+!$omp target exit data map(delete:w)
+ deallocate(w%a)
+
+ ! Explicit-shape arrays also require gaps between the outer slices. Keep
+ ! a nondefault lower bound in the innermost dimension.
+ plain = 42
+!$omp target enter data map(alloc:plain)
+!$omp target map(present,to:plain)
+ plain = device_sentinel
+!$omp end target
+!$omp target update to(plain(:,nb+1:,:,:))
+!$omp target map(present,to:plain) map(from:plain_actual)
+ plain_actual = plain
+ plain = -42
+!$omp end target
+ do s = 1, 2
+ do k = 1, 2
+ do b = 1, 5
+ do i = -2, 0
+ if (b > nb) then
+ if (plain_actual(i,b,k,s) /= 42) stop 3
+ else
+ if (plain_actual(i,b,k,s) /= device_sentinel) stop 4
+ end if
+ end do
+ end do
+ end do
+ end do
+ plain = host_sentinel
+!$omp target update from(plain(:,nb+1:,:,:))
+ do s = 1, 2
+ do k = 1, 2
+ do b = 1, 5
+ do i = -2, 0
+ if (b > nb) then
+ if (plain(i,b,k,s) /= -42) stop 5
+ else
+ if (plain(i,b,k,s) /= host_sentinel) stop 6
+ end if
+ end do
+ end do
+ end do
+ end do
+ ! A scalar subscript selects one element but retains the physical first
+ ! dimension extent when calculating the stride of subsequent dimensions.
+ plain = 84
+!$omp target map(present,to:plain)
+ plain = device_sentinel
+!$omp end target
+!$omp target update to(plain(0,nb+1:,:,:))
+!$omp target map(present,to:plain) map(from:plain_actual)
+ plain_actual = plain
+ plain = -84
+!$omp end target
+ if (any(plain_actual(0,nb+1:,:,:) /= 84)) stop 7
+ plain_actual(0,nb+1:,:,:) = device_sentinel
+ if (any(plain_actual /= device_sentinel)) stop 8
+
+ plain = host_sentinel
+!$omp target update from(plain(0,nb+1:,:,:))
+ if (any(plain(0,nb+1:,:,:) /= -84)) stop 9
+ plain(0,nb+1:,:,:) = host_sentinel
+ if (any(plain /= host_sentinel)) stop 10
+!$omp target exit data map(delete:plain)
+ call update_absent()
+ print *, "PASS"
+contains
+ subroutine update_absent(a)
+ integer, optional :: a(4,5)
+ ! An absent explicit-shape dummy has a null raw base rather than a null
+ ! descriptor. PRESENT must not request a mapping for this empty update.
+!$omp target update to(present:a(2:3,:))
+!$omp target update from(present:a(2:3,:))
+ end subroutine
+end program
+
+! CHECK: PASS
diff --git a/offload/test/offloading/fortran/target-update-wide-character.f90 b/offload/test/offloading/fortran/target-update-wide-character.f90
new file mode 100644
index 00000000000000..b3a031316f6b5d
--- /dev/null
+++ b/offload/test/offloading/fortran/target-update-wide-character.f90
@@ -0,0 +1,74 @@
+! REQUIRES: flang, amdgpu
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+
+! Assumed-length character bounds count code units, not bytes. Check every
+! code unit on the device after TO and on the host after FROM.
+program target_update_wide_character
+ implicit none
+ character(kind=2,len=8) :: a2
+ character(kind=4,len=8) :: a4
+ call check_kind2(a2)
+ call check_kind4(a4)
+ call absent_kind4()
+ print *, 'PASS'
+contains
+ subroutine check_kind2(a)
+ character(kind=2,len=*) :: a
+ character(kind=2,len=8), parameter :: to_value = 2_'A B C '
+ character(kind=2,len=8), parameter :: from_value = 2_'X Y Z '
+ integer :: actual(8), i
+!$omp target enter data map(alloc:a)
+!$omp target map(present,alloc:a)
+ a = 2_'device '
+!$omp end target
+ a = to_value
+!$omp target update to(a)
+!$omp target map(present,alloc:a) map(from:actual)
+ do i = 1, 8
+ actual(i) = iachar(a(i:i))
+ end do
+ a = from_value
+!$omp end target
+ do i = 1, 8
+ if (actual(i) /= iachar(to_value(i:i))) stop 1
+ end do
+ a = 2_'hostdata'
+!$omp target update from(a)
+ if (a /= from_value) stop 2
+!$omp target exit data map(delete:a)
+ end subroutine
+
+ subroutine check_kind4(a)
+ character(kind=4,len=*) :: a
+ character(kind=4,len=8), parameter :: to_value = 4_'A B C '
+ character(kind=4,len=8), parameter :: from_value = 4_'X Y Z '
+ integer :: actual(8), i
+!$omp target enter data map(alloc:a)
+!$omp target map(present,alloc:a)
+ a = 4_'device '
+!$omp end target
+ a = to_value
+!$omp target update to(a)
+!$omp target map(present,alloc:a) map(from:actual)
+ do i = 1, 8
+ actual(i) = iachar(a(i:i))
+ end do
+ a = from_value
+!$omp end target
+ do i = 1, 8
+ if (actual(i) /= iachar(to_value(i:i))) stop 3
+ end do
+ a = 4_'hostdata'
+!$omp target update from(a)
+ if (a /= from_value) stop 4
+!$omp target exit data map(delete:a)
+ end subroutine
+
+ subroutine absent_kind4(a)
+ character(kind=4,len=*), optional :: a
+!$omp target update to(present:a)
+!$omp target update from(present:a)
+ end subroutine
+end program
+
+! CHECK: PASS
diff --git a/offload/test/offloading/non_contiguous_update.cpp b/offload/test/offloading/non_contiguous_update.cpp
index 011e25a4a3b59b..632f3bac693e5a 100644
--- a/offload/test/offloading/non_contiguous_update.cpp
+++ b/offload/test/offloading/non_contiguous_update.cpp
@@ -130,5 +130,29 @@ int main() {
// DEBUG: offset 80 len 12
}
+ {
+ // case 4: a contiguous pair of outer dimensions must not be merged across
+ // a gap in an inner dimension. This describes the Fortran array section
+ // arr(:,3:5,:,:) for integer(4) :: arr(4,5,3,2).
+ __tgt_target_non_contig non_contig[5] = {
+ {0, 2, 240}, {0, 3, 80}, {32, 3, 16}, {0, 4, 4}, {0, 4, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
+
+ int32_t array[2][3][5][4];
+ void *base = array;
+ void *begin = &non_contig;
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ &size, &type);
+ // Each (k-point, spin) slice has a separate contiguous band section.
+ // DEBUG: Merged 2 dimensions
+ // DEBUG: offset 32 len 48
+ // DEBUG: offset 112 len 48
+ // DEBUG: offset 192 len 48
+ // DEBUG: offset 272 len 48
+ // DEBUG: offset 352 len 48
+ // DEBUG: offset 432 len 48
+ }
+
return 0;
}
diff --git a/offload/test/offloading/non_contiguous_update_empty.cpp b/offload/test/offloading/non_contiguous_update_empty.cpp
new file mode 100644
index 00000000000000..7f029b5b802d6f
--- /dev/null
+++ b/offload/test/offloading/non_contiguous_update_empty.cpp
@@ -0,0 +1,107 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <cassert>
+#include <climits>
+#include <cstdint>
+#include <cstdio>
+
+// Data structure and flags from the OpenMP runtime ABI.
+struct Dim {
+ uint64_t Offset;
+ uint64_t Count;
+ uint64_t Stride;
+};
+
+constexpr int64_t To = 0x1;
+constexpr int64_t From = 0x2;
+constexpr int64_t Present = 0x1000;
+constexpr int64_t NonContiguous = 0x100000000000;
+
+extern "C" void __tgt_target_data_update(int64_t Device, int32_t NumArgs,
+ void **Bases, void **Begins,
+ int64_t *Sizes,
+ int64_t *Types) noexcept;
+
+template <unsigned N>
+static void update(void *Base, Dim (&Dims)[N], int64_t Flags) noexcept {
+ void *Begin = Dims;
+ int64_t Size = N;
+ Flags |= NonContiguous;
+ __tgt_target_data_update(-1, 1, &Base, &Begin, &Size, &Flags);
+}
+
+static void emptyContiguousUpdate(void *Base, int64_t Flags) noexcept {
+ void *Begin = Base;
+ int64_t Size = 0;
+ __tgt_target_data_update(-1, 1, &Base, &Begin, &Size, &Flags);
+}
+
+int main() {
+ // Fortran integer(4) array(4,5,3,2), selecting array(:,3:5,:,:).
+ Dim Section[] = {{0, 2, 240}, {0, 3, 80}, {32, 3, 16}, {0, 4, 4}, {0, 4, 1}};
+ // A zero innermost extent also makes all outer physical strides zero.
+ // Merging these dimensions must not result in a zero-byte runtime copy.
+ Dim EmptyInner[] = {{0, 2, 0}, {0, 3, 0}, {0, 5, 0}, {0, 0, 4}, {0, 4, 1}};
+ Dim EmptyOuter[] = {
+ {0, 0, 240}, {0, 3, 80}, {32, 3, 16}, {0, 4, 4}, {0, 4, 1}};
+
+ int Data[120];
+ for (int &Value : Data)
+ Value = -1;
+#pragma omp target data map(to : Data)
+ {
+ for (int &Value : Data)
+ Value = 7;
+ emptyContiguousUpdate(Data, To);
+ emptyContiguousUpdate(Data, To | Present);
+ update(Data, EmptyInner, To);
+ update(Data, EmptyOuter, To);
+#pragma omp target update from(Data)
+ for (int Value : Data)
+ assert(Value == -1);
+
+ for (int &Value : Data)
+ Value = 7;
+ update(Data, Section, To);
+#pragma omp target update from(Data)
+ for (int I = 0; I < 120; ++I)
+ assert(Data[I] == (I % 20 >= 8 ? 7 : -1));
+
+ for (int I = 0; I < 120; ++I)
+ Data[I] = I + 100;
+#pragma omp target update to(Data)
+ for (int &Value : Data)
+ Value = -9;
+ emptyContiguousUpdate(Data, From);
+ emptyContiguousUpdate(Data, From | Present);
+ update(Data, EmptyInner, From);
+ update(Data, EmptyOuter, From);
+ for (int Value : Data)
+ assert(Value == -9);
+ update(Data, Section, From);
+ for (int I = 0; I < 120; ++I)
+ assert(Data[I] == (I % 20 >= 8 ? I + 100 : -9));
+ }
+
+ // Empty sections must not look up an unmapped address, even with PRESENT.
+ static int Unmapped;
+ emptyContiguousUpdate(&Unmapped, To);
+ emptyContiguousUpdate(&Unmapped, From);
+ emptyContiguousUpdate(&Unmapped, To | Present);
+ emptyContiguousUpdate(&Unmapped, From | Present);
+ update(&Unmapped, EmptyInner, To | Present);
+ update(&Unmapped, EmptyOuter, From | Present);
+
+ // An already contiguous final dimension needs one visit regardless of its
+ // count. No TO/FROM flag or allocation is needed to check the traversal.
+ Dim Large[] = {{0, uint64_t(UINT_MAX) + 1, 1}};
+ update(&Unmapped, Large, 0);
+
+ // Skip an empty section before traversing a potentially enormous outer
+ // dimension. This also checks that an inner zero suppresses PRESENT checks.
+ Dim LargeEmpty[] = {{0, uint64_t(UINT_MAX) + 1, 8}, {0, 0, 4}, {0, 4, 1}};
+ update(&Unmapped, LargeEmpty, To | Present);
+
+ std::puts("PASS");
+ // CHECK: PASS
+}
>From 8c7dea2eebeebed11cdcc6c60606162f597f2de3 Mon Sep 17 00:00:00 2001
From: Akash Banerjee <Akash.Banerjee at amd.com>
Date: Fri, 23 Oct 2026 18:44:52 +0100
Subject: [PATCH 2/2] Preserve non-contiguous update descriptors in deferred
tasks.
---
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 9 ++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 44 +++++++--
...mptarget-update-nowait-non-contiguous.mlir | 97 +++++++++++++++++++
.../fortran/dtype-char-array-map-2.f90 | 3 +-
.../fortran/target-update-empty.f90 | 3 +-
.../target-update-noncontiguous-nowait.f90 | 77 +++++++++++++++
.../fortran/target-update-noncontiguous.f90 | 3 +-
.../fortran/target-update-wide-character.f90 | 1 +
.../non_contiguous_update_empty.cpp | 1 +
9 files changed, 229 insertions(+), 9 deletions(-)
create mode 100644 mlir/test/Target/LLVMIR/omptarget-update-nowait-non-contiguous.mlir
create mode 100644 offload/test/offloading/fortran/target-update-noncontiguous-nowait.f90
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index a1ed0d2b065759..e908d2027dbc2b 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2827,6 +2827,15 @@ class OpenMPIRBuilder {
/// runtime library for debugging
Value *MapNamesArray = nullptr;
+ struct NonContiguousDescriptor {
+ /// The entry in PointersArray that refers to this descriptor.
+ unsigned MapIndex;
+ AllocaInst *Storage;
+ };
+ /// Descriptors referenced by PointersArray need their own task-private
+ /// copies when the runtime call is deferred.
+ SmallVector<NonContiguousDescriptor, 2> NonContiguousDescriptors;
+
explicit TargetDataRTArgs() = default;
explicit TargetDataRTArgs(Value *BasePointersArray, Value *PointersArray,
Value *SizesArray, Value *MapTypesArray,
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index ff35966ccfe3ab..0a912a23414d76 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9691,9 +9691,10 @@ static Type *getOffloadingArrayType(Value *V) {
// %struct.kmp_task_ompbuilder_t is returned.
static StructType *
createTaskWithPrivatesTy(OpenMPIRBuilder &OMPIRBuilder,
- ArrayRef<Value *> OffloadingArraysToPrivatize) {
+ ArrayRef<Value *> OffloadingArraysToPrivatize,
+ ArrayRef<Type *> DescriptorTypes) {
- if (OffloadingArraysToPrivatize.empty())
+ if (OffloadingArraysToPrivatize.empty() && DescriptorTypes.empty())
return OMPIRBuilder.Task;
SmallVector<Type *, 4> StructFieldTypes;
@@ -9705,6 +9706,7 @@ createTaskWithPrivatesTy(OpenMPIRBuilder &OMPIRBuilder,
assert(ArrayTy && "ArrayType cannot be nullptr");
StructFieldTypes.push_back(ArrayTy);
}
+ llvm::append_range(StructFieldTypes, DescriptorTypes);
StructType *PrivatesStructTy =
StructType::create(StructFieldTypes, "struct.privates");
return StructType::create({OMPIRBuilder.Task, PrivatesStructTy},
@@ -9908,9 +9910,14 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::emitTargetTask(
}
}
}
+ auto NonContiguousDescriptors = RTArgs.NonContiguousDescriptors;
+ if (!NeedsTargetTask)
+ NonContiguousDescriptors.clear();
OI->PostOutlineCB = [this, ToBeDeleted, Dependencies, NeedsTargetTask,
- DeviceID, OffloadingArraysToPrivatize](
- Function &OutlinedFn) mutable {
+ DeviceID, OffloadingArraysToPrivatize,
+ NonContiguousDescriptors,
+ PointersArray =
+ RTArgs.PointersArray](Function &OutlinedFn) mutable {
assert(OutlinedFn.hasOneUse() &&
"there must be a single user for the outlined function");
@@ -9934,8 +9941,13 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::emitTargetTask(
int SharedArgOperandNo =
HasShareds ? OffloadingArraysToPrivatize.size() + 1 : 0;
- StructType *TaskWithPrivatesTy =
- createTaskWithPrivatesTy(*this, OffloadingArraysToPrivatize);
+ // Descriptor storage is not an argument to the outlined function. It is
+ // reached indirectly through the task-private pointer array.
+ SmallVector<Type *, 2> DescriptorTypes;
+ for (const auto &Descriptor : NonContiguousDescriptors)
+ DescriptorTypes.push_back(Descriptor.Storage->getAllocatedType());
+ StructType *TaskWithPrivatesTy = createTaskWithPrivatesTy(
+ *this, OffloadingArraysToPrivatize, DescriptorTypes);
StructType *PrivatesTy = nullptr;
if (!OffloadingArraysToPrivatize.empty())
@@ -10034,6 +10046,7 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::emitTargetTask(
if (!OffloadingArraysToPrivatize.empty()) {
Value *Privates =
Builder.CreateStructGEP(TaskWithPrivatesTy, TaskData, 1);
+ Value *PrivatePointersArray = nullptr;
for (unsigned int i = 0; i < OffloadingArraysToPrivatize.size(); ++i) {
Value *PtrToPrivatize = OffloadingArraysToPrivatize[i];
[[maybe_unused]] Type *ArrayType =
@@ -10049,6 +10062,23 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::emitTargetTask(
Builder.CreateMemCpy(
Dst, Alignment, PtrToPrivatize, Alignment,
Builder.getInt64(M.getDataLayout().getTypeStoreSize(ElementType)));
+ if (PtrToPrivatize == PointersArray)
+ PrivatePointersArray = Dst;
+ }
+ // Snapshot the encounter-time bounds before enqueueing the task, then
+ // redirect its pointer entries to descriptors owned by that task.
+ for (auto [I, Descriptor] : llvm::enumerate(NonContiguousDescriptors)) {
+ assert(PrivatePointersArray && "missing task-private pointer array");
+ Value *Dst = Builder.CreateStructGEP(
+ PrivatesTy, Privates, OffloadingArraysToPrivatize.size() + I);
+ Builder.CreateMemCpy(
+ Dst, Alignment, Descriptor.Storage, Descriptor.Storage->getAlign(),
+ Builder.getInt64(M.getDataLayout().getTypeStoreSize(
+ Descriptor.Storage->getAllocatedType())));
+ Value *Entry = Builder.CreateConstInBoundsGEP2_32(
+ getOffloadingArrayType(PointersArray), PrivatePointersArray, 0,
+ Descriptor.MapIndex);
+ Builder.CreateAlignedStore(Dst, Entry, Alignment);
}
}
@@ -10524,6 +10554,7 @@ void OpenMPIRBuilder::emitOffloadingArraysArgument(IRBuilderBase &Builder,
auto Int64Ty = Type::getInt64Ty(M.getContext());
auto Int64PtrTy = UnqualPtrTy;
+ RTArgs.NonContiguousDescriptors = Info.RTArgs.NonContiguousDescriptors;
if (!Info.NumberOfPtrs) {
RTArgs.BasePointersArray = ConstantPointerNull::get(VoidPtrPtrTy);
RTArgs.PointersArray = ConstantPointerNull::get(VoidPtrPtrTy);
@@ -10603,6 +10634,7 @@ void OpenMPIRBuilder::emitNonContiguousDescriptor(InsertPointTy AllocaIP,
ArrayType *ArrayTy = ArrayType::get(DimTy, NonContigInfo.Dims[I]);
AllocaInst *DimsAddr =
Builder.CreateAlloca(ArrayTy, /* ArraySize = */ nullptr, "dims");
+ Info.RTArgs.NonContiguousDescriptors.push_back({I, DimsAddr});
Builder.restoreIP(CodeGenIP);
for (unsigned II = 0, EE = NonContigInfo.Dims[I]; II < EE; ++II) {
unsigned RevIdx = EE - II - 1;
diff --git a/mlir/test/Target/LLVMIR/omptarget-update-nowait-non-contiguous.mlir b/mlir/test/Target/LLVMIR/omptarget-update-nowait-non-contiguous.mlir
new file mode 100644
index 00000000000000..cd09631eebbe26
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/omptarget-update-nowait-non-contiguous.mlir
@@ -0,0 +1,97 @@
+// RUN: mlir-translate -mlir-to-llvmir -split-input-file %s | FileCheck %s
+
+// This translates host IR and does not require an AMDGPU backend.
+// Deferred updates must own their non-contiguous descriptors as well as their
+// base and begin pointer arrays. Different descriptor ranks and intervening
+// contiguous entries check that the task rewrites the correct begin pointers.
+module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
+ llvm.func @deferred_update(%first: !llvm.ptr, %array: !llvm.ptr, %middle: !llvm.ptr, %strided: !llvm.ptr, %last: !llvm.ptr) {
+ %c0 = llvm.mlir.constant(0 : i64) : i64
+ %c1 = llvm.mlir.constant(1 : i64) : i64
+ %c3 = llvm.mlir.constant(3 : i64) : i64
+ %c4 = llvm.mlir.constant(4 : i64) : i64
+ %c16 = llvm.mlir.constant(16 : i64) : i64
+ %inner = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c1 : i64) stride(%c1 : i64) start_idx(%c1 : i64)
+ %outer = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c4 : i64) stride(%c3 : i64) start_idx(%c1 : i64)
+ %strided_bound = omp.map.bounds lower_bound(%c1 : i64) upper_bound(%c3 : i64) stride(%c16 : i64) start_idx(%c1 : i64) stride_in_bytes(true)
+ %map_first = omp.map.info var_ptr(%first : !llvm.ptr, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr
+ %map_array = omp.map.info var_ptr(%array : !llvm.ptr, !llvm.array<5 x array<3 x i32>>) map_clauses(to) capture(ByRef) bounds(%inner, %outer) -> !llvm.ptr
+ %map_middle = omp.map.info var_ptr(%middle : !llvm.ptr, i64) map_clauses(to) capture(ByRef) -> !llvm.ptr
+ %map_strided = omp.map.info var_ptr(%strided : !llvm.ptr, i64) map_clauses(from) capture(ByRef) bounds(%strided_bound) -> !llvm.ptr
+ %map_last = omp.map.info var_ptr(%last : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr
+ omp.target_update map_entries(%map_first, %map_array, %map_middle, %map_strided, %map_last : !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr) nowait
+ llvm.return
+ }
+}
+
+// CHECK: %struct.[[$TASK_TYPE:[^ ]+]] = type { %struct.kmp_task_ompbuilder_t, %struct.[[$PRIVATES:[^ ]+]] }
+// CHECK: %struct.[[$PRIVATES]] = type { [5 x ptr], [5 x ptr], [3 x %struct.descriptor_dim], [2 x %struct.descriptor_dim] }
+// CHECK: @.offload_sizes = {{.*}}[5 x i64] [i64 4, i64 3, i64 8, i64 2, i64 4]
+// CHECK: @.offload_maptypes = {{.*}}[5 x i64] [i64 1, i64 17592186044417, i64 1, i64 17592186044418, i64 2]
+// CHECK-LABEL: define void @deferred_update(
+// CHECK: %[[BASES:.*]] = alloca [5 x ptr]
+// CHECK: %[[POINTERS:.*]] = alloca [5 x ptr]
+// CHECK: %[[DESC0:.*]] = alloca [3 x %struct.descriptor_dim]
+// CHECK: %[[DESC1:.*]] = alloca [2 x %struct.descriptor_dim]
+// CHECK: store ptr %[[DESC1]],
+// CHECK: %[[BASES_BEGIN:.*]] = getelementptr inbounds [5 x ptr], ptr %[[BASES]], i32 0, i32 0
+// CHECK: %[[POINTERS_BEGIN:.*]] = getelementptr inbounds [5 x ptr], ptr %[[POINTERS]], i32 0, i32 0
+// CHECK: %[[TASK:.*]] = call ptr @__kmpc_omp_target_task_alloc({{.*}}i64 240, i64 0, ptr @[[PROXY:[^,]+]], i64 -1)
+// CHECK: %[[PRIVATE_DATA:.*]] = getelementptr inbounds nuw %struct.[[$TASK_TYPE]], ptr %[[TASK]], i32 0, i32 1
+// CHECK: %[[PRIVATE_BASES:.*]] = getelementptr inbounds nuw %struct.[[$PRIVATES]], ptr %[[PRIVATE_DATA]], i32 0, i32 0
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 1 %[[PRIVATE_BASES]], ptr align 1 %[[BASES_BEGIN]], i64 40, i1 false)
+// CHECK: %[[PRIVATE_POINTERS:.*]] = getelementptr inbounds nuw %struct.[[$PRIVATES]], ptr %[[PRIVATE_DATA]], i32 0, i32 1
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 1 %[[PRIVATE_POINTERS]], ptr align 1 %[[POINTERS_BEGIN]], i64 40, i1 false)
+// CHECK: %[[PRIVATE_DESC0:.*]] = getelementptr inbounds nuw %struct.[[$PRIVATES]], ptr %[[PRIVATE_DATA]], i32 0, i32 2
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 1 %[[PRIVATE_DESC0]], ptr align 8 %[[DESC0]], i64 72, i1 false)
+// CHECK: %[[ENTRY1:.*]] = getelementptr inbounds [5 x ptr], ptr %[[PRIVATE_POINTERS]], i32 0, i32 1
+// CHECK-NEXT: store ptr %[[PRIVATE_DESC0]], ptr %[[ENTRY1]]
+// CHECK: %[[PRIVATE_DESC1:.*]] = getelementptr inbounds nuw %struct.[[$PRIVATES]], ptr %[[PRIVATE_DATA]], i32 0, i32 3
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 1 %[[PRIVATE_DESC1]], ptr align 8 %[[DESC1]], i64 48, i1 false)
+// CHECK: %[[ENTRY3:.*]] = getelementptr inbounds [5 x ptr], ptr %[[PRIVATE_POINTERS]], i32 0, i32 3
+// CHECK-NEXT: store ptr %[[PRIVATE_DESC1]], ptr %[[ENTRY3]]
+// CHECK: call i32 @__kmpc_omp_task({{.*}}ptr %[[TASK]])
+// CHECK: ret void
+
+// CHECK: define internal void @[[WORKER:[^(]+]](i32 %{{.*}}, ptr %[[WORKER_BASES:.*]], ptr %[[WORKER_POINTERS:.*]]) {
+// CHECK: call void @__tgt_target_data_update_nowait_mapper({{.*}}i64 -1, i32 5, ptr %[[WORKER_BASES]], ptr %[[WORKER_POINTERS]], ptr @.offload_sizes, ptr @.offload_maptypes,
+// CHECK: ret void
+// CHECK: define internal void @[[PROXY]](i32 %[[TID:.*]], ptr %[[TASK_ARG:.*]]) {
+// CHECK: %[[PROXY_PRIVATE_DATA:.*]] = getelementptr inbounds nuw %struct.[[$TASK_TYPE]], ptr %[[TASK_ARG]], i32 0, i32 1
+// CHECK: %[[PROXY_BASES:.*]] = getelementptr inbounds nuw %struct.[[$PRIVATES]], ptr %[[PROXY_PRIVATE_DATA]], i32 0, i32 0
+// CHECK: %[[PROXY_POINTERS:.*]] = getelementptr inbounds nuw %struct.[[$PRIVATES]], ptr %[[PROXY_PRIVATE_DATA]], i32 0, i32 1
+// CHECK: call void @[[WORKER]](i32 %[[TID]], ptr %[[PROXY_BASES]], ptr %[[PROXY_POINTERS]])
+// CHECK: ret void
+
+// -----
+
+// Synchronous updates can consume the encounter-time descriptor directly.
+module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
+ llvm.func @synchronous_update(%strided: !llvm.ptr) {
+ %c0 = llvm.mlir.constant(0 : i64) : i64
+ %c1 = llvm.mlir.constant(1 : i64) : i64
+ %c3 = llvm.mlir.constant(3 : i64) : i64
+ %c16 = llvm.mlir.constant(16 : i64) : i64
+ %bound = omp.map.bounds lower_bound(%c0 : i64) upper_bound(%c3 : i64) stride(%c16 : i64) start_idx(%c1 : i64) stride_in_bytes(true)
+ %map = omp.map.info var_ptr(%strided : !llvm.ptr, i64) map_clauses(to) capture(ByRef) bounds(%bound) -> !llvm.ptr
+ omp.target_update map_entries(%map : !llvm.ptr)
+ llvm.return
+ }
+}
+
+// CHECK: @.offload_sizes = {{.*}}[1 x i64] [i64 2]
+// CHECK: @.offload_maptypes = {{.*}}[1 x i64] [i64 17592186044417]
+// CHECK-LABEL: define void @synchronous_update(
+// CHECK-NOT: @__kmpc
+// CHECK: alloca [1 x ptr]
+// CHECK: %[[SYNC_POINTERS:.*]] = alloca [1 x ptr]
+// CHECK: %[[SYNC_DESC:.*]] = alloca [2 x %struct.descriptor_dim]
+// CHECK-NOT: @__kmpc
+// CHECK-NOT: @llvm.memcpy
+// CHECK: getelementptr inbounds [2 x %struct.descriptor_dim], ptr %[[SYNC_DESC]], i64 0, i64 1
+// CHECK: %[[SYNC_ENTRY:.*]] = getelementptr inbounds [1 x ptr], ptr %[[SYNC_POINTERS]], i32 0, i32 0
+// CHECK-NEXT: store ptr %[[SYNC_DESC]], ptr %[[SYNC_ENTRY]]
+// CHECK-NOT: @__kmpc
+// CHECK-NOT: @llvm.memcpy
+// CHECK: call void @__tgt_target_data_update_mapper({{.*}}ptr @.offload_sizes, ptr @.offload_maptypes,
+// CHECK-NEXT: ret void
diff --git a/offload/test/offloading/fortran/dtype-char-array-map-2.f90 b/offload/test/offloading/fortran/dtype-char-array-map-2.f90
index a0748dbe3f40d6..c340491242a0f4 100644
--- a/offload/test/offloading/fortran/dtype-char-array-map-2.f90
+++ b/offload/test/offloading/fortran/dtype-char-array-map-2.f90
@@ -1,5 +1,6 @@
! Check whole-array and noncontiguous updates of a character-array component.
-! REQUIRES: flang, amdgpu
+! REQUIRES: flang, flang-rt, libc, amdgpu
+! UNSUPPORTED: apu
! RUN: %libomptarget-compile-fortran-run-and-check-generic
program main
diff --git a/offload/test/offloading/fortran/target-update-empty.f90 b/offload/test/offloading/fortran/target-update-empty.f90
index 41042ec09cad79..9694a3ded45f35 100644
--- a/offload/test/offloading/fortran/target-update-empty.f90
+++ b/offload/test/offloading/fortran/target-update-empty.f90
@@ -1,4 +1,5 @@
-! REQUIRES: flang, amdgpu
+! REQUIRES: flang, flang-rt, libc, amdgpu
+! UNSUPPORTED: apu
! RUN: %libomptarget-compile-fortran-run-and-check-generic
! Empty dimensions must remain empty when contiguous dimensions are merged.
diff --git a/offload/test/offloading/fortran/target-update-noncontiguous-nowait.f90 b/offload/test/offloading/fortran/target-update-noncontiguous-nowait.f90
new file mode 100644
index 00000000000000..7a25a915062b4d
--- /dev/null
+++ b/offload/test/offloading/fortran/target-update-noncontiguous-nowait.f90
@@ -0,0 +1,77 @@
+! REQUIRES: flang, amdgpu
+! UNSUPPORTED: apu
+! RUN: %libomptarget-compile-fortran-generic
+! RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=0 %libomptarget-run-generic | %fcheck-generic
+
+! Outstanding updates from the same call site must retain their own bounds.
+program target_update_noncontiguous_nowait
+ use omp_lib
+ implicit none
+ integer :: a(4,8,3), actual(4,8,3), i, b, k, release, worker_ready, ready
+
+ do k = 1, 3
+ do b = 1, 8
+ do i = 1, 4
+ a(i,b,k) = i + 10*b + 100*k
+ end do
+ end do
+ end do
+!$omp target enter data map(alloc:a)
+!$omp target map(present,alloc:a)
+ do k = 1, 3
+ do b = 1, 8
+ do i = 1, 4
+ a(i,b,k) = -9999
+ end do
+ end do
+ end do
+!$omp end target
+
+ release = 0
+ worker_ready = 0
+ call omp_set_dynamic(.false.)
+!$omp parallel num_threads(2) shared(a,release,worker_ready) private(ready,b)
+ if (omp_get_num_threads() /= 2) stop 2
+ if (omp_get_thread_num() == 0) then
+ ! With hidden helper threads disabled, keep the other worker outside the
+ ! scheduler until all updates are queued. Each iteration has distinct bounds.
+ do
+!$omp atomic read
+ ready = worker_ready
+ if (ready /= 0) exit
+ end do
+ do b = 1, 8
+!$omp target update to(a(:,b:b,:)) nowait
+ end do
+!$omp atomic write
+ release = 1
+!$omp taskwait
+ else
+!$omp atomic write
+ worker_ready = 1
+ do
+!$omp atomic read
+ ready = release
+ if (ready /= 0) exit
+ end do
+ end if
+!$omp end parallel
+
+!$omp target map(present,alloc:a) map(from:actual)
+ do k = 1, 3
+ do b = 1, 8
+ do i = 1, 4
+ actual(i,b,k) = a(i,b,k)
+ end do
+ end do
+ end do
+!$omp end target
+!$omp target exit data map(delete:a)
+ if (any(actual /= a)) then
+ print *, 'FAIL: incorrect elements:', count(actual /= a)
+ stop 1
+ end if
+ print *, 'PASS'
+end program
+
+! CHECK: PASS
diff --git a/offload/test/offloading/fortran/target-update-noncontiguous.f90 b/offload/test/offloading/fortran/target-update-noncontiguous.f90
index f9e2e134e8c0da..cc767d97ed38ea 100644
--- a/offload/test/offloading/fortran/target-update-noncontiguous.f90
+++ b/offload/test/offloading/fortran/target-update-noncontiguous.f90
@@ -1,4 +1,5 @@
-! REQUIRES: flang, amdgpu
+! REQUIRES: flang, flang-rt, libc, amdgpu
+! UNSUPPORTED: apu
! RUN: %libomptarget-compile-fortran-run-and-check-generic
! A section of the second dimension has gaps between the (k, s) slices.
diff --git a/offload/test/offloading/fortran/target-update-wide-character.f90 b/offload/test/offloading/fortran/target-update-wide-character.f90
index b3a031316f6b5d..eeed3efd7702c1 100644
--- a/offload/test/offloading/fortran/target-update-wide-character.f90
+++ b/offload/test/offloading/fortran/target-update-wide-character.f90
@@ -1,4 +1,5 @@
! REQUIRES: flang, amdgpu
+! UNSUPPORTED: apu
! RUN: %libomptarget-compile-fortran-run-and-check-generic
! Assumed-length character bounds count code units, not bytes. Check every
diff --git a/offload/test/offloading/non_contiguous_update_empty.cpp b/offload/test/offloading/non_contiguous_update_empty.cpp
index 7f029b5b802d6f..8064d0d97ef557 100644
--- a/offload/test/offloading/non_contiguous_update_empty.cpp
+++ b/offload/test/offloading/non_contiguous_update_empty.cpp
@@ -1,4 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: apu
#include <cassert>
#include <climits>
More information about the Mlir-commits
mailing list