Author: Erich Keane Date: 2025-06-26T09:25:59-07:00 New Revision: 69bbf2135e7dcd685acf6111e148e620cfe10280
URL: https://github.com/llvm/llvm-project/commit/69bbf2135e7dcd685acf6111e148e620cfe10280 DIFF: https://github.com/llvm/llvm-project/commit/69bbf2135e7dcd685acf6111e148e620cfe10280.diff LOG: [OpenACC][CIR] Implement 'modifier-list' lowering (#145770) Some of the 'data' clauses can have a 'modifier-list' which specifies one of a few keywords from a list. This patch adds support for lowering them following #144806. We have to keep a separate enum from MLIR, since we have to keep 'always' around for semantic reasons, whereas the dialect doesn't differentiate these. This patch ensures we get these right for the only applicable clause so far, which is 'copy'. Added: Modified: clang/include/clang/Basic/OpenACCKinds.h clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp clang/test/CIR/CodeGenOpenACC/combined-copy.c clang/test/CIR/CodeGenOpenACC/compute-copy.c Removed: ################################################################################ diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h index a7a29e2add20a..938385679e3ab 100644 --- a/clang/include/clang/Basic/OpenACCKinds.h +++ b/clang/include/clang/Basic/OpenACCKinds.h @@ -634,16 +634,19 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out, } // Represents the 'modifier' of a 'modifier-list', as applied to copy, copyin, -// copyout, and create. Implemented as a 'bitmask' +// copyout, and create. Implemented as a 'bitmask'. +// Note: This attempts to synchronize with mlir::acc::DataClauseModifier, +// however has to store `Always` separately(whereas MLIR has it as AlwaysIn & +// AlwaysOut). However, we keep them in sync so that we can cast between them. enum class OpenACCModifierKind : uint8_t { Invalid = 0, - Always = 1 << 0, - AlwaysIn = 1 << 1, - AlwaysOut = 1 << 2, - Readonly = 1 << 3, - Zero = 1 << 4, - Capture = 1 << 5, - LLVM_MARK_AS_BITMASK_ENUM(Capture) + Zero = 1 << 0, + Readonly = 1 << 1, + AlwaysIn = 1 << 2, + AlwaysOut = 1 << 3, + Capture = 1 << 4, + Always = 1 << 5, + LLVM_MARK_AS_BITMASK_ENUM(Always) }; inline bool isOpenACCModifierBitSet(OpenACCModifierKind List, diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 899e91574e917..1454cee336a09 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -286,16 +286,39 @@ class OpenACCClauseCIREmitter final std::move(bounds)}; } + mlir::acc::DataClauseModifier + convertModifiers(OpenACCModifierKind modifiers) { + using namespace mlir::acc; + static_assert(static_cast<int>(OpenACCModifierKind::Zero) == + static_cast<int>(DataClauseModifier::zero) && + static_cast<int>(OpenACCModifierKind::Readonly) == + static_cast<int>(DataClauseModifier::readonly) && + static_cast<int>(OpenACCModifierKind::AlwaysIn) == + static_cast<int>(DataClauseModifier::alwaysin) && + static_cast<int>(OpenACCModifierKind::AlwaysOut) == + static_cast<int>(DataClauseModifier::alwaysout) && + static_cast<int>(OpenACCModifierKind::Capture) == + static_cast<int>(DataClauseModifier::capture)); + + DataClauseModifier mlirModifiers{}; + + // The MLIR representation of this represents `always` as `alwaysin` + + // `alwaysout`. So do a small fixup here. + if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) { + mlirModifiers = mlirModifiers | DataClauseModifier::always; + modifiers &= ~OpenACCModifierKind::Always; + } + + mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers); + return mlirModifiers; + } + template <typename BeforeOpTy, typename AfterOpTy> void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause, - bool structured, bool implicit) { + OpenACCModifierKind modifiers, bool structured, + bool implicit) { DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand); - // TODO: OpenACC: we should comprehend the 'modifier-list' here for the data - // operand. At the moment, we don't have a uniform way to assign these - // properly, and the dialect cannot represent anything other than 'readonly' - // and 'zero' on copyin/copyout/create, so for now, we skip it. - auto beforeOp = builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured, implicit, opInfo.name, opInfo.bounds); @@ -323,6 +346,8 @@ class OpenACCClauseCIREmitter final // Set the 'rest' of the info for both operations. beforeOp.setDataClause(dataClause); afterOp.setDataClause(dataClause); + beforeOp.setModifiers(convertModifiers(modifiers)); + afterOp.setModifiers(convertModifiers(modifiers)); // Make sure we record these, so 'async' values can be updated later. dataOperands.push_back(beforeOp.getOperation()); @@ -331,7 +356,8 @@ class OpenACCClauseCIREmitter final template <typename BeforeOpTy> void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause, - bool structured, bool implicit) { + OpenACCModifierKind modifiers, bool structured, + bool implicit) { DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand); auto beforeOp = builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured, @@ -340,6 +366,8 @@ class OpenACCClauseCIREmitter final // Set the 'rest' of the info for the operation. beforeOp.setDataClause(dataClause); + beforeOp.setModifiers(convertModifiers(modifiers)); + // Make sure we record these, so 'async' values can be updated later. dataOperands.push_back(beforeOp.getOperation()); } @@ -818,7 +846,8 @@ class OpenACCClauseCIREmitter final mlir::acc::KernelsOp>) { for (auto var : clause.getVarList()) addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>( - var, mlir::acc::DataClause::acc_copy, /*structured=*/true, + var, mlir::acc::DataClause::acc_copy, clause.getModifierList(), + /*structured=*/true, /*implicit=*/false); } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); @@ -833,8 +862,8 @@ class OpenACCClauseCIREmitter final if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) { for (auto var : clause.getVarList()) addDataOperand<mlir::acc::UseDeviceOp>( - var, mlir::acc::DataClause::acc_use_device, - /*structured=*/true, /*implicit=*/false); + var, mlir::acc::DataClause::acc_use_device, {}, /*structured=*/true, + /*implicit=*/false); } else { llvm_unreachable("Unknown construct kind in VisitUseDeviceClause"); } @@ -845,7 +874,8 @@ class OpenACCClauseCIREmitter final mlir::acc::KernelsOp>) { for (auto var : clause.getVarList()) addDataOperand<mlir::acc::DevicePtrOp>( - var, mlir::acc::DataClause::acc_deviceptr, /*structured=*/true, + var, mlir::acc::DataClause::acc_deviceptr, {}, + /*structured=*/true, /*implicit=*/false); } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); @@ -861,7 +891,7 @@ class OpenACCClauseCIREmitter final mlir::acc::KernelsOp>) { for (auto var : clause.getVarList()) addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>( - var, mlir::acc::DataClause::acc_no_create, /*structured=*/true, + var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true, /*implicit=*/false); } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); @@ -877,7 +907,7 @@ class OpenACCClauseCIREmitter final mlir::acc::KernelsOp>) { for (auto var : clause.getVarList()) addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>( - var, mlir::acc::DataClause::acc_present, /*structured=*/true, + var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true, /*implicit=*/false); } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); @@ -893,7 +923,7 @@ class OpenACCClauseCIREmitter final mlir::acc::KernelsOp>) { for (auto var : clause.getVarList()) addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>( - var, mlir::acc::DataClause::acc_attach, /*structured=*/true, + var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true, /*implicit=*/false); } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c index 1c94fa8238ce8..16df179d900da 100644 --- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c @@ -77,29 +77,29 @@ void acc_compute(int parmVar) { // these do nothing to the IR. #pragma acc parallel loop copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3) for(int i = 0; i < 5; ++i); - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc + // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc + // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc #pragma acc serial loop copy(always, alwaysin, alwaysout: localVar1) for(int i = 0; i < 5; ++i); - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc short *localPointer; float localArray[5]; @@ -1102,3 +1102,60 @@ void copy_member_of_array_element_member() { // CHECK-NEXT: } loc // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[GETB]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "outer.inner[2].b"} } + +void modifier_list() { + // CHECK: cir.func @modifier_list() { + int localVar; + // CHECK-NEXT: %[[LOCALVAR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar"] + +#pragma acc parallel loop copy(always:localVar) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"} +#pragma acc serial loop copy(alwaysin:localVar) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"} +#pragma acc kernels loop copy(alwaysout:localVar) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"} +#pragma acc parallel loop copy(capture:localVar) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"} +#pragma acc serial loop copy(capture, always, alwaysin, alwaysout:localVar) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"} +} diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copy.c b/clang/test/CIR/CodeGenOpenACC/compute-copy.c index 0fb150475bc72..b3010ab6e4b69 100644 --- a/clang/test/CIR/CodeGenOpenACC/compute-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/compute-copy.c @@ -69,23 +69,23 @@ void acc_compute(int parmVar) { // these do nothing to the IR. #pragma acc parallel copy(alwaysin: localVar1) copy(alwaysout: localVar2) copy(always: localVar3) ; - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc - // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc + // CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) -> !cir.ptr<!s16i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc + // CHECK-NEXT: %[[COPYIN3:.*]] = acc.copyin varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]], %[[COPYIN2]], %[[COPYIN3]] : !cir.ptr<!s32i>, !cir.ptr<!s16i>, !cir.ptr<!cir.float>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, name = "localVar3"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar2"} loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN3]] : !cir.ptr<!cir.float>) to varPtr(%[[LOCAL3]] : !cir.ptr<!cir.float>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar3"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s16i>) to varPtr(%[[LOCAL2]] : !cir.ptr<!s16i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar2"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar1"} loc #pragma acc serial copy(always, alwaysin, alwaysout: localVar1) ; - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar1"} loc short *localPointer; float localArray[5]; @@ -897,3 +897,46 @@ void acc_compute_members() { // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr<!cir.ptr<!cir.ptr<!cir.double>>>) {dataClause = #acc<data_clause acc_copy>, name = "localStruct.ptrPtrMember[1:3][1:1]"} } + +void modifier_list() { + // CHECK: cir.func @modifier_list() { + int localVar; + // CHECK-NEXT: %[[LOCALVAR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar"] + +#pragma acc parallel copy(always:localVar) + ; + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "localVar"} + +#pragma acc serial copy(alwaysin:localVar) + ; + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysin>, name = "localVar"} +#pragma acc kernels copy(alwaysout:localVar) + ; + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier alwaysout>, name = "localVar"} +#pragma acc parallel copy(capture:localVar) + ; + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "localVar"} +#pragma acc serial copy(capture, always, alwaysin, alwaysout:localVar) + ; + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[LOCALVAR]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "localVar"} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
