Skip to content

[OpenACC][CIR] Implement 'modifier-list' lowering #145770

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Jun 26, 2025
Merged

[OpenACC][CIR] Implement 'modifier-list' lowering #145770

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
188633e
[OpenACC][CIR] Implement 'modifier-list' lowering
erichkeane Jun 25, 2025
6a4a71e
Add static assert to make sure types match for modifier kind
erichkeane Jun 26, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 11 additions & 8 deletions clang/include/clang/Basic/OpenACCKinds.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
58 changes: 44 additions & 14 deletions clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This feels very fragile. Is there no way to use the same type in both places? Failing that, maybe you could statically assert the equality of the bits that you think are the same.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Unfortunately not on the using the same type. The FE and the MLIR have different needs when it comes to tracking always. CFE needs to differentiate always from someone typing alwaysin and alwaysout for diagnostics purposes. The MLIR passes need always->alwaysin|alwaysout to ensure they don't miss opt.

The static assert however is a great idea!

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);
Expand Down Expand Up @@ -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());
Expand All @@ -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,
Expand All @@ -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());
}
Expand Down Expand Up @@ -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);
Expand All @@ -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");
}
Expand All @@ -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);
Expand All @@ -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);
Expand All @@ -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);
Expand All @@ -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);
Expand Down
73 changes: 65 additions & 8 deletions clang/test/CIR/CodeGenOpenACC/combined-copy.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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];
Expand Down Expand Up @@ -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"}
}
Loading