From 9d36ed7475abbc485c25ebb34b173107a235d13f Mon Sep 17 00:00:00 2001 From: ouankou Date: Sun, 21 Dec 2025 20:28:14 +0000 Subject: [PATCH 1/6] Refactor AST to Enum-based & Implement OpenACC 3.4 'indirect' Clause --- src/OpenACCASTConstructor.cpp | 291 +++++++----- src/OpenACCASTConstructor.h | 29 +- src/OpenACCIR.cpp | 139 +++--- src/OpenACCIR.h | 107 +++-- src/OpenACCIRToString.cpp | 278 ++++++++++-- src/OpenACCKinds.h | 32 ++ src/acclexer.g4 | 548 +++++++++++++++++++---- src/accparser.g4 | 79 ++-- tests/acc_tester.cpp | 3 +- tests/builtin/CMakeLists.txt | 1 + tests/builtin/indirect.txt | 6 + tests/builtin/reference/ref_indirect.txt | 3 + 12 files changed, 1113 insertions(+), 403 deletions(-) create mode 100644 tests/builtin/indirect.txt create mode 100644 tests/builtin/reference/ref_indirect.txt diff --git a/src/OpenACCASTConstructor.cpp b/src/OpenACCASTConstructor.cpp index 6fa9d9d..990eb1c 100644 --- a/src/OpenACCASTConstructor.cpp +++ b/src/OpenACCASTConstructor.cpp @@ -221,11 +221,12 @@ void OpenACCIRConstructor::enterBind_clause( void OpenACCIRConstructor::exitName_or_string( accparser::Name_or_stringContext *ctx) { - std::string expression = trimEnclosingWhiteSpace(ctx->getText()); - bool is_string = - !expression.empty() && (expression.front() == '"' || expression.front() == '\''); - static_cast(current_clause) - ->setBinding(expression, is_string); + if (current_clause && current_clause->getKind() == ACCC_bind) { + std::string expression = trimEnclosingWhiteSpace(ctx->getText()); + bool is_string = ctx->STRING_LITERAL() != nullptr; + static_cast(current_clause) + ->setBinding(expression, is_string); + } } void OpenACCIRConstructor::exitBind_clause(accparser::Bind_clauseContext *ctx) { @@ -247,6 +248,9 @@ void OpenACCIRConstructor::enterCollapse_clause( void OpenACCIRConstructor::exitCollapse_clause( accparser::Collapse_clauseContext *ctx) { + if (ctx->FORCE()) { + static_cast(current_clause)->setForce(true); + } ((OpenACCCollapseClause *)current_clause) ->mergeClause(current_directive, current_clause); } @@ -259,19 +263,31 @@ void OpenACCIRConstructor::enterAttach_clause( void OpenACCIRConstructor::enterCopy_clause( accparser::Copy_clauseContext *ctx) { current_clause = current_directive->addOpenACCClause(ACCC_copy); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); + if (ctx->PCOPY()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPY_pcopy); + } else if (ctx->PRESENT_OR_COPY()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPY_present_or_copy); + } else { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPY_copy); + } } void OpenACCIRConstructor::enterCopyin_clause( accparser::Copyin_clauseContext *ctx) { current_clause = current_directive->addOpenACCClause(ACCC_copyin); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); - ((OpenACCCopyinClause *)current_clause)->setModifier(ACCC_COPYIN_unspecified); -} - -void OpenACCIRConstructor::exitCopyin_clause_modifier( - accparser::Copyin_clause_modifierContext *ctx) { - ((OpenACCCopyinClause *)current_clause)->setModifier(ACCC_COPYIN_readonly); + if (ctx->PCOPYIN()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPYIN_pcopyin); + } else if (ctx->PRESENT_OR_COPYIN()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPYIN_present_or_copyin); + } else { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPYIN_copyin); + } } void OpenACCIRConstructor::exitCopyin_clause( @@ -280,30 +296,19 @@ void OpenACCIRConstructor::exitCopyin_clause( ->mergeClause(current_directive, current_clause); } -void OpenACCIRConstructor::enterCopyin_no_modifier_clause( - accparser::Copyin_no_modifier_clauseContext *ctx) { - current_clause = current_directive->addOpenACCClause(ACCC_copyin); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); - ((OpenACCCopyinClause *)current_clause)->setModifier(ACCC_COPYIN_unspecified); -} - -void OpenACCIRConstructor::exitCopyin_no_modifier_clause( - accparser::Copyin_no_modifier_clauseContext *ctx) { - ((OpenACCCopyinClause *)current_clause) - ->mergeClause(current_directive, current_clause); -} - void OpenACCIRConstructor::enterCopyout_clause( accparser::Copyout_clauseContext *ctx) { current_clause = current_directive->addOpenACCClause(ACCC_copyout); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); - ((OpenACCCopyoutClause *)current_clause) - ->setModifier(ACCC_COPYOUT_unspecified); -} - -void OpenACCIRConstructor::exitCopyout_clause_modifier( - accparser::Copyout_clause_modifierContext *ctx) { - ((OpenACCCopyoutClause *)current_clause)->setModifier(ACCC_COPYOUT_zero); + if (ctx->PCOPYOUT()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPYOUT_pcopyout); + } else if (ctx->PRESENT_OR_COPYOUT()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPYOUT_present_or_copyout); + } else { + static_cast(current_clause) + ->setVariant(ACCC_DATA_COPYOUT_copyout); + } } void OpenACCIRConstructor::exitCopyout_clause( @@ -312,25 +317,19 @@ void OpenACCIRConstructor::exitCopyout_clause( ->mergeClause(current_directive, current_clause); } -void OpenACCIRConstructor::enterCopyout_no_modifier_clause( - accparser::Copyout_no_modifier_clauseContext *ctx) { - current_clause = current_directive->addOpenACCClause(ACCC_copyout); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); - ((OpenACCCopyoutClause *)current_clause) - ->setModifier(ACCC_COPYOUT_unspecified); -} - -void OpenACCIRConstructor::exitCopyout_no_modifier_clause( - accparser::Copyout_no_modifier_clauseContext *ctx) { - ((OpenACCCopyoutClause *)current_clause) - ->mergeClause(current_directive, current_clause); -} - void OpenACCIRConstructor::enterCreate_clause( accparser::Create_clauseContext *ctx) { current_clause = current_directive->addOpenACCClause(ACCC_create); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); - ((OpenACCCreateClause *)current_clause)->setModifier(ACCC_CREATE_unspecified); + if (ctx->PCREATE()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_CREATE_pcreate); + } else if (ctx->PRESENT_OR_CREATE()) { + static_cast(current_clause) + ->setVariant(ACCC_DATA_CREATE_present_or_create); + } else { + static_cast(current_clause) + ->setVariant(ACCC_DATA_CREATE_create); + } } void OpenACCIRConstructor::exitCreate_clause( @@ -339,22 +338,32 @@ void OpenACCIRConstructor::exitCreate_clause( ->mergeClause(current_directive, current_clause); } -void OpenACCIRConstructor::exitCreate_clause_modifier( - accparser::Create_clause_modifierContext *ctx) { - ((OpenACCCreateClause *)current_clause)->setModifier(ACCC_CREATE_zero); -} +void OpenACCIRConstructor::exitData_clause_modifier( + accparser::Data_clause_modifierContext *ctx) { + if (!current_clause) { + return; + } + auto *data_clause = dynamic_cast(current_clause); + if (!data_clause) { + return; + } -void OpenACCIRConstructor::enterCreate_no_modifier_clause( - accparser::Create_no_modifier_clauseContext *ctx) { - current_clause = current_directive->addOpenACCClause(ACCC_create); - current_clause->setOriginalKeyword(ctx->getStart()->getText()); - ((OpenACCCreateClause *)current_clause)->setModifier(ACCC_CREATE_unspecified); -} + OpenACCDataClauseModifierKind modifier = ACCC_DATA_MOD_unknown; + if (ctx->ALWAYS()) { + modifier = ACCC_DATA_MOD_always; + } else if (ctx->ALWAYSIN()) { + modifier = ACCC_DATA_MOD_alwaysin; + } else if (ctx->ALWAYSOUT()) { + modifier = ACCC_DATA_MOD_alwaysout; + } else if (ctx->CAPTURE()) { + modifier = ACCC_DATA_MOD_capture; + } else if (ctx->READONLY()) { + modifier = ACCC_DATA_MOD_readonly; + } else if (ctx->ZERO()) { + modifier = ACCC_DATA_MOD_zero; + } -void OpenACCIRConstructor::exitCreate_no_modifier_clause( - accparser::Create_no_modifier_clauseContext *ctx) { - ((OpenACCCreateClause *)current_clause) - ->mergeClause(current_directive, current_clause); + data_clause->addModifier(modifier); } void OpenACCIRConstructor::enterDefault_clause( @@ -364,12 +373,12 @@ void OpenACCIRConstructor::enterDefault_clause( void OpenACCIRConstructor::exitDefault_kind( accparser::Default_kindContext *ctx) { - std::string expression = trimEnclosingWhiteSpace(ctx->getText()); - OpenACCDefaultClauseKind kind = ACCC_DEFAULT_unspecified; - if (expression == "none") + OpenACCDefaultClauseKind kind = ACCC_DEFAULT_unknown; + if (ctx->NONE()) { kind = ACCC_DEFAULT_none; - else if (expression == "present") + } else if (ctx->PRESENT()) { kind = ACCC_DEFAULT_present; + } ((OpenACCDefaultClause *)current_clause)->setKind(kind); }; @@ -432,6 +441,26 @@ void OpenACCIRConstructor::enterDevice_type_clause( current_clause = current_directive->addOpenACCClause(ACCC_device_type); } +void OpenACCIRConstructor::exitDevice_type_item( + accparser::Device_type_itemContext *ctx) { + if (!current_clause || current_clause->getKind() != ACCC_device_type) { + return; + } + auto *clause = static_cast(current_clause); + if (ctx->HOST()) { + clause->addDeviceType(ACCC_DEVICE_TYPE_host); + } else if (ctx->ANY()) { + clause->addDeviceType(ACCC_DEVICE_TYPE_any); + } else if (ctx->MULTICORE()) { + clause->addDeviceType(ACCC_DEVICE_TYPE_multicore); + } else if (ctx->DEFAULT()) { + clause->addDeviceType(ACCC_DEVICE_TYPE_default); + } else if (ctx->EXPR()) { + clause->addUnknownDeviceType( + trimEnclosingWhiteSpace(ctx->EXPR()->getText())); + } +} + void OpenACCIRConstructor::enterDeviceptr_clause( accparser::Deviceptr_clauseContext *ctx) { current_clause = current_directive->addOpenACCClause(ACCC_deviceptr); @@ -458,6 +487,38 @@ void OpenACCIRConstructor::exitGang_clause(accparser::Gang_clauseContext *ctx) { ->mergeClause(current_directive, current_clause); }; +void OpenACCIRConstructor::exitGang_arg(accparser::Gang_argContext *ctx) { + if (!current_clause || current_clause->getKind() != ACCC_gang) { + return; + } + auto *gang = static_cast(current_clause); + + OpenACCGangArgKind kind = ACCC_GANG_ARG_other; + std::string value_text; + + if (ctx->NUM()) { + kind = ACCC_GANG_ARG_num; + if (ctx->int_expr()) { + value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); + } + } else if (ctx->DIM()) { + kind = ACCC_GANG_ARG_dim; + if (ctx->int_expr()) { + value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); + } + } else if (ctx->STATIC()) { + kind = ACCC_GANG_ARG_static; + if (ctx->int_expr()) { + value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); + } + } else if (ctx->int_expr()) { + kind = ACCC_GANG_ARG_num_no_keyword; + value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); + } + + gang->addArg(kind, OpenACCExpressionItem{value_text, ACCC_CLAUSE_SEP_comma}); +} + void OpenACCIRConstructor::enterNohost_clause( accparser::Nohost_clauseContext *ctx) { std::string expression = trimEnclosingWhiteSpace(ctx->getText()); @@ -555,43 +616,42 @@ void OpenACCIRConstructor::enterReduction_clause( void OpenACCIRConstructor::exitReduction_operator( accparser::Reduction_operatorContext *ctx) { - std::string expression = trimEnclosingWhiteSpace(ctx->getText()); - OpenACCReductionClauseOperator reduction_operator = - ACCC_REDUCTION_unspecified; - if (expression == "+") + OpenACCReductionClauseOperator reduction_operator = ACCC_REDUCTION_unknown; + if (ctx->ADD()) { reduction_operator = ACCC_REDUCTION_add; - else if (expression == "-") + } else if (ctx->SUB()) { reduction_operator = ACCC_REDUCTION_sub; - else if (expression == "*") + } else if (ctx->MUL()) { reduction_operator = ACCC_REDUCTION_mul; - else if (expression == "max") + } else if (ctx->MAX()) { reduction_operator = ACCC_REDUCTION_max; - else if (expression == "min") + } else if (ctx->MIN()) { reduction_operator = ACCC_REDUCTION_min; - else if (expression == "&") + } else if (ctx->BITAND()) { reduction_operator = ACCC_REDUCTION_bitand; - else if (expression == "|") + } else if (ctx->BITOR()) { reduction_operator = ACCC_REDUCTION_bitor; - else if (expression == "^") + } else if (ctx->BITXOR()) { reduction_operator = ACCC_REDUCTION_bitxor; - else if (expression == "&&") + } else if (ctx->LOGAND()) { reduction_operator = ACCC_REDUCTION_logand; - else if (expression == "||") + } else if (ctx->LOGOR()) { reduction_operator = ACCC_REDUCTION_logor; - else if (expression == ".and." || expression == ".AND.") + } else if (ctx->FORT_AND()) { reduction_operator = ACCC_REDUCTION_fort_and; - else if (expression == ".or." || expression == ".OR.") + } else if (ctx->FORT_OR()) { reduction_operator = ACCC_REDUCTION_fort_or; - else if (expression == ".eqv." || expression == ".EQV.") + } else if (ctx->FORT_EQV()) { reduction_operator = ACCC_REDUCTION_fort_eqv; - else if (expression == ".neqv." || expression == ".NEQV.") + } else if (ctx->FORT_NEQV()) { reduction_operator = ACCC_REDUCTION_fort_neqv; - else if (expression == "iand" || expression == "IAND") + } else if (ctx->FORT_IAND()) { reduction_operator = ACCC_REDUCTION_fort_iand; - else if (expression == "ior" || expression == "IOR") + } else if (ctx->FORT_IOR()) { reduction_operator = ACCC_REDUCTION_fort_ior; - else if (expression == "ieor" || expression == "IEOR") + } else if (ctx->FORT_IEOR()) { reduction_operator = ACCC_REDUCTION_fort_ieor; + } ((OpenACCReductionClause *)current_clause)->setOperator(reduction_operator); }; @@ -828,6 +888,9 @@ void OpenACCIRConstructor::exitInt_expr(accparser::Int_exprContext *ctx) { kind == ACCC_vector_length) { return; } + if (kind == ACCC_gang) { + return; + } if (kind == ACCC_wait) { static_cast(current_clause) ->addAsyncId(OpenACCExpressionItem{ @@ -880,11 +943,6 @@ void OpenACCIRConstructor::exitVar(accparser::VarContext *ctx) { ((OpenACCCacheDirective *)current_directive)->addVar(expression); } else { if (current_clause) { - if (current_clause->getKind() == ACCC_device_type) { - static_cast(current_clause) - ->addDeviceTypeString(expression); - return; - } if (current_clause->getKind() == ACCC_device) { static_cast(current_clause) ->addDevice(expression); @@ -911,29 +969,6 @@ void OpenACCIRConstructor::exitVar(accparser::VarContext *ctx) { static_cast(current_clause)->addVar(expression); return; } - if (current_clause->getKind() == ACCC_gang) { - auto *gang = static_cast(current_clause); - OpenACCGangArgKind kind = ACCC_GANG_ARG_other; - OpenACCExpressionItem value{expression, ACCC_CLAUSE_SEP_comma}; - size_t colon = expression.find(':'); - if (colon != std::string::npos) { - std::string key = trimEnclosingWhiteSpace(expression.substr(0, colon)); - value.text = trimEnclosingWhiteSpace(expression.substr(colon + 1)); - std::transform(key.begin(), key.end(), key.begin(), ::tolower); - if (key == "num") { - kind = ACCC_GANG_ARG_num; - } else if (key == "dim") { - kind = ACCC_GANG_ARG_dim; - } else if (key == "static") { - kind = ACCC_GANG_ARG_static; - } else { - kind = ACCC_GANG_ARG_other; - value.text = expression; - } - } - gang->addArg(kind, value); - return; - } if (current_clause->getKind() == ACCC_tile) { static_cast(current_clause) ->addTileSize(expression); @@ -981,3 +1016,27 @@ OpenACCDirective *parseOpenACC(std::string source) { return current_directive; } + +// OpenACCIRConstructor::enterIndirect_clause implementation +void OpenACCIRConstructor::enterIndirect_clause( + accparser::Indirect_clauseContext *ctx) { + current_clause = current_directive->addOpenACCClause(ACCC_indirect); + auto *indirect = static_cast(current_clause); + indirect->setPresent(true); + if (ctx->name_or_string()) { + std::string text = + trimEnclosingWhiteSpace(ctx->name_or_string()->getText()); + bool is_str = ctx->name_or_string()->STRING_LITERAL() != nullptr; + indirect->setValue(text, is_str); + } +} + +void OpenACCIRConstructor::exitIndirect_clause( + accparser::Indirect_clauseContext *ctx) { + if (current_clause && current_clause->getKind() == ACCC_indirect) { + if (auto *indirect = + dynamic_cast(current_clause)) { + indirect->mergeClause(current_directive, current_clause); + } + } +} diff --git a/src/OpenACCASTConstructor.h b/src/OpenACCASTConstructor.h index e128163..62a5b80 100644 --- a/src/OpenACCASTConstructor.h +++ b/src/OpenACCASTConstructor.h @@ -21,7 +21,7 @@ class OpenACCIRConstructor : public accparserBaseListener { virtual void enterCache_directive(accparser::Cache_directiveContext * /*ctx*/) override; virtual void exitCache_directive_modifier( - accparser::accparser::Cache_directive_modifierContext * /*ctx*/) override; + accparser::Cache_directive_modifierContext * /*ctx*/) override; virtual void enterData_directive(accparser::Data_directiveContext * /*ctx*/) override; virtual void enterDeclare_directive( @@ -90,34 +90,18 @@ class OpenACCIRConstructor : public accparserBaseListener { enterCopy_clause(accparser::Copy_clauseContext * /*ctx*/) override; virtual void enterCopyin_clause(accparser::Copyin_clauseContext * /*ctx*/) override; - virtual void exitCopyin_clause_modifier( - accparser::Copyin_clause_modifierContext * /*ctx*/) override; virtual void exitCopyin_clause(accparser::Copyin_clauseContext * /*ctx*/) override; - virtual void enterCopyin_no_modifier_clause( - accparser::Copyin_no_modifier_clauseContext * /*ctx*/) override; - virtual void exitCopyin_no_modifier_clause( - accparser::Copyin_no_modifier_clauseContext * /*ctx*/) override; virtual void enterCopyout_clause(accparser::Copyout_clauseContext * /*ctx*/) override; - virtual void exitCopyout_clause_modifier( - accparser::Copyout_clause_modifierContext * /*ctx*/) override; virtual void exitCopyout_clause(accparser::Copyout_clauseContext * /*ctx*/) override; - virtual void enterCopyout_no_modifier_clause( - accparser::Copyout_no_modifier_clauseContext * /*ctx*/) override; - virtual void exitCopyout_no_modifier_clause( - accparser::Copyout_no_modifier_clauseContext * /*ctx*/) override; virtual void enterCreate_clause(accparser::Create_clauseContext * /*ctx*/) override; - virtual void exitCreate_clause_modifier( - accparser::Create_clause_modifierContext * /*ctx*/) override; virtual void exitCreate_clause(accparser::Create_clauseContext * /*ctx*/) override; - virtual void enterCreate_no_modifier_clause( - accparser::Create_no_modifier_clauseContext * /*ctx*/) override; - virtual void exitCreate_no_modifier_clause( - accparser::Create_no_modifier_clauseContext * /*ctx*/) override; + virtual void exitData_clause_modifier( + accparser::Data_clause_modifierContext * /*ctx*/) override; virtual void enterDefault_clause(accparser::Default_clauseContext * /*ctx*/) override; virtual void @@ -143,6 +127,8 @@ class OpenACCIRConstructor : public accparserBaseListener { virtual void enterDevice_type_clause( accparser::Device_type_clauseContext * /*ctx*/) override; virtual void + exitDevice_type_item(accparser::Device_type_itemContext * /*ctx*/) override; + virtual void enterDeviceptr_clause(accparser::Deviceptr_clauseContext * /*ctx*/) override; virtual void enterFinalize_clause(accparser::Finalize_clauseContext * /*ctx*/) override; @@ -152,6 +138,7 @@ class OpenACCIRConstructor : public accparserBaseListener { enterGang_clause(accparser::Gang_clauseContext * /*ctx*/) override; virtual void exitGang_clause(accparser::Gang_clauseContext * /*ctx*/) override; + virtual void exitGang_arg(accparser::Gang_argContext * /*ctx*/) override; virtual void enterNohost_clause(accparser::Nohost_clauseContext * /*ctx*/) override; virtual void enterGang_no_list_clause( @@ -165,6 +152,10 @@ class OpenACCIRConstructor : public accparserBaseListener { virtual void enterIndependent_clause( accparser::Independent_clauseContext * /*ctx*/) override; virtual void + enterIndirect_clause(accparser::Indirect_clauseContext *ctx) override; + virtual void + exitIndirect_clause(accparser::Indirect_clauseContext *ctx) override; + virtual void enterLink_clause(accparser::Link_clauseContext * /*ctx*/) override; virtual void enterNo_create_clause(accparser::No_create_clauseContext * /*ctx*/) override; diff --git a/src/OpenACCIR.cpp b/src/OpenACCIR.cpp index 05d5c69..d781c49 100644 --- a/src/OpenACCIR.cpp +++ b/src/OpenACCIR.cpp @@ -6,6 +6,18 @@ // occurrences and ordering. bool OpenACCDirective::enable_clause_merging = false; +void OpenACCDataClause::addModifier(OpenACCDataClauseModifierKind modifier) { + if (modifier == ACCC_DATA_MOD_unknown) { + return; + } + for (const auto &existing : modifiers) { + if (existing == modifier) { + return; + } + } + modifiers.push_back(modifier); +} + void OpenACCClause::addLangExpr(const OpenACCExpressionItem &expression, int line, int col) { if (isClauseMergingEnabled()) { @@ -188,6 +200,10 @@ OpenACCClause *OpenACCDirective::addOpenACCClause(int k, ...) { new_clause = OpenACCNumGangsClause::addClause(this); break; } + case ACCC_indirect: { + new_clause = OpenACCIndirectClause::addClause(this); + break; + } case ACCC_link: { new_clause = OpenACCLinkClause::addClause(this); break; @@ -369,6 +385,9 @@ void OpenACCCollapseClause::mergeClause(OpenACCDirective *directive, for (auto it = current_clauses->begin(); it != current_clauses->end() - 1; ++it) { auto *existing = static_cast(*it); + if (existing->isForce() != incoming->isForce()) { + continue; + } bool merged = false; if (existing->getCounts().empty() && incoming->getCounts().empty()) { @@ -569,28 +588,25 @@ OpenACCClause *OpenACCCopyinClause::addClause(OpenACCDirective *directive) { void OpenACCCopyinClause::mergeClause(OpenACCDirective *directive, OpenACCClause *current_clause) { - // Respect the global clause merging flag if (!OpenACCDirective::getClauseMerging()) { return; } - std::vector *current_clauses = directive->getClauses(ACCC_copyin); for (std::vector::iterator it = current_clauses->begin(); it != current_clauses->end() - 1; it++) { - if (((OpenACCCopyinClause *)(*it))->getModifier() == - ((OpenACCCopyinClause *)current_clause)->getModifier() && - ((OpenACCClause *)(*it))->getOriginalKeyword() == - ((OpenACCClause *)current_clause)->getOriginalKeyword()) { - auto *existing = static_cast(*it); - auto *incoming = static_cast(current_clause); + auto *existing = static_cast(*it); + auto *incoming = static_cast(current_clause); + if (existing->getModifiers() == incoming->getModifiers() && + existing->getVariant() == incoming->getVariant()) { mergeVarList(existing, incoming); current_clauses->pop_back(); directive->getClausesInOriginalOrder()->pop_back(); + delete incoming; break; } } -}; +} OpenACCClause *OpenACCCopyoutClause::addClause(OpenACCDirective *directive) { @@ -615,28 +631,25 @@ OpenACCClause *OpenACCCopyoutClause::addClause(OpenACCDirective *directive) { void OpenACCCopyoutClause::mergeClause(OpenACCDirective *directive, OpenACCClause *current_clause) { - // Respect the global clause merging flag if (!OpenACCDirective::getClauseMerging()) { return; } - std::vector *current_clauses = directive->getClauses(ACCC_copyout); for (std::vector::iterator it = current_clauses->begin(); it != current_clauses->end() - 1; it++) { - if (((OpenACCCopyoutClause *)(*it))->getModifier() == - ((OpenACCCopyoutClause *)current_clause)->getModifier() && - ((OpenACCClause *)(*it))->getOriginalKeyword() == - ((OpenACCClause *)current_clause)->getOriginalKeyword()) { - auto *existing = static_cast(*it); - auto *incoming = static_cast(current_clause); + auto *existing = static_cast(*it); + auto *incoming = static_cast(current_clause); + if (existing->getModifiers() == incoming->getModifiers() && + existing->getVariant() == incoming->getVariant()) { mergeVarList(existing, incoming); current_clauses->pop_back(); directive->getClausesInOriginalOrder()->pop_back(); + delete incoming; break; } } -}; +} OpenACCClause *OpenACCCreateClause::addClause(OpenACCDirective *directive) { @@ -661,28 +674,61 @@ OpenACCClause *OpenACCCreateClause::addClause(OpenACCDirective *directive) { void OpenACCCreateClause::mergeClause(OpenACCDirective *directive, OpenACCClause *current_clause) { - // Respect the global clause merging flag if (!OpenACCDirective::getClauseMerging()) { return; } - std::vector *current_clauses = directive->getClauses(ACCC_create); for (std::vector::iterator it = current_clauses->begin(); it != current_clauses->end() - 1; it++) { - if (((OpenACCCreateClause *)(*it))->getModifier() == - ((OpenACCCreateClause *)current_clause)->getModifier() && - ((OpenACCClause *)(*it))->getOriginalKeyword() == - ((OpenACCClause *)current_clause)->getOriginalKeyword()) { - auto *existing = static_cast(*it); - auto *incoming = static_cast(current_clause); + auto *existing = static_cast(*it); + auto *incoming = static_cast(current_clause); + if (existing->getModifiers() == incoming->getModifiers() && + existing->getVariant() == incoming->getVariant()) { mergeVarList(existing, incoming); current_clauses->pop_back(); directive->getClausesInOriginalOrder()->pop_back(); + delete incoming; break; } } -}; +} + +OpenACCClause *OpenACCIndirectClause::addClause(OpenACCDirective *directive) { + auto *current_clauses = directive->getClauses(ACCC_indirect); + OpenACCIndirectClause *new_clause = nullptr; + if (current_clauses->empty()) { + new_clause = new OpenACCIndirectClause(); + current_clauses->push_back(new_clause); + // Ensure the map entry exists (though getClauses usually handles it, + // addOpenACCClause does the insert) + } else { + // If multiple indirect clauses are present, we might want to return the + // existing one if we support merging, or create a new one. For routine, + // usually one is enough, but to be safe and match patterns: + if (OpenACCDirective::getClauseMerging()) { + new_clause = + static_cast(current_clauses->front()); + } else { + new_clause = new OpenACCIndirectClause(); + current_clauses->push_back(new_clause); + } + } + return new_clause; +} + +void OpenACCIndirectClause::mergeClause(OpenACCDirective *directive, + OpenACCClause *clause) { + OpenACCIndirectClause *indirect_clause = + static_cast(clause); + if (indirect_clause->isPresent()) { + this->setPresent(true); + if (!indirect_clause->getValue().text.empty()) { + this->setValue(indirect_clause->getValue().text, + indirect_clause->getValue().is_string_literal); + } + } +} OpenACCClause *OpenACCNoCreateClause::addClause(OpenACCDirective *directive) { auto *all_clauses = directive->getAllClauses(); @@ -1810,24 +1856,6 @@ void OpenACCWorkerClause::mergeClause(OpenACCDirective *directive, } }; -static OpenACCDeviceTypeKind parseDeviceTypeKind(const std::string &value) { - std::string lower = value; - std::transform(lower.begin(), lower.end(), lower.begin(), ::tolower); - if (lower == "host") { - return ACCC_DEVICE_TYPE_host; - } - if (lower == "any") { - return ACCC_DEVICE_TYPE_any; - } - if (lower == "multicore") { - return ACCC_DEVICE_TYPE_multicore; - } - if (lower == "default") { - return ACCC_DEVICE_TYPE_default; - } - return ACCC_DEVICE_TYPE_unknown; -} - void OpenACCDeviceTypeClause::addDeviceType(OpenACCDeviceTypeKind kind) { if (kind == ACCC_DEVICE_TYPE_unknown) { return; @@ -1840,16 +1868,19 @@ void OpenACCDeviceTypeClause::addDeviceType(OpenACCDeviceTypeKind kind) { device_types.push_back(kind); } -void OpenACCDeviceTypeClause::addDeviceTypeString(const std::string &value) { - OpenACCDeviceTypeKind kind = parseDeviceTypeKind(value); - if (kind == ACCC_DEVICE_TYPE_unknown) { - if (std::find(unknown_types.begin(), unknown_types.end(), value) == - unknown_types.end()) { - unknown_types.push_back(value); - } - } else { - addDeviceType(kind); +void OpenACCDeviceTypeClause::addUnknownDeviceType(const std::string &value) { + if (value.empty()) { + return; + } + if (std::find(unknown_types.begin(), unknown_types.end(), value) != + unknown_types.end()) { + return; } + unknown_types.push_back(value); +} + +void OpenACCDeviceTypeClause::addDeviceTypeString(const std::string &value) { + addUnknownDeviceType(value); } OpenACCClause *OpenACCDeviceTypeClause::addClause(OpenACCDirective *directive) { diff --git a/src/OpenACCIR.h b/src/OpenACCIR.h index 4d0026c..c8ed1c8 100644 --- a/src/OpenACCIR.h +++ b/src/OpenACCIR.h @@ -61,17 +61,6 @@ class OpenACCClause : public ACC_SourceLocation { // the clause position in the vector of clauses in original order int clause_position = -1; - // Store original keyword text for alias preservation (e.g., pcreate vs create) - std::string original_keyword = ""; - - /* consider this is a struct of array, i.e. - * the expression/localtionLine/locationColumn are the same index are one - * record for an expression and its location - */ - std::vector expressions; - - std::vector locations; - public: OpenACCClause(OpenACCClauseKind k, int _line = 0, int _col = 0) : ACC_SourceLocation(_line, _col), kind(k){}; @@ -82,9 +71,15 @@ class OpenACCClause : public ACC_SourceLocation { clause_position = _clause_position; }; - void setOriginalKeyword(std::string keyword) { original_keyword = keyword; }; - std::string getOriginalKeyword() { return original_keyword; }; +protected: + /* consider this is a struct of array, i.e. + * the expression/localtionLine/locationColumn are the same index are one + * record for an expression and its location + */ + std::vector expressions; + std::vector locations; +public: // a list of expressions or variables that are language-specific for the // clause, accparser does not parse them, instead, it only stores them as // strings @@ -144,6 +139,26 @@ class OpenACCVarListClause : public OpenACCClause { } }; +// Base for OpenACC data clauses that support modifier-lists (e.g., +// copyin(always, readonly: ...)). +class OpenACCDataClause : public OpenACCVarListClause { +protected: + OpenACCDataClauseVariant variant = ACCC_DATA_COPY_unspecified; + std::vector modifiers; + +public: + OpenACCDataClause(OpenACCClauseKind k, int _line = 0, int _col = 0) + : OpenACCVarListClause(k, _line, _col) {} + + void setVariant(OpenACCDataClauseVariant v) { variant = v; } + OpenACCDataClauseVariant getVariant() const { return variant; } + + void addModifier(OpenACCDataClauseModifierKind modifier); + const std::vector &getModifiers() const { + return modifiers; + } +}; + /** * The class for all the OpenACC directives */ @@ -414,6 +429,7 @@ class OpenACCCollapseClause : public OpenACCClause { protected: std::vector counts; + bool force = false; public: OpenACCCollapseClause() : OpenACCClause(ACCC_collapse){}; @@ -426,6 +442,8 @@ class OpenACCCollapseClause : public OpenACCClause { const std::vector &getCounts() const { return counts; } + void setForce(bool f) { force = f; } + bool isForce() const { return force; } static OpenACCClause *addClause(OpenACCDirective *); std::string toString(); @@ -433,10 +451,10 @@ class OpenACCCollapseClause : public OpenACCClause { }; // Copy Clause -class OpenACCCopyClause : public OpenACCVarListClause { +class OpenACCCopyClause : public OpenACCDataClause { public: - OpenACCCopyClause() : OpenACCVarListClause(ACCC_copy) {} + OpenACCCopyClause() : OpenACCDataClause(ACCC_copy) {} static OpenACCClause *addClause(OpenACCDirective *); std::string toString(); @@ -444,19 +462,10 @@ class OpenACCCopyClause : public OpenACCVarListClause { }; // Copyin Clause -class OpenACCCopyinClause : public OpenACCVarListClause { - -protected: - OpenACCCopyinClauseModifier modifier = ACCC_COPYIN_unspecified; +class OpenACCCopyinClause : public OpenACCDataClause { public: - OpenACCCopyinClause() : OpenACCVarListClause(ACCC_copyin){}; - - OpenACCCopyinClauseModifier getModifier() { return modifier; }; - - void setModifier(OpenACCCopyinClauseModifier _modifier) { - modifier = _modifier; - }; + OpenACCCopyinClause() : OpenACCDataClause(ACCC_copyin) {} static OpenACCClause *addClause(OpenACCDirective *); std::string toString(); @@ -464,19 +473,10 @@ class OpenACCCopyinClause : public OpenACCVarListClause { }; // Copyout Clause -class OpenACCCopyoutClause : public OpenACCVarListClause { - -protected: - OpenACCCopyoutClauseModifier modifier = ACCC_COPYOUT_unspecified; +class OpenACCCopyoutClause : public OpenACCDataClause { public: - OpenACCCopyoutClause() : OpenACCVarListClause(ACCC_copyout){}; - - OpenACCCopyoutClauseModifier getModifier() { return modifier; }; - - void setModifier(OpenACCCopyoutClauseModifier _modifier) { - modifier = _modifier; - }; + OpenACCCopyoutClause() : OpenACCDataClause(ACCC_copyout) {} static OpenACCClause *addClause(OpenACCDirective *); std::string toString(); @@ -484,19 +484,10 @@ class OpenACCCopyoutClause : public OpenACCVarListClause { }; // Create Clause -class OpenACCCreateClause : public OpenACCVarListClause { - -protected: - OpenACCCreateClauseModifier modifier = ACCC_CREATE_unspecified; +class OpenACCCreateClause : public OpenACCDataClause { public: - OpenACCCreateClause() : OpenACCVarListClause(ACCC_create){}; - - OpenACCCreateClauseModifier getModifier() { return modifier; }; - - void setModifier(OpenACCCreateClauseModifier _modifier) { - modifier = _modifier; - }; + OpenACCCreateClause() : OpenACCDataClause(ACCC_create) {} static OpenACCClause *addClause(OpenACCDirective *); std::string toString(); @@ -963,6 +954,7 @@ class OpenACCDeviceTypeClause : public OpenACCClause { OpenACCDeviceTypeClause() : OpenACCClause(ACCC_device_type) {} void addDeviceType(OpenACCDeviceTypeKind kind); + void addUnknownDeviceType(const std::string &value); void addDeviceTypeString(const std::string &value); const std::vector &getDeviceTypes() const { return device_types; @@ -1001,3 +993,22 @@ class OpenACCWorkerClause : public OpenACCClause { std::string toString(); void mergeClause(OpenACCDirective *, OpenACCClause *); }; +// Indirect Clause +class OpenACCIndirectClause : public OpenACCClause { +protected: + bool is_present = false; + OpenACCIdentifier val_id; + +public: + OpenACCIndirectClause() : OpenACCClause(ACCC_indirect) {} + void setPresent(bool val) { is_present = val; } + void setValue(const std::string &val, bool is_str = false) { + val_id = {val, is_str}; + is_present = true; + } + const OpenACCIdentifier &getValue() const { return val_id; } + bool isPresent() const { return is_present; } + static OpenACCClause *addClause(OpenACCDirective *); + std::string toString(); + void mergeClause(OpenACCDirective *, OpenACCClause *); +}; diff --git a/src/OpenACCIRToString.cpp b/src/OpenACCIRToString.cpp index b44538e..15b6011 100644 --- a/src/OpenACCIRToString.cpp +++ b/src/OpenACCIRToString.cpp @@ -5,7 +5,7 @@ std::string OpenACCDirective::generatePragmaString(std::string prefix, std::string beginning_symbol, std::string ending_symbol) { - if (this->getBaseLang() == ACC_Lang_Fortran && prefix == "#pragma acc ") { + if (this->getBaseLang() == ACC_Lang_Fortran && !prefix.empty()) { prefix = "!$acc "; }; std::string result = prefix; @@ -214,6 +214,9 @@ std::string OpenACCClause::toString() { case ACCC_independent: result += "independent "; break; + case ACCC_indirect: + return static_cast(this) + ->OpenACCIndirectClause::toString(); case ACCC_link: return static_cast(this)->OpenACCLinkClause::toString(); case ACCC_nohost: @@ -288,7 +291,7 @@ static std::string deviceTypeToString(OpenACCDeviceTypeKind kind) { case ACCC_DEVICE_TYPE_host: return "host"; case ACCC_DEVICE_TYPE_any: - return "any"; + return "*"; case ACCC_DEVICE_TYPE_multicore: return "multicore"; case ACCC_DEVICE_TYPE_default: @@ -364,7 +367,11 @@ std::string OpenACCCollapseClause::toString() { for (const auto &val : vals) { result += "collapse"; - result += "(" + val.text + ") "; + result += "("; + if (isForce()) { + result += "force:"; + } + result += val.text + ") "; } return result; } @@ -489,11 +496,18 @@ std::string OpenACCGangClause::toString() { case ACCC_GANG_ARG_num: parameter_string += "num:" + it->value.text; break; + case ACCC_GANG_ARG_num_no_keyword: + parameter_string += it->value.text; + break; case ACCC_GANG_ARG_dim: parameter_string += "dim:" + it->value.text; break; case ACCC_GANG_ARG_static: - parameter_string += "static:" + it->value.text; + if (it->value.text.empty()) { + parameter_string += "static"; + } else { + parameter_string += "static:" + it->value.text; + } break; default: parameter_string += it->value.text; @@ -565,18 +579,58 @@ std::string OpenACCBindClause::toString() { std::string OpenACCCopyinClause::toString() { - std::string keyword = original_keyword.empty() ? "copyin" : original_keyword; - std::string result = keyword; - std::string parameter_string = ""; - OpenACCCopyinClauseModifier modifier = this->getModifier(); - switch (modifier) { - case ACCC_COPYIN_readonly: - parameter_string = "readonly: "; + std::string keyword = "copyin"; + switch (getVariant()) { + case ACCC_DATA_COPYIN_pcopyin: + keyword = "pcopyin"; break; - default:; - }; - parameter_string += this->varsToString(); - if (!parameter_string.empty()) { + case ACCC_DATA_COPYIN_present_or_copyin: + keyword = "present_or_copyin"; + break; + default: + break; + } + std::string result = keyword; + if (!getVars().empty()) { + std::string parameter_string; + bool first = true; + for (auto mod : getModifiers()) { + std::string name; + switch (mod) { + case ACCC_DATA_MOD_always: + name = "always"; + break; + case ACCC_DATA_MOD_alwaysin: + name = "alwaysin"; + break; + case ACCC_DATA_MOD_alwaysout: + name = "alwaysout"; + break; + case ACCC_DATA_MOD_capture: + name = "capture"; + break; + case ACCC_DATA_MOD_readonly: + name = "readonly"; + break; + case ACCC_DATA_MOD_zero: + name = "zero"; + break; + default: + name.clear(); + } + if (name.empty()) { + continue; + } + if (!first) { + parameter_string += ", "; + } + parameter_string += name; + first = false; + } + if (!parameter_string.empty()) { + parameter_string += ": "; + } + parameter_string += varsToString(); result += "(" + parameter_string + ") "; } else { result += " "; @@ -587,18 +641,58 @@ std::string OpenACCCopyinClause::toString() { std::string OpenACCCopyoutClause::toString() { - std::string keyword = original_keyword.empty() ? "copyout" : original_keyword; - std::string result = keyword; - std::string parameter_string = ""; - OpenACCCopyoutClauseModifier modifier = this->getModifier(); - switch (modifier) { - case ACCC_COPYOUT_zero: - parameter_string = "zero: "; + std::string keyword = "copyout"; + switch (getVariant()) { + case ACCC_DATA_COPYOUT_pcopyout: + keyword = "pcopyout"; break; - default:; - }; - parameter_string += this->varsToString(); - if (!parameter_string.empty()) { + case ACCC_DATA_COPYOUT_present_or_copyout: + keyword = "present_or_copyout"; + break; + default: + break; + } + std::string result = keyword; + if (!getVars().empty()) { + std::string parameter_string; + bool first = true; + for (auto mod : getModifiers()) { + std::string name; + switch (mod) { + case ACCC_DATA_MOD_always: + name = "always"; + break; + case ACCC_DATA_MOD_alwaysin: + name = "alwaysin"; + break; + case ACCC_DATA_MOD_alwaysout: + name = "alwaysout"; + break; + case ACCC_DATA_MOD_capture: + name = "capture"; + break; + case ACCC_DATA_MOD_readonly: + name = "readonly"; + break; + case ACCC_DATA_MOD_zero: + name = "zero"; + break; + default: + name.clear(); + } + if (name.empty()) { + continue; + } + if (!first) { + parameter_string += ", "; + } + parameter_string += name; + first = false; + } + if (!parameter_string.empty()) { + parameter_string += ": "; + } + parameter_string += varsToString(); result += "(" + parameter_string + ") "; } else { result += " "; @@ -609,18 +703,58 @@ std::string OpenACCCopyoutClause::toString() { std::string OpenACCCreateClause::toString() { - std::string keyword = original_keyword.empty() ? "create" : original_keyword; - std::string result = keyword; - std::string parameter_string = ""; - OpenACCCreateClauseModifier modifier = this->getModifier(); - switch (modifier) { - case ACCC_CREATE_zero: - parameter_string = "zero: "; + std::string keyword = "create"; + switch (getVariant()) { + case ACCC_DATA_CREATE_pcreate: + keyword = "pcreate"; break; - default:; - }; - parameter_string += this->varsToString(); - if (!parameter_string.empty()) { + case ACCC_DATA_CREATE_present_or_create: + keyword = "present_or_create"; + break; + default: + break; + } + std::string result = keyword; + if (!getVars().empty()) { + std::string parameter_string; + bool first = true; + for (auto mod : getModifiers()) { + std::string name; + switch (mod) { + case ACCC_DATA_MOD_always: + name = "always"; + break; + case ACCC_DATA_MOD_alwaysin: + name = "alwaysin"; + break; + case ACCC_DATA_MOD_alwaysout: + name = "alwaysout"; + break; + case ACCC_DATA_MOD_capture: + name = "capture"; + break; + case ACCC_DATA_MOD_readonly: + name = "readonly"; + break; + case ACCC_DATA_MOD_zero: + name = "zero"; + break; + default: + name.clear(); + } + if (name.empty()) { + continue; + } + if (!first) { + parameter_string += ", "; + } + parameter_string += name; + first = false; + } + if (!parameter_string.empty()) { + parameter_string += ": "; + } + parameter_string += varsToString(); result += "(" + parameter_string + ") "; } else { result += " "; @@ -649,10 +783,59 @@ varClauseToString(const std::string &keyword, } std::string OpenACCCopyClause::toString() { - std::string keyword = original_keyword.empty() ? "copy" : original_keyword; + std::string keyword = "copy"; + switch (getVariant()) { + case ACCC_DATA_COPY_pcopy: + keyword = "pcopy"; + break; + case ACCC_DATA_COPY_present_or_copy: + keyword = "present_or_copy"; + break; + default: + break; + } std::string result = keyword; if (!getVars().empty()) { - result += "(" + varsToString() + ") "; + std::string parameter_string; + bool first = true; + for (auto mod : getModifiers()) { + std::string name; + switch (mod) { + case ACCC_DATA_MOD_always: + name = "always"; + break; + case ACCC_DATA_MOD_alwaysin: + name = "alwaysin"; + break; + case ACCC_DATA_MOD_alwaysout: + name = "alwaysout"; + break; + case ACCC_DATA_MOD_capture: + name = "capture"; + break; + case ACCC_DATA_MOD_readonly: + name = "readonly"; + break; + case ACCC_DATA_MOD_zero: + name = "zero"; + break; + default: + name.clear(); + } + if (name.empty()) { + continue; + } + if (!first) { + parameter_string += ", "; + } + parameter_string += name; + first = false; + } + if (!parameter_string.empty()) { + parameter_string += ": "; + } + parameter_string += varsToString(); + result += "(" + parameter_string + ") "; } else { result += " "; } @@ -939,3 +1122,20 @@ std::string OpenACCWorkerClause::toString() { return result; }; + +std::string OpenACCIndirectClause::toString() { + std::string result = "indirect"; + if (isPresent()) { + if (!getValue().text.empty()) { + std::string val = getValue().text; + if (getValue().is_string_literal) { + val = "\"" + val + + "\""; // rough handling, relying on stored text being clean + } + result += "(" + val + ") "; + } else { + result += " "; + } + } + return result; +} diff --git a/src/OpenACCKinds.h b/src/OpenACCKinds.h index 67f818e..74a6de5 100644 --- a/src/OpenACCKinds.h +++ b/src/OpenACCKinds.h @@ -65,6 +65,7 @@ enum OpenACCClauseKind { OPENACC_CLAUSE(if) OPENACC_CLAUSE(if_present) OPENACC_CLAUSE(independent) + OPENACC_CLAUSE(indirect) OPENACC_CLAUSE(link) OPENACC_CLAUSE(nohost) OPENACC_CLAUSE(no_create) @@ -95,6 +96,19 @@ enum OpenACCClauseSeparator { ACCC_CLAUSE_SEP_comma }; +// OpenACC data clause modifiers (e.g., copyin(always, readonly: ...)). +enum OpenACCDataClauseModifierKind { +#define OPENACC_DATA_CLAUSE_MODIFIER(Name) ACCC_DATA_MOD_##Name, + OPENACC_DATA_CLAUSE_MODIFIER(always) + OPENACC_DATA_CLAUSE_MODIFIER(alwaysin) + OPENACC_DATA_CLAUSE_MODIFIER(alwaysout) + OPENACC_DATA_CLAUSE_MODIFIER(capture) + OPENACC_DATA_CLAUSE_MODIFIER(readonly) + OPENACC_DATA_CLAUSE_MODIFIER(zero) + OPENACC_DATA_CLAUSE_MODIFIER(unknown) +#undef OPENACC_DATA_CLAUSE_MODIFIER +}; + enum OpenACCDeviceTypeKind { ACCC_DEVICE_TYPE_unknown, ACCC_DEVICE_TYPE_host, @@ -130,6 +144,23 @@ enum OpenACCCopyoutClauseModifier { #undef OPENACC_COPYOUT_MODIFIER }; +// OpenACC data clause variants (e.g., copy vs pcopy vs present_or_copy) +enum OpenACCDataClauseVariant { + ACCC_DATA_COPY_unspecified, + ACCC_DATA_COPY_copy, + ACCC_DATA_COPY_pcopy, + ACCC_DATA_COPY_present_or_copy, + ACCC_DATA_COPYIN_copyin, + ACCC_DATA_COPYIN_pcopyin, + ACCC_DATA_COPYIN_present_or_copyin, + ACCC_DATA_COPYOUT_copyout, + ACCC_DATA_COPYOUT_pcopyout, + ACCC_DATA_COPYOUT_present_or_copyout, + ACCC_DATA_CREATE_create, + ACCC_DATA_CREATE_pcreate, + ACCC_DATA_CREATE_present_or_create +}; + // OpenACC attributes for 'create' clause. enum OpenACCCreateClauseModifier { #define OPENACC_CREATE_MODIFIER(Name) ACCC_CREATE_##Name, @@ -203,6 +234,7 @@ enum OpenACCAsyncModifier { enum OpenACCGangArgKind { ACCC_GANG_ARG_unknown, ACCC_GANG_ARG_num, + ACCC_GANG_ARG_num_no_keyword, ACCC_GANG_ARG_dim, ACCC_GANG_ARG_static, ACCC_GANG_ARG_other diff --git a/src/acclexer.g4 b/src/acclexer.g4 index daa798a..38fd9bc 100644 --- a/src/acclexer.g4 +++ b/src/acclexer.g4 @@ -16,9 +16,6 @@ lexer grammar acclexer; @ lexer :: postinclude { /* lexer postinclude section */ -#ifndef _WIN32 -#pragma GCC diagnostic ignored "-Wunused-parameter" -#endif } // Directly preceds the lexer class declaration in the h file (e.g. for additional types etc.). @@ -48,6 +45,8 @@ lexer grammar acclexer; int bracket_count = 0; int colon_count = 0; bool paren_processed = false; + bool data_clause_in_modifier_list = false; + bool gang_clause_allows_keyword_args = false; } // Appears in line with the other class member definitions in the cpp file. @@ -183,11 +182,27 @@ AUTO : 'auto' ; +ALWAYS + : 'always' + ; + +ALWAYSIN + : 'alwaysin' + ; + +ALWAYSOUT + : 'alwaysout' + ; + +FORCE + : 'force' + ; + BIND : 'bind' [\p{White_Space}]* { if (_input->LA(1) == '(') { - pushMode(expr_clause); + pushMode(bind_clause); } } ; @@ -197,24 +212,19 @@ CAPTURE ; COLLAPSE - : 'collapse' - { - // Allow colon-delimited modifiers like collapse(force:2) to stay inside the expression. - colon_count = 1; -} - -> pushMode (expr_clause) + : 'collapse' -> pushMode (collapse_clause) ; PCOPY - : 'pcopy' -> pushMode (expr_clause) + : 'pcopy' -> pushMode (copy_clause) ; PRESENT_OR_COPY - : 'present_or_copy' -> pushMode (expr_clause) + : 'present_or_copy' -> pushMode (copy_clause) ; COPY - : 'copy' -> pushMode (expr_clause) + : 'copy' -> pushMode (copy_clause) ; PCOPYIN @@ -282,11 +292,11 @@ DEVICE_RESIDENT ; DEVICE_TYPE - : 'device_type' -> pushMode (expr_clause) + : 'device_type' -> pushMode (device_type_clause) ; DTYPE - : 'dtype' -> type (DEVICE_TYPE) , pushMode (expr_clause) + : 'dtype' -> type (DEVICE_TYPE) , pushMode (device_type_clause) ; DEVICEPTR @@ -304,8 +314,10 @@ FIRSTPRIVATE GANG : 'gang' [\p{White_Space}]* { - if (_input->LA(1) == '(') + if (_input->LA(1) == '(') { + gang_clause_allows_keyword_args = true; pushMode(gang_clause); + } } ; @@ -329,6 +341,15 @@ LINK : 'link' -> pushMode (expr_clause) ; +INDIRECT + : 'indirect' [\p{White_Space}]* + { + if (_input->LA(1) == '(') { + pushMode(expr_clause); + } + } + ; + NOHOST : 'nohost' ; @@ -342,7 +363,12 @@ NUM ; NUM_GANGS - : 'num_gangs' -> pushMode (gang_clause) + : 'num_gangs' + { + // Uses the same lexer mode as gang(...) but only carries expressions. + gang_clause_allows_keyword_args = false; +} + -> pushMode (gang_clause) ; NUM_WORKERS @@ -378,7 +404,12 @@ SEQ ; TILE - : 'tile' -> pushMode (gang_clause) + : 'tile' + { + // Uses the same lexer mode as gang(...) but only carries expressions. + gang_clause_allows_keyword_args = false; +} + -> pushMode (gang_clause) ; USE_DEVICE @@ -507,57 +538,165 @@ ROUTINE_LINE_END : [\n\r] -> skip ; -mode copyin_clause; -COPYIN_LEFT_PAREN +mode copy_clause; +COPY_LEFT_PAREN : '(' [\p{White_Space}]* { setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("readonly") == false) { - colon_count = 0; + data_clause_in_modifier_list = + lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || + lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; pushMode(expression_mode); } } ; -COPYIN_RIGHT_PAREN - : ')' -> type (RIGHT_PAREN) , popMode +COPY_RIGHT_PAREN + : ')' + { + data_clause_in_modifier_list = false; +} + -> type (RIGHT_PAREN) , popMode ; -READONLY - : 'readonly' [\p{White_Space}]* +COPY_COLON + : ':' [\p{White_Space}]* { - if ((_input->LA(1) == ':' && _input->LA(2) == ':') || - (_input->LA(1) != ':')) { + setType(COLON); + data_clause_in_modifier_list = false; + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; +} + ; + +COPY_COMMA + : ',' [\p{White_Space}]* + { + skip(); + if (!data_clause_in_modifier_list) { + bracket_count = 0; colon_count = 1; - more(); pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; } } ; -COPYIN_COLON - : ':' [\p{White_Space}]* +COPY_ALWAYSIN + : 'alwaysin' -> type(ALWAYSIN) + ; + +COPY_ALWAYSOUT + : 'alwaysout' -> type(ALWAYSOUT) + ; + +COPY_ALWAYS + : 'always' -> type(ALWAYS) + ; + +COPY_CAPTURE + : 'capture' -> type(CAPTURE) + ; + +COPY_READONLY + : 'readonly' -> type(READONLY) + ; + +COPY_ZERO + : 'zero' -> type(ZERO) + ; + +COPY_BLANK + : [\p{White_Space}]+ -> skip + ; + +COPY_LINE_END + : [\n\r] -> skip + ; + +mode copyin_clause; +COPYIN_LEFT_PAREN + : '(' [\p{White_Space}]* { - if (_input->LA(1) == ':') - more(); - else { - colon_count = 0; - setType(COLON); + setType(LEFT_PAREN); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + data_clause_in_modifier_list = + lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || + lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; pushMode(expression_mode); } } ; -COPYIN_COMMA - : ',' [\p{White_Space}]* +COPYIN_RIGHT_PAREN + : ')' { - skip(); + data_clause_in_modifier_list = false; +} + -> type (RIGHT_PAREN) , popMode + ; + +READONLY + : 'readonly' + ; + +COPYIN_ALWAYSIN + : 'alwaysin' -> type(ALWAYSIN) + ; + +COPYIN_ALWAYSOUT + : 'alwaysout' -> type(ALWAYSOUT) + ; + +COPYIN_ALWAYS + : 'always' -> type(ALWAYS) + ; + +COPYIN_CAPTURE + : 'capture' -> type(CAPTURE) + ; + +COPYIN_ZERO + : 'zero' -> type(ZERO) + ; + +COPYIN_COLON + : ':' [\p{White_Space}]* + { + setType(COLON); + data_clause_in_modifier_list = false; + bracket_count = 0; + colon_count = 1; pushMode(expression_mode); parenthesis_global_count = 1; parenthesis_local_count = 0; - colon_count = 0; +} + ; + +COPYIN_COMMA + : ',' [\p{White_Space}]* + { + skip(); + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + } } ; @@ -576,39 +715,59 @@ COPYOUT_LEFT_PAREN setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("zero") == false) { - colon_count = 0; + data_clause_in_modifier_list = + lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || + lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; pushMode(expression_mode); } } ; COPYOUT_RIGHT_PAREN - : ')' -> type (RIGHT_PAREN) , popMode + : ')' + { + data_clause_in_modifier_list = false; +} + -> type (RIGHT_PAREN) , popMode ; ZERO - : 'zero' [\p{White_Space}]* - { - if ((_input->LA(1) == ':' && _input->LA(2) == ':') || - (_input->LA(1) != ':')) { - colon_count = 1; - more(); - pushMode(expression_mode); - } -} + : 'zero' + ; + +COPYOUT_ALWAYSIN + : 'alwaysin' -> type(ALWAYSIN) + ; + +COPYOUT_ALWAYSOUT + : 'alwaysout' -> type(ALWAYSOUT) + ; + +COPYOUT_ALWAYS + : 'always' -> type(ALWAYS) + ; + +COPYOUT_CAPTURE + : 'capture' -> type(CAPTURE) + ; + +COPYOUT_READONLY + : 'readonly' -> type(READONLY) ; COPYOUT_COLON : ':' [\p{White_Space}]* { - if (_input->LA(1) == ':') - more(); - else { - colon_count = 0; - setType(COLON); - pushMode(expression_mode); - } + setType(COLON); + data_clause_in_modifier_list = false; + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; } ; @@ -616,10 +775,13 @@ COPYOUT_COMMA : ',' [\p{White_Space}]* { skip(); - pushMode(expression_mode); - parenthesis_global_count = 1; - parenthesis_local_count = 0; - colon_count = 0; + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + } } ; @@ -638,39 +800,59 @@ CREATE_LEFT_PAREN setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("zero") == false) { + data_clause_in_modifier_list = + lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || + lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; pushMode(expression_mode); } } ; CREATE_RIGHT_PAREN - : ')' -> type (RIGHT_PAREN) , popMode + : ')' + { + data_clause_in_modifier_list = false; +} + -> type (RIGHT_PAREN) , popMode ; CREATE_ZERO - : 'zero' [\p{White_Space}]* - { - if ((_input->LA(1) == ':' && _input->LA(2) == ':') || - (_input->LA(1) != ':')) { - colon_count = 1; - more(); - pushMode(expression_mode); - } - setType(ZERO); -} + : 'zero' -> type(ZERO) + ; + +CREATE_ALWAYSIN + : 'alwaysin' -> type(ALWAYSIN) + ; + +CREATE_ALWAYSOUT + : 'alwaysout' -> type(ALWAYSOUT) + ; + +CREATE_ALWAYS + : 'always' -> type(ALWAYS) + ; + +CREATE_CAPTURE + : 'capture' -> type(CAPTURE) + ; + +CREATE_READONLY + : 'readonly' -> type(READONLY) ; CREATE_COLON : ':' [\p{White_Space}]* { - if (_input->LA(1) == ':') - more(); - else { - colon_count = 0; - setType(COLON); - pushMode(expression_mode); - } + setType(COLON); + data_clause_in_modifier_list = false; + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; } ; @@ -678,9 +860,13 @@ CREATE_COMMA : ',' [\p{White_Space}]* { skip(); - pushMode(expression_mode); - parenthesis_global_count = 1; - parenthesis_local_count = 0; + if (!data_clause_in_modifier_list) { + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + } } ; @@ -717,6 +903,49 @@ EXPR_LINE_END : [\n\r] -> skip ; +mode collapse_clause; +COLLAPSE_LEFT_PAREN + : '(' [\p{White_Space}]* + { + setType(LEFT_PAREN); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + if (lookAhead("force") == false) { + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + } +} + ; + +COLLAPSE_RIGHT_PAREN + : ')' -> type (RIGHT_PAREN) , popMode + ; + +COLLAPSE_FORCE + : 'force' -> type(FORCE) + ; + +COLLAPSE_COLON + : ':' [\p{White_Space}]* + { + setType(COLON); + bracket_count = 0; + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; +} + ; + +COLLAPSE_BLANK + : [\p{White_Space}]+ -> skip + ; + +COLLAPSE_LINE_END + : [\n\r] -> skip + ; + mode vector_clause; VECTOR_LEFT_PAREN : '(' [\p{White_Space}]* @@ -1147,15 +1376,50 @@ WORKER_LINE_END : [\n\r] -> skip ; +mode bind_clause; +BIND_LEFT_PAREN + : '(' [\p{White_Space}]* -> type(LEFT_PAREN) + ; + +BIND_RIGHT_PAREN + : ')' -> type(RIGHT_PAREN) , popMode + ; + +STRING_LITERAL + : '"' ( '""' | '\\' . | ~["\r\n] )* '"' + | '\'' ( '\'' '\'' | '\\' . | ~['\r\n] )* '\'' + ; + +BIND_NAME + : ~[ \t\f\r\n()] (~[ \t\f\r\n()])* -> type(EXPR) + ; + +BIND_BLANK + : [\p{White_Space}]+ -> skip + ; + +BIND_LINE_END + : [\n\r] -> skip + ; + mode gang_clause; GANG_LEFT_PAREN : '(' [\p{White_Space}]* { setType(LEFT_PAREN); - pushMode(expression_mode); parenthesis_global_count = 1; parenthesis_local_count = 0; - colon_count = 1; // Allow colons in gang/tile/num_gangs expressions + // If this is a gang clause with key:value args, let the key tokens match. + // Otherwise, fall back to expression capture (tile/num_gangs or positional args). + if (gang_clause_allows_keyword_args && + ((lookAhead("num") && _input->LA(4) == ':') || + (lookAhead("dim") && _input->LA(4) == ':') || + lookAhead("static"))) { + // Do not enter expression_mode yet; consume the key token first. + } else { + colon_count = 1; // Allow colons in gang/tile/num_gangs expressions + pushMode(expression_mode); + } } ; @@ -1167,10 +1431,64 @@ GANG_COMMA : ',' [\p{White_Space}]* { skip(); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + if (gang_clause_allows_keyword_args && + ((lookAhead("num") && _input->LA(4) == ':') || + (lookAhead("dim") && _input->LA(4) == ':') || + lookAhead("static"))) { + // Do not enter expression_mode yet; consume the key token first. + } else { + colon_count = 1; // Allow colons in gang/tile/num_gangs expressions + pushMode(expression_mode); + } +} + ; + +GANG_NUM + : 'num' [\p{White_Space}]* + { + if (_input->LA(1) != ':') { + // Not a key:value pair; treat as part of an expression. + setType(EXPR); + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + more(); + } else { + setType(NUM); + } +} + ; + +DIM + : 'dim' [\p{White_Space}]* + { + if (_input->LA(1) != ':') { + setType(EXPR); + colon_count = 1; + pushMode(expression_mode); + parenthesis_global_count = 1; + parenthesis_local_count = 0; + more(); + } +} + ; + +STATIC + : 'static' + ; + +GANG_COLON + : ':' [\p{White_Space}]* + { + setType(COLON); + // After key:value separator, capture the value expression as EXPR. + colon_count = 1; pushMode(expression_mode); parenthesis_global_count = 1; parenthesis_local_count = 0; - colon_count = 1; // Allow colons in gang/tile/num_gangs expressions } ; @@ -1220,6 +1538,52 @@ EXPRESSION_LEFT_PAREN } ; +mode device_type_clause; +DEVICE_TYPE_LEFT_PAREN + : '(' [\p{White_Space}]* -> type(LEFT_PAREN) + ; + +DEVICE_TYPE_RIGHT_PAREN + : ')' -> type(RIGHT_PAREN), popMode + ; + +DEVICE_TYPE_COMMA + : ',' [\p{White_Space}]* -> skip + ; + +DEVICE_TYPE_HOST + : 'host' -> type(HOST) + ; + +ANY + : 'any' -> type(ANY) + ; + +DEVICE_TYPE_ANY_STAR + : '*' -> type(ANY) + ; + +MULTICORE + : 'multicore' -> type(MULTICORE) + ; + +DEVICE_TYPE_DEFAULT + : 'default' -> type(DEFAULT) + ; + +DEVICE_TYPE_ID + : [A-Za-z_] [A-Za-z0-9_]* -> type(EXPR) + ; + +DEVICE_TYPE_BLANK + : [\p{White_Space}]+ -> skip + ; + +DEVICE_TYPE_LINE_END + : [\n\r] -> skip + ; + +mode expression_mode; EXPRESSION_RIGHT_PAREN : ')' [\p{White_Space}]* { @@ -1295,8 +1659,11 @@ EXPRESSION_CHAR // Track colon count for clause modifier separators (e.g., "readonly:") // Only track when NOT inside brackets (C array slices [1:10]) if (bracket_count == 0) { - if (colon_count == 0) { - colon_count += 1; + // Preserve C++ scope operator '::' regardless of the current colon_count. + if (_input->LA(2) == ':') { + colon_count = 1; + } else if (colon_count == 0) { + colon_count = 1; } else { colon_count = 0; } @@ -1311,3 +1678,4 @@ EXPRESSION_CHAR } } ; + diff --git a/src/accparser.g4 b/src/accparser.g4 index 9bfdf1c..461e514 100644 --- a/src/accparser.g4 +++ b/src/accparser.g4 @@ -17,9 +17,6 @@ options { tokenVocab = acclexer; } @ parser :: postinclude { /* parser postinclude section */ -#ifndef _WIN32 -#pragma GCC diagnostic ignored "-Wunused-parameter" -#endif } // Directly preceeds the parser class declaration in the h file (e.g. for additional types etc.). @@ -136,8 +133,7 @@ openacc_directive ; atomic_directive - : ATOMIC - | ATOMIC atomic_clause + : ATOMIC atomic_clause? if_clause? ; atomic_clause @@ -190,8 +186,8 @@ declare_clause_list declare_clauses : copy_clause | copyin_clause - | copyout_no_modifier_clause - | create_no_modifier_clause + | copyout_clause + | create_clause | device_resident_clause | deviceptr_clause | link_clause @@ -270,7 +266,7 @@ enter_data_clause_list enter_data_clauses : async_clause | attach_clause - | copyin_no_modifier_clause + | copyin_clause | create_clause | if_clause | wait_argument_clause @@ -286,7 +282,7 @@ exit_data_clause_list exit_data_clauses : async_clause - | copyout_no_modifier_clause + | copyout_clause | delete_clause | detach_clause | finalize_clause @@ -493,6 +489,7 @@ routine_clauses : bind_clause | device_type_clause | gang_no_list_clause + | indirect_clause | nohost_clause | seq_clause | vector_no_modifier_clause @@ -648,6 +645,7 @@ bind_clause name_or_string : EXPR + | STRING_LITERAL ; capture_clause @@ -655,50 +653,40 @@ capture_clause ; collapse_clause - : COLLAPSE LEFT_PAREN const_int RIGHT_PAREN + : COLLAPSE LEFT_PAREN (FORCE COLON)? const_int RIGHT_PAREN ; copy_clause : (PCOPY | PRESENT_OR_COPY | COPY) LEFT_PAREN var_list RIGHT_PAREN + | (PCOPY | PRESENT_OR_COPY | COPY) LEFT_PAREN data_clause_modifier_list COLON var_list RIGHT_PAREN ; copyin_clause : (PCOPYIN | PRESENT_OR_COPYIN | COPYIN) LEFT_PAREN var_list RIGHT_PAREN - | (PCOPYIN | PRESENT_OR_COPYIN | COPYIN) LEFT_PAREN copyin_clause_modifier COLON var_list RIGHT_PAREN - ; - -copyin_clause_modifier - : READONLY - ; - -copyin_no_modifier_clause - : (PCOPYIN | PRESENT_OR_COPYIN | COPYIN) LEFT_PAREN var_list RIGHT_PAREN + | (PCOPYIN | PRESENT_OR_COPYIN | COPYIN) LEFT_PAREN data_clause_modifier_list COLON var_list RIGHT_PAREN ; copyout_clause : (PCOPYOUT | PRESENT_OR_COPYOUT | COPYOUT) LEFT_PAREN var_list RIGHT_PAREN - | (PCOPYOUT | PRESENT_OR_COPYOUT | COPYOUT) LEFT_PAREN copyout_clause_modifier COLON var_list RIGHT_PAREN - ; - -copyout_clause_modifier - : ZERO - ; - -copyout_no_modifier_clause - : (PCOPYOUT | PRESENT_OR_COPYOUT | COPYOUT) LEFT_PAREN var_list RIGHT_PAREN + | (PCOPYOUT | PRESENT_OR_COPYOUT | COPYOUT) LEFT_PAREN data_clause_modifier_list COLON var_list RIGHT_PAREN ; create_clause : (PCREATE | PRESENT_OR_CREATE | CREATE) LEFT_PAREN var_list RIGHT_PAREN - | (PCREATE | PRESENT_OR_CREATE | CREATE) LEFT_PAREN create_clause_modifier COLON var_list RIGHT_PAREN + | (PCREATE | PRESENT_OR_CREATE | CREATE) LEFT_PAREN data_clause_modifier_list COLON var_list RIGHT_PAREN ; -create_clause_modifier - : ZERO +data_clause_modifier_list + : (data_clause_modifier COMMA | data_clause_modifier)+ ; -create_no_modifier_clause - : (PCREATE | PRESENT_OR_CREATE | CREATE) LEFT_PAREN var_list RIGHT_PAREN +data_clause_modifier + : ALWAYS + | ALWAYSIN + | ALWAYSOUT + | CAPTURE + | READONLY + | ZERO ; default_clause @@ -739,7 +727,15 @@ device_type_clause ; device_type_list - : (var COMMA | var)+ + : (device_type_item COMMA | device_type_item)+ + ; + +device_type_item + : HOST + | ANY + | MULTICORE + | DEFAULT + | EXPR ; deviceptr_clause @@ -760,7 +756,14 @@ gang_clause ; gang_arg_list - : (var COMMA | var)+ + : (gang_arg COMMA | gang_arg)+ + ; + +gang_arg + : NUM COLON int_expr + | DIM COLON int_expr + | STATIC (COLON int_expr)? + | int_expr ; gang_no_list_clause @@ -783,6 +786,11 @@ independent_clause : INDEPENDENT ; +indirect_clause + : INDIRECT + | INDIRECT LEFT_PAREN name_or_string RIGHT_PAREN + ; + link_clause : LINK LEFT_PAREN var_list RIGHT_PAREN ; @@ -971,4 +979,3 @@ var { cleanUp(); } - diff --git a/tests/acc_tester.cpp b/tests/acc_tester.cpp index 44b74c7..7b83a09 100644 --- a/tests/acc_tester.cpp +++ b/tests/acc_tester.cpp @@ -80,7 +80,8 @@ int main(int argc, char **argv) { savePragmaList(acc_ast_list, filename_string); for (OpenACCDirective *directive : *acc_ast_list) { - delete directive; + if (directive) + delete directive; } delete acc_ast_list; diff --git a/tests/builtin/CMakeLists.txt b/tests/builtin/CMakeLists.txt index 395fbe7..25251bf 100644 --- a/tests/builtin/CMakeLists.txt +++ b/tests/builtin/CMakeLists.txt @@ -22,6 +22,7 @@ set(TEST_FILES host_data_fortran.txt init.txt init_fortran.txt + indirect.txt loop.txt loop_fortran.txt kernels.txt diff --git a/tests/builtin/indirect.txt b/tests/builtin/indirect.txt new file mode 100644 index 0000000..a04048a --- /dev/null +++ b/tests/builtin/indirect.txt @@ -0,0 +1,6 @@ + +#pragma acc routine indirect + +#pragma acc routine indirect(foo) + +#pragma acc routine indirect(foo) diff --git a/tests/builtin/reference/ref_indirect.txt b/tests/builtin/reference/ref_indirect.txt new file mode 100644 index 0000000..8e841ab --- /dev/null +++ b/tests/builtin/reference/ref_indirect.txt @@ -0,0 +1,3 @@ +#pragma acc routine indirect +#pragma acc routine indirect(foo) +#pragma acc routine indirect(foo) From 401ffe47ab7361178902c06ad3c1792a2e7fd7b5 Mon Sep 17 00:00:00 2001 From: ouankou Date: Sun, 21 Dec 2025 21:00:52 +0000 Subject: [PATCH 2/6] Address code review comments: Fix GANG whitespace, Indirect quotes, and refactoring --- src/OpenACCASTConstructor.cpp | 22 ++++++++-------- src/acclexer.g4 | 25 ++++++++++++++++--- tests/builtin/CMakeLists.txt | 1 + tests/builtin/gang_whitespace.txt | 3 +++ .../builtin/reference/ref_gang_whitespace.txt | 3 +++ 5 files changed, 39 insertions(+), 15 deletions(-) create mode 100644 tests/builtin/gang_whitespace.txt create mode 100644 tests/builtin/reference/ref_gang_whitespace.txt diff --git a/src/OpenACCASTConstructor.cpp b/src/OpenACCASTConstructor.cpp index 990eb1c..f57a1f0 100644 --- a/src/OpenACCASTConstructor.cpp +++ b/src/OpenACCASTConstructor.cpp @@ -251,7 +251,7 @@ void OpenACCIRConstructor::exitCollapse_clause( if (ctx->FORCE()) { static_cast(current_clause)->setForce(true); } - ((OpenACCCollapseClause *)current_clause) + static_cast(current_clause) ->mergeClause(current_directive, current_clause); } @@ -496,24 +496,18 @@ void OpenACCIRConstructor::exitGang_arg(accparser::Gang_argContext *ctx) { OpenACCGangArgKind kind = ACCC_GANG_ARG_other; std::string value_text; + if (ctx->int_expr()) { + value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); + } + if (ctx->NUM()) { kind = ACCC_GANG_ARG_num; - if (ctx->int_expr()) { - value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); - } } else if (ctx->DIM()) { kind = ACCC_GANG_ARG_dim; - if (ctx->int_expr()) { - value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); - } } else if (ctx->STATIC()) { kind = ACCC_GANG_ARG_static; - if (ctx->int_expr()) { - value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); - } } else if (ctx->int_expr()) { kind = ACCC_GANG_ARG_num_no_keyword; - value_text = trimEnclosingWhiteSpace(ctx->int_expr()->getText()); } gang->addArg(kind, OpenACCExpressionItem{value_text, ACCC_CLAUSE_SEP_comma}); @@ -1027,6 +1021,12 @@ void OpenACCIRConstructor::enterIndirect_clause( std::string text = trimEnclosingWhiteSpace(ctx->name_or_string()->getText()); bool is_str = ctx->name_or_string()->STRING_LITERAL() != nullptr; + if (is_str && text.length() >= 2) { + if ((text.front() == '"' && text.back() == '"') || + (text.front() == '\'' && text.back() == '\'')) { + text = text.substr(1, text.length() - 2); + } + } indirect->setValue(text, is_str); } } diff --git a/src/acclexer.g4 b/src/acclexer.g4 index 38fd9bc..8eb0ff8 100644 --- a/src/acclexer.g4 +++ b/src/acclexer.g4 @@ -35,6 +35,23 @@ lexer grammar acclexer; } return true; } + bool lookAheadKWArgs(std::string keyword) { + size_t i; + for (i = 0; i < keyword.size(); i++) { + if (_input->LA(i + 1) != (size_t)keyword[i]) { + return false; + } + } + size_t offset = keyword.size() + 1; + while (true) { + size_t nextChar = _input->LA(offset); + if (nextChar == ' ' || nextChar == '\t' || nextChar == '\r' || nextChar == '\n') { + offset++; + } else { + return nextChar == ':'; + } + } + } } // Appears in the private part of the lexer in the h file. @@ -1412,8 +1429,8 @@ GANG_LEFT_PAREN // If this is a gang clause with key:value args, let the key tokens match. // Otherwise, fall back to expression capture (tile/num_gangs or positional args). if (gang_clause_allows_keyword_args && - ((lookAhead("num") && _input->LA(4) == ':') || - (lookAhead("dim") && _input->LA(4) == ':') || + (lookAheadKWArgs("num") || + lookAheadKWArgs("dim") || lookAhead("static"))) { // Do not enter expression_mode yet; consume the key token first. } else { @@ -1434,8 +1451,8 @@ GANG_COMMA parenthesis_global_count = 1; parenthesis_local_count = 0; if (gang_clause_allows_keyword_args && - ((lookAhead("num") && _input->LA(4) == ':') || - (lookAhead("dim") && _input->LA(4) == ':') || + (lookAheadKWArgs("num") || + lookAheadKWArgs("dim") || lookAhead("static"))) { // Do not enter expression_mode yet; consume the key token first. } else { diff --git a/tests/builtin/CMakeLists.txt b/tests/builtin/CMakeLists.txt index 25251bf..0365dcb 100644 --- a/tests/builtin/CMakeLists.txt +++ b/tests/builtin/CMakeLists.txt @@ -23,6 +23,7 @@ set(TEST_FILES init.txt init_fortran.txt indirect.txt + gang_whitespace.txt loop.txt loop_fortran.txt kernels.txt diff --git a/tests/builtin/gang_whitespace.txt b/tests/builtin/gang_whitespace.txt new file mode 100644 index 0000000..08df56b --- /dev/null +++ b/tests/builtin/gang_whitespace.txt @@ -0,0 +1,3 @@ +#pragma acc parallel loop gang(num: 1) +#pragma acc parallel loop gang(num : 1) +#pragma acc parallel loop gang( dim : 2 ) diff --git a/tests/builtin/reference/ref_gang_whitespace.txt b/tests/builtin/reference/ref_gang_whitespace.txt new file mode 100644 index 0000000..3ac1317 --- /dev/null +++ b/tests/builtin/reference/ref_gang_whitespace.txt @@ -0,0 +1,3 @@ +#pragma acc parallel loop gang(num:1) +#pragma acc parallel loop gang(num:1) +#pragma acc parallel loop gang(dim:2) From 76cacb7b4579927b30992bf85d627b0bd82c342f Mon Sep 17 00:00:00 2001 From: ouankou Date: Sun, 21 Dec 2025 22:48:17 +0000 Subject: [PATCH 3/6] Address review: Fix indirect string literal detection and restore gang test --- src/OpenACCASTConstructor.cpp | 9 +++++++++ tests/builtin/indirect.txt | 1 + tests/builtin/reference/ref_indirect.txt | 1 + 3 files changed, 11 insertions(+) diff --git a/src/OpenACCASTConstructor.cpp b/src/OpenACCASTConstructor.cpp index f57a1f0..0484851 100644 --- a/src/OpenACCASTConstructor.cpp +++ b/src/OpenACCASTConstructor.cpp @@ -1021,6 +1021,15 @@ void OpenACCIRConstructor::enterIndirect_clause( std::string text = trimEnclosingWhiteSpace(ctx->name_or_string()->getText()); bool is_str = ctx->name_or_string()->STRING_LITERAL() != nullptr; + // Fallback: IF lexer returns EXPR for a quoted string (e.g. in expr_clause mode), + // detect it manually to preserve semantics. + if (!is_str && text.length() >= 2) { + if ((text.front() == '"' && text.back() == '"') || + (text.front() == '\'' && text.back() == '\'')) { + is_str = true; + } + } + if (is_str && text.length() >= 2) { if ((text.front() == '"' && text.back() == '"') || (text.front() == '\'' && text.back() == '\'')) { diff --git a/tests/builtin/indirect.txt b/tests/builtin/indirect.txt index a04048a..bd656ef 100644 --- a/tests/builtin/indirect.txt +++ b/tests/builtin/indirect.txt @@ -4,3 +4,4 @@ #pragma acc routine indirect(foo) #pragma acc routine indirect(foo) +#pragma acc routine indirect("foo") diff --git a/tests/builtin/reference/ref_indirect.txt b/tests/builtin/reference/ref_indirect.txt index 8e841ab..3564d0e 100644 --- a/tests/builtin/reference/ref_indirect.txt +++ b/tests/builtin/reference/ref_indirect.txt @@ -1,3 +1,4 @@ #pragma acc routine indirect #pragma acc routine indirect(foo) #pragma acc routine indirect(foo) +#pragma acc routine indirect("foo") From 45dacb80d300b883f4c54ea7bd93d248ba5c7998 Mon Sep 17 00:00:00 2001 From: ouankou Date: Sun, 21 Dec 2025 23:03:41 +0000 Subject: [PATCH 4/6] Address review: guard modifier detection against prefix matches --- src/acclexer.g4 | 46 ++++++++++++------- tests/builtin/CMakeLists.txt | 1 + tests/builtin/modifier_prefix.txt | 3 ++ .../builtin/reference/ref_modifier_prefix.txt | 3 ++ 4 files changed, 37 insertions(+), 16 deletions(-) create mode 100644 tests/builtin/modifier_prefix.txt create mode 100644 tests/builtin/reference/ref_modifier_prefix.txt diff --git a/src/acclexer.g4 b/src/acclexer.g4 index 8eb0ff8..909c5dd 100644 --- a/src/acclexer.g4 +++ b/src/acclexer.g4 @@ -52,6 +52,20 @@ lexer grammar acclexer; } } } + bool lookAheadToken(std::string keyword) { + if (!lookAhead(keyword)) { + return false; + } + size_t nextChar = _input->LA(keyword.size() + 1); + // Check if next char is identifier part (alphanumeric or underscore) + if ((nextChar >= 'a' && nextChar <= 'z') || + (nextChar >= 'A' && nextChar <= 'Z') || + (nextChar >= '0' && nextChar <= '9') || + nextChar == '_') { + return false; + } + return true; + } } // Appears in the private part of the lexer in the h file. @@ -476,7 +490,7 @@ CACHE_LEFT_PAREN setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("readonly") == false) { + if (lookAheadToken("readonly") == false) { colon_count = 0; pushMode(expression_mode); } @@ -563,8 +577,8 @@ COPY_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || - lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -648,8 +662,8 @@ COPYIN_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || - lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -733,8 +747,8 @@ COPYOUT_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || - lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -818,8 +832,8 @@ CREATE_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAhead("always") || lookAhead("alwaysin") || lookAhead("alwaysout") || - lookAhead("capture") || lookAhead("readonly") || lookAhead("zero"); + lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -927,7 +941,7 @@ COLLAPSE_LEFT_PAREN setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("force") == false) { + if (lookAheadToken("force") == false) { bracket_count = 0; colon_count = 1; pushMode(expression_mode); @@ -970,7 +984,7 @@ VECTOR_LEFT_PAREN setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("length") == false) { + if (lookAheadToken("length") == false) { colon_count = 0; pushMode(expression_mode); } @@ -1271,7 +1285,7 @@ WAIT_LEFT_PAREN bracket_count = 0; parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("devnum") == false && lookAhead("queues") == false) { + if (lookAheadToken("devnum") == false && lookAheadToken("queues") == false) { colon_count = 0; pushMode(expression_mode); } @@ -1311,7 +1325,7 @@ QUEUES WAIT_COLON : ':' [\p{White_Space}]* { - if (lookAhead("queues") == false) { + if (lookAheadToken("queues") == false) { bracket_count = 0; colon_count = 0; pushMode(expression_mode); @@ -1349,7 +1363,7 @@ WORKER_LEFT_PAREN setType(LEFT_PAREN); parenthesis_global_count = 1; parenthesis_local_count = 0; - if (lookAhead("num") == false) { + if (lookAheadToken("num") == false) { pushMode(expression_mode); } } @@ -1431,7 +1445,7 @@ GANG_LEFT_PAREN if (gang_clause_allows_keyword_args && (lookAheadKWArgs("num") || lookAheadKWArgs("dim") || - lookAhead("static"))) { + lookAheadToken("static"))) { // Do not enter expression_mode yet; consume the key token first. } else { colon_count = 1; // Allow colons in gang/tile/num_gangs expressions @@ -1453,7 +1467,7 @@ GANG_COMMA if (gang_clause_allows_keyword_args && (lookAheadKWArgs("num") || lookAheadKWArgs("dim") || - lookAhead("static"))) { + lookAheadToken("static"))) { // Do not enter expression_mode yet; consume the key token first. } else { colon_count = 1; // Allow colons in gang/tile/num_gangs expressions diff --git a/tests/builtin/CMakeLists.txt b/tests/builtin/CMakeLists.txt index 0365dcb..7cb7417 100644 --- a/tests/builtin/CMakeLists.txt +++ b/tests/builtin/CMakeLists.txt @@ -24,6 +24,7 @@ set(TEST_FILES init_fortran.txt indirect.txt gang_whitespace.txt + modifier_prefix.txt loop.txt loop_fortran.txt kernels.txt diff --git a/tests/builtin/modifier_prefix.txt b/tests/builtin/modifier_prefix.txt new file mode 100644 index 0000000..0301316 --- /dev/null +++ b/tests/builtin/modifier_prefix.txt @@ -0,0 +1,3 @@ +#pragma acc data copy(alwaysVar) +#pragma acc data copyin(readonlyVar) +#pragma acc parallel loop private(forceVar) diff --git a/tests/builtin/reference/ref_modifier_prefix.txt b/tests/builtin/reference/ref_modifier_prefix.txt new file mode 100644 index 0000000..0301316 --- /dev/null +++ b/tests/builtin/reference/ref_modifier_prefix.txt @@ -0,0 +1,3 @@ +#pragma acc data copy(alwaysVar) +#pragma acc data copyin(readonlyVar) +#pragma acc parallel loop private(forceVar) From e2236e42631aa9c456dae01c0eb0c6acf239020e Mon Sep 17 00:00:00 2001 From: ouankou Date: Sun, 21 Dec 2025 23:20:44 +0000 Subject: [PATCH 5/6] Add agent config --- .gemini/config.yaml | 10 ++++++++++ AGENTS.md | 8 ++++++++ CLAUDE.md | 1 + GEMINI.md | 1 + 4 files changed, 20 insertions(+) create mode 100644 .gemini/config.yaml create mode 100644 AGENTS.md create mode 100644 CLAUDE.md create mode 100644 GEMINI.md diff --git a/.gemini/config.yaml b/.gemini/config.yaml new file mode 100644 index 0000000..a09c824 --- /dev/null +++ b/.gemini/config.yaml @@ -0,0 +1,10 @@ +have_fun: false +code_review: + disable: false + comment_severity_threshold: 'HIGH' + max_review_comments: -1 + pull_request_opened: + help: false + summary: false + code_review: false +ignore_patterns: [] diff --git a/AGENTS.md b/AGENTS.md new file mode 100644 index 0000000..96cebd0 --- /dev/null +++ b/AGENTS.md @@ -0,0 +1,8 @@ +# Repository Agent Instructions + +- Always consult the latest OpenACC 3.4 specification when updating or extending the parser to ensure semantic accuracy. +- Preserve LLVM coding style across all source files, including C, C++, ANTLR .g4 sources. +- To run the regression tests via CMake/CTest: + 1. Ensure ANTLR4 is installed and visible to CMake (e.g., `which antlr4` should succeed). + 2. Configure the build directory: `cmake -S . -B build -DCMAKE_BUILD_TYPE=Debug`. + 3. From the `build` directory invoke `ctest --output-on-failure`. diff --git a/CLAUDE.md b/CLAUDE.md new file mode 100644 index 0000000..c31b640 --- /dev/null +++ b/CLAUDE.md @@ -0,0 +1 @@ +For full project context and agent instructions, please read [AGENTS.md](AGENTS.md). \ No newline at end of file diff --git a/GEMINI.md b/GEMINI.md new file mode 100644 index 0000000..c31b640 --- /dev/null +++ b/GEMINI.md @@ -0,0 +1 @@ +For full project context and agent instructions, please read [AGENTS.md](AGENTS.md). \ No newline at end of file From ea1fa8b398084020821725ab2522f07420da541a Mon Sep 17 00:00:00 2001 From: ouankou Date: Sun, 21 Dec 2025 23:24:16 +0000 Subject: [PATCH 6/6] Address review: fix keyword-named variable regression in data clauses --- src/acclexer.g4 | 43 ++++++++++++++++---- tests/builtin/CMakeLists.txt | 1 + tests/builtin/keyword_vars.txt | 5 +++ tests/builtin/reference/ref_keyword_vars.txt | 5 +++ 4 files changed, 46 insertions(+), 8 deletions(-) create mode 100644 tests/builtin/keyword_vars.txt create mode 100644 tests/builtin/reference/ref_keyword_vars.txt diff --git a/src/acclexer.g4 b/src/acclexer.g4 index 909c5dd..d0c2f3d 100644 --- a/src/acclexer.g4 +++ b/src/acclexer.g4 @@ -66,6 +66,29 @@ lexer grammar acclexer; } return true; } + bool hasColonInClause() { + size_t offset = 1; + int p_depth = 0; + int b_depth = 0; + while (true) { + size_t c = _input->LA(offset); + if (c == 0 || c == (size_t)-1) break; // EOF + if (c == '(') p_depth++; + else if (c == ')') { + if (p_depth == 0) return false; // End of clause + p_depth--; + } + else if (c == '[') b_depth++; + else if (c == ']') { + if (b_depth > 0) b_depth--; + } + else if (c == ':') { + if (p_depth == 0 && b_depth == 0) return true; + } + offset++; + } + return false; + } } // Appears in the private part of the lexer in the h file. @@ -577,8 +600,9 @@ COPY_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || - lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); + (lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero")) && + hasColonInClause(); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -662,8 +686,9 @@ COPYIN_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || - lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); + (lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero")) && + hasColonInClause(); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -747,8 +772,9 @@ COPYOUT_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || - lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); + (lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero")) && + hasColonInClause(); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; @@ -832,8 +858,9 @@ CREATE_LEFT_PAREN parenthesis_global_count = 1; parenthesis_local_count = 0; data_clause_in_modifier_list = - lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || - lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero"); + (lookAheadToken("always") || lookAheadToken("alwaysin") || lookAheadToken("alwaysout") || + lookAheadToken("capture") || lookAheadToken("readonly") || lookAheadToken("zero")) && + hasColonInClause(); if (!data_clause_in_modifier_list) { bracket_count = 0; colon_count = 1; diff --git a/tests/builtin/CMakeLists.txt b/tests/builtin/CMakeLists.txt index 7cb7417..4b0b560 100644 --- a/tests/builtin/CMakeLists.txt +++ b/tests/builtin/CMakeLists.txt @@ -25,6 +25,7 @@ set(TEST_FILES indirect.txt gang_whitespace.txt modifier_prefix.txt + keyword_vars.txt loop.txt loop_fortran.txt kernels.txt diff --git a/tests/builtin/keyword_vars.txt b/tests/builtin/keyword_vars.txt new file mode 100644 index 0000000..646c01d --- /dev/null +++ b/tests/builtin/keyword_vars.txt @@ -0,0 +1,5 @@ +#pragma acc data copy(always) +#pragma acc data copyin(readonly) +#pragma acc data copyout(zero) +#pragma acc data copy(always, zero: x) +#pragma acc parallel loop private(always) diff --git a/tests/builtin/reference/ref_keyword_vars.txt b/tests/builtin/reference/ref_keyword_vars.txt new file mode 100644 index 0000000..646c01d --- /dev/null +++ b/tests/builtin/reference/ref_keyword_vars.txt @@ -0,0 +1,5 @@ +#pragma acc data copy(always) +#pragma acc data copyin(readonly) +#pragma acc data copyout(zero) +#pragma acc data copy(always, zero: x) +#pragma acc parallel loop private(always)