Skip to content

Commit 7a06703

Browse files
skc7skc7
committed
[Flang] Update assign_omp logic
1 parent 6aa3346 commit 7a06703

File tree

4 files changed

+126
-105
lines changed

4 files changed

+126
-105
lines changed

‎flang-rt/lib/runtime/assign_omp.cpp

Lines changed: 31 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -18,9 +18,23 @@
1818
#include <omp.h>
1919

2020
namespace Fortran::runtime {
21+
namespace omp {
22+
23+
typedef int32_t OMPDeviceTy;
24+
25+
template <typename T> static T *getDevicePtr(T *anyPtr, OMPDeviceTy ompDevice) {
26+
auto voidAnyPtr = reinterpret_cast<void *>(anyPtr);
27+
// If not present on the device it should already be a device ptr
28+
if (!omp_target_is_present(voidAnyPtr, ompDevice))
29+
return anyPtr;
30+
T *device_ptr = nullptr;
31+
#pragma omp target data use_device_ptr(anyPtr) device(ompDevice)
32+
device_ptr = anyPtr;
33+
return device_ptr;
34+
}
2135

2236
RT_API_ATTRS static void Assign(Descriptor &to, const Descriptor &from,
23-
Terminator &terminator, int flags, int32_t omp_device) {
37+
Terminator &terminator, int flags, OMPDeviceTy omp_device) {
2438
std::size_t toElementBytes{to.ElementBytes()};
2539
std::size_t fromElementBytes{from.ElementBytes()};
2640
std::size_t toElements{to.Elements()};
@@ -31,42 +45,34 @@ RT_API_ATTRS static void Assign(Descriptor &to, const Descriptor &from,
3145
if (toElements != fromElements)
3246
terminator.Crash("Assign: toElements != fromElements");
3347

34-
void *host_to_ptr = to.raw().base_addr;
35-
void *host_from_ptr = from.raw().base_addr;
48+
// Get base addresses and calculate length
49+
void *to_base = to.raw().base_addr;
50+
void *from_base = from.raw().base_addr;
3651
size_t length = toElements * toElementBytes;
3752

38-
printf("assign length: %zu\n", length);
53+
// Get device pointers after ensuring data is on device
54+
void *to_ptr = getDevicePtr(to_base, omp_device);
55+
void *from_ptr = getDevicePtr(from_base, omp_device);
3956

40-
if (!omp_target_is_present(host_to_ptr, omp_device))
41-
terminator.Crash("Assign: !omp_target_is_present(host_to_ptr, omp_device)");
42-
if (!omp_target_is_present(host_from_ptr, omp_device))
43-
terminator.Crash(
44-
"Assign: !omp_target_is_present(host_from_ptr, omp_device)");
45-
46-
printf("host_to_ptr: %p\n", host_to_ptr);
47-
#pragma omp target data use_device_ptr(host_to_ptr, host_from_ptr) device(omp_device)
48-
{
49-
printf("device_to_ptr: %p\n", host_to_ptr);
50-
// TODO do we need to handle overlapping memory? does this function do that?
51-
omp_target_memcpy(host_to_ptr, host_from_ptr, length, /*dst_offset*/ 0,
52-
/*src_offset*/ 0, /*dst*/ omp_device, /*src*/ omp_device);
53-
}
57+
// Perform copy between device pointers
58+
int result = omp_target_memcpy(to_ptr, from_ptr, length,
59+
/*dst_offset*/ 0, /*src_offset*/ 0, omp_device, omp_device);
5460

61+
if (result != 0)
62+
terminator.Crash("Assign: omp_target_memcpy failed");
5563
return;
5664
}
5765

5866
extern "C" {
5967
RT_EXT_API_GROUP_BEGIN
6068
void RTDEF(Assign_omp)(Descriptor &to, const Descriptor &from,
61-
const char *sourceFile, int sourceLine, int32_t omp_device) {
69+
const char *sourceFile, int sourceLine, omp::OMPDeviceTy omp_device) {
6270
Terminator terminator{sourceFile, sourceLine};
63-
// All top-level defined assignments can be recognized in semantics and
64-
// will have been already been converted to calls, so don't check for
65-
// defined assignment apart from components.
66-
Assign(to, from, terminator,
71+
omp::Assign(to, from, terminator,
6772
MaybeReallocate | NeedFinalization | ComponentCanBeDefinedAssignment,
6873
omp_device);
6974
}
70-
} // extern "C"
7175

72-
}
76+
} // extern "C"
77+
} // namespace omp
78+
} // namespace Fortran::runtime

‎flang/lib/Optimizer/OpenMP/LowerWorkdistribute.cpp

Lines changed: 51 additions & 48 deletions
Original file line numberDiff line numberDiff line change
@@ -357,55 +357,77 @@ struct SplitTargetResult {
357357
/// original data region and avoid unnecessary data movement at each of the
358358
/// subkernels - we split the target region into a target_data{target}
359359
/// nest where only the outer one moves the data
360-
std::optional<SplitTargetResult> splitTargetData(omp::TargetOp targetOp, RewriterBase &rewriter) {
361-
360+
std::optional<SplitTargetResult> splitTargetData(omp::TargetOp targetOp,
361+
RewriterBase &rewriter) {
362362
auto loc = targetOp->getLoc();
363363
if (targetOp.getMapVars().empty()) {
364364
LLVM_DEBUG(llvm::dbgs() << DEBUG_TYPE << " target region has no data maps\n");
365365
return std::nullopt;
366366
}
367367

368-
// Collect all map_entries with capture(ByRef)
369-
SmallVector<mlir::Value> byRefMapInfos;
370-
SmallVector<omp::MapInfoOp> MapInfos;
368+
SmallVector<omp::MapInfoOp> mapInfos;
371369
for (auto opr : targetOp.getMapVars()) {
372370
auto mapInfo = cast<omp::MapInfoOp>(opr.getDefiningOp());
373-
MapInfos.push_back(mapInfo);
374-
if (mapInfo.getMapCaptureType() == omp::VariableCaptureKind::ByRef)
375-
byRefMapInfos.push_back(opr);
371+
mapInfos.push_back(mapInfo);
372+
}
373+
374+
rewriter.setInsertionPoint(targetOp);
375+
SmallVector<Value> innerMapInfos;
376+
SmallVector<Value> outerMapInfos;
377+
378+
for (auto mapInfo : mapInfos) {
379+
auto originalMapType =
380+
(llvm::omp::OpenMPOffloadMappingFlags)(mapInfo.getMapType());
381+
auto originalCaptureType = mapInfo.getMapCaptureType();
382+
llvm::omp::OpenMPOffloadMappingFlags newMapType;
383+
mlir::omp::VariableCaptureKind newCaptureType;
384+
385+
if (originalCaptureType == mlir::omp::VariableCaptureKind::ByCopy) {
386+
newMapType = originalMapType;
387+
newCaptureType = originalCaptureType;
388+
} else if (originalCaptureType == mlir::omp::VariableCaptureKind::ByRef) {
389+
newMapType = llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE;
390+
newCaptureType = originalCaptureType;
391+
outerMapInfos.push_back(mapInfo);
392+
} else {
393+
llvm_unreachable("Unhandled case");
394+
}
395+
auto innerMapInfo = cast<omp::MapInfoOp>(rewriter.clone(*mapInfo));
396+
innerMapInfo.setMapTypeAttr(rewriter.getIntegerAttr(
397+
rewriter.getIntegerType(64, false),
398+
static_cast<
399+
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
400+
newMapType)));
401+
innerMapInfo.setMapCaptureType(newCaptureType);
402+
innerMapInfos.push_back(innerMapInfo.getResult());
376403
}
377404

378-
// Create the new omp.target_data op with these collected map_entries
379405
rewriter.setInsertionPoint(targetOp);
380406
auto device = targetOp.getDevice();
381407
auto ifExpr = targetOp.getIfExpr();
382408
auto deviceAddrVars = targetOp.getHasDeviceAddrVars();
383409
auto devicePtrVars = targetOp.getIsDevicePtrVars();
384-
auto targetDataOp = rewriter.create<omp::TargetDataOp>(loc, device, ifExpr,
385-
mlir::ValueRange{byRefMapInfos},
386-
deviceAddrVars,
387-
devicePtrVars);
388-
410+
auto targetDataOp = rewriter.create<omp::TargetDataOp>(
411+
loc, device, ifExpr, outerMapInfos, deviceAddrVars, devicePtrVars);
389412
auto taregtDataBlock = rewriter.createBlock(&targetDataOp.getRegion());
390413
rewriter.create<mlir::omp::TerminatorOp>(loc);
391414
rewriter.setInsertionPointToStart(taregtDataBlock);
392415

393-
// Clone mapInfo ops inside omp.target_data region
394-
IRMapping mapping;
395-
for (auto mapInfo : MapInfos) {
396-
rewriter.clone(*mapInfo, mapping);
397-
}
398-
// Clone omp.target from exisiting targetOp inside target_data region.
399-
auto newTargetOp = rewriter.clone(*targetOp, mapping);
416+
auto newTargetOp = rewriter.create<omp::TargetOp>(
417+
targetOp.getLoc(), targetOp.getAllocateVars(),
418+
targetOp.getAllocatorVars(), targetOp.getBareAttr(),
419+
targetOp.getDependKindsAttr(), targetOp.getDependVars(),
420+
targetOp.getDevice(), targetOp.getHasDeviceAddrVars(),
421+
targetOp.getHostEvalVars(), targetOp.getIfExpr(),
422+
targetOp.getInReductionVars(), targetOp.getInReductionByrefAttr(),
423+
targetOp.getInReductionSymsAttr(), targetOp.getIsDevicePtrVars(),
424+
innerMapInfos, targetOp.getNowaitAttr(), targetOp.getPrivateVars(),
425+
targetOp.getPrivateSymsAttr(), targetOp.getThreadLimit(),
426+
targetOp.getPrivateMapsAttr());
427+
rewriter.inlineRegionBefore(targetOp.getRegion(), newTargetOp.getRegion(),
428+
newTargetOp.getRegion().begin());
400429

401-
// Erase TargetOp and its MapInfoOps
402-
rewriter.eraseOp(targetOp);
403-
404-
for (auto mapInfo : MapInfos) {
405-
auto mapInfoRes = mapInfo.getResult();
406-
if (mapInfoRes.getUsers().empty())
407-
rewriter.eraseOp(mapInfo);
408-
}
430+
rewriter.replaceOp(targetOp, newTargetOp);
409431
return SplitTargetResult{cast<omp::TargetOp>(newTargetOp), targetDataOp};
410432
}
411433

@@ -521,25 +543,6 @@ static bool isRecomputableAfterFission(Operation *op, Operation *splitBefore) {
521543
if (isa<fir::DeclareOp>(op))
522544
return true;
523545

524-
if (auto loadOp = dyn_cast<fir::LoadOp>(op)) {
525-
Value memref = loadOp.getMemref();
526-
if (auto blockArg = dyn_cast<BlockArgument>(memref)) {
527-
// 'op' is an operation within the targetOp that 'splitBefore' is also in.
528-
Operation *parentOpOfLoadBlock = op->getBlock()->getParentOp();
529-
// Ensure the blockArg belongs to the entry block of this parent omp.TargetOp.
530-
// This implies the load is from a variable directly mapped into the target region.
531-
if (isa<omp::TargetOp>(parentOpOfLoadBlock) &&
532-
!parentOpOfLoadBlock->getRegions().empty()) {
533-
Block *targetOpEntryBlock = &parentOpOfLoadBlock->getRegions().front().front();
534-
if (blockArg.getOwner() == targetOpEntryBlock) {
535-
// This load is from a direct argument of the target op.
536-
// It's safe to recompute.
537-
return true;
538-
}
539-
}
540-
}
541-
}
542-
543546
llvm::SmallVector<MemoryEffects::EffectInstance> effects;
544547
MemoryEffectOpInterface interface = dyn_cast<MemoryEffectOpInterface>(op);
545548
if (!interface) {

‎flang/test/Transforms/OpenMP/lower-workdistribute-fission-target.mlir

Lines changed: 43 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -11,34 +11,46 @@
1111
// CHECK: %[[VAL_4:.*]] = omp.map.info var_ptr(%[[VAL_1]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "ub"}
1212
// CHECK: %[[VAL_5:.*]] = omp.map.info var_ptr(%[[VAL_2]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "step"}
1313
// CHECK: %[[VAL_6:.*]] = omp.map.info var_ptr(%[[ARG3:.*]] : !fir.ref<index>, index) map_clauses(tofrom) capture(ByRef) -> !fir.ref<index> {name = "addr"}
14+
// CHECK: %[[VAL_7:.*]] = omp.map.info var_ptr(%[[VAL_0]] : !fir.ref<index>, index) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !fir.ref<index> {name = "lb"}
15+
// CHECK: %[[VAL_8:.*]] = omp.map.info var_ptr(%[[VAL_1]] : !fir.ref<index>, index) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !fir.ref<index> {name = "ub"}
16+
// CHECK: %[[VAL_9:.*]] = omp.map.info var_ptr(%[[VAL_2]] : !fir.ref<index>, index) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !fir.ref<index> {name = "step"}
17+
// CHECK: %[[VAL_10:.*]] = omp.map.info var_ptr(%[[ARG3:.*]] : !fir.ref<index>, index) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !fir.ref<index> {name = "addr"}
1418
// CHECK: omp.target_data map_entries(%[[VAL_3]], %[[VAL_4]], %[[VAL_5]], %[[VAL_6]] : !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<index>) {
15-
// CHECK: %[[VAL_7:.*]] = omp.map.info var_ptr(%[[VAL_0]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "lb"}
16-
// CHECK: %[[VAL_8:.*]] = omp.map.info var_ptr(%[[VAL_1]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "ub"}
17-
// CHECK: %[[VAL_9:.*]] = omp.map.info var_ptr(%[[VAL_2]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "step"}
18-
// CHECK: %[[VAL_10:.*]] = omp.map.info var_ptr(%[[ARG3:.*]] : !fir.ref<index>, index) map_clauses(tofrom) capture(ByRef) -> !fir.ref<index> {name = "addr"}
19-
// CHECK: %[[VAL_11:.*]] = fir.alloca !fir.heap<index>
20-
// CHECK: %[[VAL_12:.*]] = omp.map.info var_ptr(%[[VAL_11]] : !fir.ref<!fir.heap<index>>, !fir.heap<index>) map_clauses(from) capture(ByRef) -> !fir.ref<!fir.heap<index>> {name = "__flang_workdistribute_from"}
21-
// CHECK: %[[VAL_13:.*]] = omp.map.info var_ptr(%[[VAL_11]] : !fir.ref<!fir.heap<index>>, !fir.heap<index>) map_clauses(to) capture(ByRef) -> !fir.ref<!fir.heap<index>> {name = "__flang_workdistribute_to"}
22-
// CHECK: %[[VAL_14:.*]] = arith.constant 1 : index
23-
// CHECK: %[[VAL_15:.*]] = fir.load %[[VAL_0]] : !fir.ref<index>
24-
// CHECK: %[[VAL_16:.*]] = fir.load %[[VAL_1]] : !fir.ref<index>
25-
// CHECK: %[[VAL_17:.*]] = fir.load %[[VAL_2]] : !fir.ref<index>
26-
// CHECK: %[[VAL_18:.*]] = arith.addi %[[VAL_16]], %[[VAL_16]] : index
27-
// CHECK: %[[VAL_19:.*]] = llvm.mlir.constant(0 : i32) : i32
28-
// CHECK: %[[VAL_20:.*]] = "fir.omp_target_allocmem"(%[[VAL_19]], %[[VAL_14]]) <{in_type = index, operandSegmentSizes = array<i32: 1, 0, 1>, uniq_name = "dev_buf"}> : (i32, index) -> !fir.heap<index>
29-
// CHECK: fir.store %[[VAL_20]] to %[[VAL_11]] : !fir.ref<!fir.heap<index>>
30-
// CHECK: omp.target map_entries(%[[VAL_7]] -> %[[VAL_21:.*]], %[[VAL_8]] -> %[[VAL_22:.*]], %[[VAL_9]] -> %[[VAL_23:.*]], %[[VAL_10]] -> %[[VAL_24:.*]], %[[VAL_13]] -> %[[VAL_25:.*]] : !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<!fir.heap<index>>) {
31-
// CHECK: %[[VAL_26:.*]] = fir.load %[[VAL_25]] : !fir.llvm_ptr<!fir.heap<index>>
32-
// CHECK: %[[VAL_27:.*]] = fir.load %[[VAL_21]] : !fir.ref<index>
33-
// CHECK: %[[VAL_28:.*]] = fir.load %[[VAL_22]] : !fir.ref<index>
34-
// CHECK: %[[VAL_29:.*]] = fir.load %[[VAL_23]] : !fir.ref<index>
35-
// CHECK: %[[VAL_30:.*]] = arith.addi %[[VAL_28]], %[[VAL_28]] : index
19+
// CHECK: %[[VAL_11:.*]] = fir.alloca index
20+
// CHECK: %[[VAL_12:.*]] = omp.map.info var_ptr(%[[VAL_11]] : !fir.ref<index>, index) map_clauses(from) capture(ByRef) -> !fir.ref<index> {name = "__flang_workdistribute_from"}
21+
// CHECK: %[[VAL_13:.*]] = omp.map.info var_ptr(%[[VAL_11]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "__flang_workdistribute_to"}
22+
// CHECK: %[[VAL_14:.*]] = fir.alloca index
23+
// CHECK: %[[VAL_15:.*]] = omp.map.info var_ptr(%[[VAL_14]] : !fir.ref<index>, index) map_clauses(from) capture(ByRef) -> !fir.ref<index> {name = "__flang_workdistribute_from"}
24+
// CHECK: %[[VAL_16:.*]] = omp.map.info var_ptr(%[[VAL_14]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "__flang_workdistribute_to"}
25+
// CHECK: %[[VAL_17:.*]] = fir.alloca index
26+
// CHECK: %[[VAL_18:.*]] = omp.map.info var_ptr(%[[VAL_17]] : !fir.ref<index>, index) map_clauses(from) capture(ByRef) -> !fir.ref<index> {name = "__flang_workdistribute_from"}
27+
// CHECK: %[[VAL_19:.*]] = omp.map.info var_ptr(%[[VAL_17]] : !fir.ref<index>, index) map_clauses(to) capture(ByRef) -> !fir.ref<index> {name = "__flang_workdistribute_to"}
28+
// CHECK: %[[VAL_20:.*]] = fir.alloca !fir.heap<index>
29+
// CHECK: %[[VAL_21:.*]] = omp.map.info var_ptr(%[[VAL_20]] : !fir.ref<!fir.heap<index>>, !fir.heap<index>) map_clauses(from) capture(ByRef) -> !fir.ref<!fir.heap<index>> {name = "__flang_workdistribute_from"}
30+
// CHECK: %[[VAL_22:.*]] = omp.map.info var_ptr(%[[VAL_20]] : !fir.ref<!fir.heap<index>>, !fir.heap<index>) map_clauses(to) capture(ByRef) -> !fir.ref<!fir.heap<index>> {name = "__flang_workdistribute_to"}
31+
// CHECK: %[[VAL_23:.*]] = arith.constant 1 : index
32+
// CHECK: %[[VAL_24:.*]] = fir.load %[[VAL_0]] : !fir.ref<index>
33+
// CHECK: %[[VAL_25:.*]] = fir.load %[[VAL_1]] : !fir.ref<index>
34+
// CHECK: %[[VAL_26:.*]] = fir.load %[[VAL_2]] : !fir.ref<index>
35+
// CHECK: %[[VAL_27:.*]] = arith.addi %[[VAL_25]], %[[VAL_25]] : index
36+
// CHECK: %[[VAL_28:.*]] = llvm.mlir.constant(0 : i32) : i32
37+
// CHECK: %[[VAL_29:.*]] = "fir.omp_target_allocmem"(%[[VAL_28]], %[[VAL_23]]) <{in_type = index, operandSegmentSizes = array<i32: 1, 0, 1>, uniq_name = "dev_buf"}> : (i32, index) -> !fir.heap<index>
38+
// CHECK: fir.store %[[VAL_24]] to %[[VAL_11]] : !fir.ref<index>
39+
// CHECK: fir.store %[[VAL_25]] to %[[VAL_14]] : !fir.ref<index>
40+
// CHECK: fir.store %[[VAL_26]] to %[[VAL_17]] : !fir.ref<index>
41+
// CHECK: fir.store %[[VAL_29]] to %[[VAL_20]] : !fir.ref<!fir.heap<index>>
42+
// CHECK: omp.target map_entries(%[[VAL_7]] -> %[[VAL_30:.*]], %[[VAL_8]] -> %[[VAL_31:.*]], %[[VAL_9]] -> %[[VAL_32:.*]], %[[VAL_10]] -> %[[VAL_33:.*]], %[[VAL_13]] -> %[[VAL_34:.*]], %[[VAL_16]] -> %[[VAL_35:.*]], %[[VAL_19]] -> %[[VAL_36:.*]], %[[VAL_22]] -> %[[VAL_37:.*]] : !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<index>, !fir.ref<!fir.heap<index>>) {
43+
// CHECK: %[[VAL_38:.*]] = fir.load %[[VAL_34]] : !fir.llvm_ptr<index>
44+
// CHECK: %[[VAL_39:.*]] = fir.load %[[VAL_35]] : !fir.llvm_ptr<index>
45+
// CHECK: %[[VAL_40:.*]] = fir.load %[[VAL_36]] : !fir.llvm_ptr<index>
46+
// CHECK: %[[VAL_41:.*]] = fir.load %[[VAL_37]] : !fir.llvm_ptr<!fir.heap<index>>
47+
// CHECK: %[[VAL_42:.*]] = arith.addi %[[VAL_39]], %[[VAL_39]] : index
3648
// CHECK: omp.teams {
3749
// CHECK: omp.parallel {
3850
// CHECK: omp.distribute {
3951
// CHECK: omp.wsloop {
40-
// CHECK: omp.loop_nest (%[[VAL_31:.*]]) : index = (%[[VAL_27]]) to (%[[VAL_28]]) inclusive step (%[[VAL_29]]) {
41-
// CHECK: fir.store %[[VAL_30]] to %[[VAL_26]] : !fir.heap<index>
52+
// CHECK: omp.loop_nest (%[[VAL_43:.*]]) : index = (%[[VAL_38]]) to (%[[VAL_39]]) inclusive step (%[[VAL_40]]) {
53+
// CHECK: fir.store %[[VAL_42]] to %[[VAL_41]] : !fir.heap<index>
4254
// CHECK: omp.yield
4355
// CHECK: }
4456
// CHECK: } {omp.composite}
@@ -49,14 +61,14 @@
4961
// CHECK: }
5062
// CHECK: omp.terminator
5163
// CHECK: }
52-
// CHECK: %[[VAL_32:.*]] = fir.load %[[VAL_11]] : !fir.ref<!fir.heap<index>>
53-
// CHECK: %[[VAL_33:.*]] = fir.load %[[VAL_0]] : !fir.ref<index>
54-
// CHECK: %[[VAL_34:.*]] = fir.load %[[VAL_1]] : !fir.ref<index>
55-
// CHECK: %[[VAL_35:.*]] = fir.load %[[VAL_2]] : !fir.ref<index>
56-
// CHECK: %[[VAL_36:.*]] = arith.addi %[[VAL_34]], %[[VAL_34]] : index
57-
// CHECK: fir.store %[[VAL_33]] to %[[VAL_32]] : !fir.heap<index>
58-
// CHECK: %[[VAL_37:.*]] = llvm.mlir.constant(0 : i32) : i32
59-
// CHECK: "fir.omp_target_freemem"(%[[VAL_37]], %[[VAL_32]]) : (i32, !fir.heap<index>) -> ()
64+
// CHECK: %[[VAL_44:.*]] = fir.load %[[VAL_11]] : !fir.ref<index>
65+
// CHECK: %[[VAL_45:.*]] = fir.load %[[VAL_14]] : !fir.ref<index>
66+
// CHECK: %[[VAL_46:.*]] = fir.load %[[VAL_17]] : !fir.ref<index>
67+
// CHECK: %[[VAL_47:.*]] = fir.load %[[VAL_20]] : !fir.ref<!fir.heap<index>>
68+
// CHECK: %[[VAL_48:.*]] = arith.addi %[[VAL_45]], %[[VAL_45]] : index
69+
// CHECK: fir.store %[[VAL_44]] to %[[VAL_47]] : !fir.heap<index>
70+
// CHECK: %[[VAL_49:.*]] = llvm.mlir.constant(0 : i32) : i32
71+
// CHECK: "fir.omp_target_freemem"(%[[VAL_49]], %[[VAL_47]]) : (i32, !fir.heap<index>) -> ()
6072
// CHECK: omp.terminator
6173
// CHECK: }
6274
// CHECK: return

0 commit comments

Comments
 (0)