diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index 3687af76a559f..681567228cbb0 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -430,6 +430,11 @@ class OpenACCSelfClause final } bool isConditionExprClause() const { return HasConditionExpr.has_value(); } + bool isVarListClause() const { return !isConditionExprClause(); } + bool isEmptySelfClause() const { + return (isConditionExprClause() && !hasConditionExpr()) || + (!isConditionExprClause() && getVarList().empty()); + } bool hasConditionExpr() const { assert(HasConditionExpr.has_value() && diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index 152f996ed0fed..3bcc6f908a841 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -32,46 +32,51 @@ constexpr bool isOneOfTypes = template constexpr bool isOneOfTypes = std::is_same_v; +template class OpenACCClauseCIREmitter final - : public OpenACCClauseVisitor { - CIRGenModule &cgm; + : public OpenACCClauseVisitor> { + OpTy &operation; + CIRGenFunction &cgf; + CIRGenBuilderTy &builder; + // This is necessary since a few of the clauses emit differently based on the // directive kind they are attached to. OpenACCDirectiveKind dirKind; + // TODO(cir): This source location should be able to go away once the NYI + // diagnostics are gone. SourceLocation dirLoc; - struct AttributeData { - // Value of the 'default' attribute, added on 'data' and 'compute'/etc - // constructs as a 'default-attr'. - std::optional defaultVal = std::nullopt; - // For directives that have their device type architectures listed in - // attributes (init/shutdown/etc), the list of architectures to be emitted. - llvm::SmallVector deviceTypeArchs{}; - } attrData; - void clauseNotImplemented(const OpenACCClause &c) { - cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind()); + cgf.cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind()); } public: - OpenACCClauseCIREmitter(CIRGenModule &cgm, OpenACCDirectiveKind dirKind, - SourceLocation dirLoc) - : cgm(cgm), dirKind(dirKind), dirLoc(dirLoc) {} + OpenACCClauseCIREmitter(OpTy &operation, CIRGenFunction &cgf, + CIRGenBuilderTy &builder, + OpenACCDirectiveKind dirKind, SourceLocation dirLoc) + : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind), + dirLoc(dirLoc) {} void VisitClause(const OpenACCClause &clause) { clauseNotImplemented(clause); } void VisitDefaultClause(const OpenACCDefaultClause &clause) { - switch (clause.getDefaultClauseKind()) { - case OpenACCDefaultClauseKind::None: - attrData.defaultVal = ClauseDefaultValue::None; - break; - case OpenACCDefaultClauseKind::Present: - attrData.defaultVal = ClauseDefaultValue::Present; - break; - case OpenACCDefaultClauseKind::Invalid: - break; + // This type-trait checks if 'op'(the first arg) is one of the mlir::acc + // operations listed in the rest of the arguments. + if constexpr (isOneOfTypes) { + switch (clause.getDefaultClauseKind()) { + case OpenACCDefaultClauseKind::None: + operation.setDefaultAttr(ClauseDefaultValue::None); + break; + case OpenACCDefaultClauseKind::Present: + operation.setDefaultAttr(ClauseDefaultValue::Present); + break; + case OpenACCDefaultClauseKind::Invalid: + break; + } + } else { + return clauseNotImplemented(clause); } } @@ -89,64 +94,70 @@ class OpenACCClauseCIREmitter final } void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) { + if constexpr (isOneOfTypes) { + llvm::SmallVector deviceTypes; + std::optional existingDeviceTypes = + operation.getDeviceTypes(); + + // Ensure we keep the existing ones, and in the correct 'new' order. + if (existingDeviceTypes) { + for (const mlir::Attribute &Attr : *existingDeviceTypes) + deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get( + builder.getContext(), + cast(Attr).getValue())); + } - switch (dirKind) { - case OpenACCDirectiveKind::Init: - case OpenACCDirectiveKind::Set: - case OpenACCDirectiveKind::Shutdown: { - // Device type has a list that is either a 'star' (emitted as 'star'), - // or an identifer list, all of which get added for attributes. - - for (const DeviceTypeArgument &arg : clause.getArchitectures()) - attrData.deviceTypeArchs.push_back(decodeDeviceType(arg.first)); - break; - } - default: + for (const DeviceTypeArgument &arg : clause.getArchitectures()) { + deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get( + builder.getContext(), decodeDeviceType(arg.first))); + } + operation.removeDeviceTypesAttr(); + operation.setDeviceTypesAttr( + mlir::ArrayAttr::get(builder.getContext(), deviceTypes)); + } else if constexpr (isOneOfTypes) { + assert(!operation.getDeviceTypeAttr() && "already have device-type?"); + assert(clause.getArchitectures().size() <= 1); + + if (!clause.getArchitectures().empty()) + operation.setDeviceType( + decodeDeviceType(clause.getArchitectures()[0].first)); + } else { return clauseNotImplemented(clause); } } - // Apply any of the clauses that resulted in an 'attribute'. - template - void applyAttributes(CIRGenBuilderTy &builder, Op &op) { - - if (attrData.defaultVal.has_value()) { - // FIXME: OpenACC: as we implement this for other directive kinds, we have - // to expand this list. - // This type-trait checks if 'op'(the first arg) is one of the mlir::acc - // operations listed in the rest of the arguments. - if constexpr (isOneOfTypes) - op.setDefaultAttr(*attrData.defaultVal); - else - cgm.errorNYI(dirLoc, "OpenACC 'default' clause lowering for ", dirKind); - } - - if (!attrData.deviceTypeArchs.empty()) { - // FIXME: OpenACC: as we implement this for other directive kinds, we have - // to expand this list, or more likely, have a 'noop' branch as most other - // uses of this apply to the operands instead. - // This type-trait checks if 'op'(the first arg) is one of the mlir::acc - if constexpr (isOneOfTypes) { - llvm::SmallVector deviceTypes; - for (mlir::acc::DeviceType DT : attrData.deviceTypeArchs) - deviceTypes.push_back( - mlir::acc::DeviceTypeAttr::get(builder.getContext(), DT)); - - op.setDeviceTypesAttr( - mlir::ArrayAttr::get(builder.getContext(), deviceTypes)); - } else if constexpr (isOneOfTypes) { - assert(attrData.deviceTypeArchs.size() <= 1 && - "Set can only have a single architecture"); - if (!attrData.deviceTypeArchs.empty()) - op.setDeviceType(attrData.deviceTypeArchs[0]); + void VisitSelfClause(const OpenACCSelfClause &clause) { + if constexpr (isOneOfTypes) { + if (clause.isEmptySelfClause()) { + operation.setSelfAttr(true); + } else if (clause.isConditionExprClause()) { + assert(clause.hasConditionExpr()); + mlir::Value condition = + cgf.evaluateExprAsBool(clause.getConditionExpr()); + + mlir::Location exprLoc = + cgf.cgm.getLoc(clause.getConditionExpr()->getBeginLoc()); + mlir::IntegerType targetType = mlir::IntegerType::get( + &cgf.getMLIRContext(), /*width=*/1, + mlir::IntegerType::SignednessSemantics::Signless); + auto conversionOp = builder.create( + exprLoc, targetType, condition); + operation.getSelfCondMutable().append(conversionOp.getResult(0)); } else { - cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ", - dirKind); + llvm_unreachable("var-list version of self shouldn't get here"); } + } else { + return clauseNotImplemented(clause); } } }; +template +auto makeClauseEmitter(OpTy &op, CIRGenFunction &cgf, CIRGenBuilderTy &builder, + OpenACCDirectiveKind dirKind, SourceLocation dirLoc) { + return OpenACCClauseCIREmitter(op, cgf, builder, dirKind, dirLoc); +} + } // namespace template @@ -158,24 +169,27 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt( llvm::SmallVector retTy; llvm::SmallVector operands; - - // Clause-emitter must be here because it might modify operands. - OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc); - clauseEmitter.VisitClauseList(clauses); - auto op = builder.create(start, retTy, operands); - // Apply the attributes derived from the clauses. - clauseEmitter.applyAttributes(builder, op); + { + mlir::OpBuilder::InsertionGuard guardCase(builder); + // Sets insertion point before the 'op', since every new expression needs to + // be before the operation. + builder.setInsertionPoint(op); + makeClauseEmitter(op, *this, builder, dirKind, dirLoc) + .VisitClauseList(clauses); + } - mlir::Block &block = op.getRegion().emplaceBlock(); - mlir::OpBuilder::InsertionGuard guardCase(builder); - builder.setInsertionPointToEnd(&block); + { + mlir::Block &block = op.getRegion().emplaceBlock(); + mlir::OpBuilder::InsertionGuard guardCase(builder); + builder.setInsertionPointToEnd(&block); - LexicalScope ls{*this, start, builder.getInsertionBlock()}; - res = emitStmt(associatedStmt, /*useCurrentScope=*/true); + LexicalScope ls{*this, start, builder.getInsertionBlock()}; + res = emitStmt(associatedStmt, /*useCurrentScope=*/true); - builder.create(end); + builder.create(end); + } return res; } @@ -187,14 +201,16 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOp( llvm::SmallVector retTy; llvm::SmallVector operands; - - // Clause-emitter must be here because it might modify operands. - OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc); - clauseEmitter.VisitClauseList(clauses); - auto op = builder.create(start, retTy, operands); - // Apply the attributes derived from the clauses. - clauseEmitter.applyAttributes(builder, op); + + { + mlir::OpBuilder::InsertionGuard guardCase(builder); + // Sets insertion point before the 'op', since every new expression needs to + // be before the operation. + builder.setInsertionPoint(op); + makeClauseEmitter(op, *this, builder, dirKind, dirLoc) + .VisitClauseList(clauses); + } return res; } @@ -254,46 +270,46 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct( mlir::LogicalResult CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Loop Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct( const OpenACCCombinedConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Combined Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Combined Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct( const OpenACCEnterDataConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC EnterData Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct( const OpenACCExitDataConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC ExitData Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct( const OpenACCHostDataConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC HostData Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Wait Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Wait Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Update Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Atomic Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Cache Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Cache Construct"); return mlir::failure(); } diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c index 0c950fe3d0f9c..934daf9e8ecc0 100644 --- a/clang/test/CIR/CodeGenOpenACC/kernels.c +++ b/clang/test/CIR/CodeGenOpenACC/kernels.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_kernels(void) { - // CHECK: cir.func @acc_kernels() { +void acc_kernels(int cond) { + // CHECK: cir.func @acc_kernels(%[[ARG:.*]]: !s32i{{.*}}) { + // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr, ["cond", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr #pragma acc kernels {} @@ -38,5 +40,29 @@ void acc_kernels(void) { // CHECK-NEXT: acc.terminator // CHECK-NEXT:} +#pragma acc kernels self + {} + // CHECK-NEXT: acc.kernels { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } attributes {selfAttr} + +#pragma acc kernels self(cond) + {} + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr, !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + +#pragma acc kernels self(0) + {} + // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: cir.return } diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c index e18270435460c..c7a4bda6faa74 100644 --- a/clang/test/CIR/CodeGenOpenACC/parallel.c +++ b/clang/test/CIR/CodeGenOpenACC/parallel.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_parallel(void) { - // CHECK: cir.func @acc_parallel() { +void acc_parallel(int cond) { + // CHECK: cir.func @acc_parallel(%[[ARG:.*]]: !s32i{{.*}}) { + // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr, ["cond", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr #pragma acc parallel {} // CHECK-NEXT: acc.parallel { @@ -37,5 +39,29 @@ void acc_parallel(void) { // CHECK-NEXT: acc.yield // CHECK-NEXT:} +#pragma acc parallel self + {} + // CHECK-NEXT: acc.parallel { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } attributes {selfAttr} + +#pragma acc parallel self(cond) + {} + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr, !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + +#pragma acc parallel self(0) + {} + // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: cir.return } diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c index 72a0995549da3..38a38ad6c9514 100644 --- a/clang/test/CIR/CodeGenOpenACC/serial.c +++ b/clang/test/CIR/CodeGenOpenACC/serial.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_serial(void) { - // CHECK: cir.func @acc_serial() { +void acc_serial(int cond) { + // CHECK: cir.func @acc_serial(%[[ARG:.*]]: !s32i{{.*}}) { + // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr, ["cond", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr #pragma acc serial {} @@ -38,5 +40,29 @@ void acc_serial(void) { // CHECK-NEXT: acc.yield // CHECK-NEXT:} +#pragma acc serial self + {} + // CHECK-NEXT: acc.serial { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } attributes {selfAttr} + +#pragma acc serial self(cond) + {} + // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr, !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + +#pragma acc serial self(0) + {} + // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: cir.return }