llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) <details> <summary>Changes</summary> 'host_data' has its own Op kind, so this handles the lowering there, it looks exactly like the other ones we've done so far, so nothing novel here. host_data takes 3 clauses, 1 of which is required. 'use_device' is required, and results in an acc.use_device operation, which then feeds into the dataOperands list on acc.host_data. 'if_present' is a simple attribute on the operand. 'if' is a condition on the operand, identical to our other handling of 'if'. This patch handles all of these. --- Full diff: https://github.com/llvm/llvm-project/pull/143136.diff 3 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+40-1) - (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+11-5) - (added) clang/test/CIR/CodeGenOpenACC/host_data.c (+55) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index f41f776225152..e3657e9014121 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -319,6 +319,21 @@ class OpenACCClauseCIREmitter final dataOperands.push_back(afterOp.getOperation()); } + template <typename BeforeOpTy> + void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause, + bool structured, bool implicit) { + DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand); + auto beforeOp = + builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured, + implicit, opInfo.name, opInfo.bounds); + operation.getDataClauseOperandsMutable().append(beforeOp.getResult()); + + // Set the 'rest' of the info for the operation. + beforeOp.setDataClause(dataClause); + // Make sure we record these, so 'async' values can be updated later. + dataOperands.push_back(beforeOp.getOperation()); + } + // Helper function that covers for the fact that we don't have this function // on all operation types. mlir::ArrayAttr getAsyncOnlyAttr() { @@ -550,7 +565,8 @@ class OpenACCClauseCIREmitter final if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp, mlir::acc::KernelsOp, mlir::acc::InitOp, mlir::acc::ShutdownOp, mlir::acc::SetOp, - mlir::acc::DataOp, mlir::acc::WaitOp>) { + mlir::acc::DataOp, mlir::acc::WaitOp, + mlir::acc::HostDataOp>) { operation.getIfCondMutable().append( createCondition(clause.getConditionExpr())); } else if constexpr (isCombinedType<OpTy>) { @@ -566,6 +582,17 @@ class OpenACCClauseCIREmitter final } } + void VisitIfPresentClause(const OpenACCIfPresentClause &clause) { + if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) { + operation.setIfPresent(true); + } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) { + // Last unimplemented one here, so just put it in this way instead. + return clauseNotImplemented(clause); + } else { + llvm_unreachable("unknown construct kind in VisitIfPresentClause"); + } + } + void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) { if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp, mlir::acc::SetOp>) { @@ -791,6 +818,17 @@ class OpenACCClauseCIREmitter final return clauseNotImplemented(clause); } } + + void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) { + if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) { + for (auto var : clause.getVarList()) + addDataOperand<mlir::acc::UseDeviceOp>( + var, mlir::acc::DataClause::acc_use_device, + /*structured=*/true, /*implicit=*/false); + } else { + llvm_unreachable("Unknown construct kind in VisitUseDeviceClause"); + } + } }; template <typename OpTy> @@ -826,6 +864,7 @@ EXPL_SPEC(mlir::acc::InitOp) EXPL_SPEC(mlir::acc::ShutdownOp) EXPL_SPEC(mlir::acc::SetOp) EXPL_SPEC(mlir::acc::WaitOp) +EXPL_SPEC(mlir::acc::HostDataOp) #undef EXPL_SPEC template <typename ComputeOp, typename LoopOp> diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index d922ca0c74d5d..2aab9cecf93d8 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -235,6 +235,17 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct( llvm_unreachable("invalid compute construct kind"); } } + +mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct( + const OpenACCHostDataConstruct &s) { + mlir::Location start = getLoc(s.getSourceRange().getBegin()); + mlir::Location end = getLoc(s.getSourceRange().getEnd()); + + return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>( + start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(), + s.getStructuredBlock()); +} + mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct( const OpenACCEnterDataConstruct &s) { cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct"); @@ -245,11 +256,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct( cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct"); return mlir::failure(); } -mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct( - const OpenACCHostDataConstruct &s) { - cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct"); - return mlir::failure(); -} mlir::LogicalResult CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) { cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct"); diff --git a/clang/test/CIR/CodeGenOpenACC/host_data.c b/clang/test/CIR/CodeGenOpenACC/host_data.c new file mode 100644 index 0000000000000..4c3f7dd092a2f --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/host_data.c @@ -0,0 +1,55 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +void acc_host_data(int cond, int var1, int var2) { + // CHECK: cir.func @acc_host_data(%[[ARG_COND:.*]]: !s32i {{.*}}, %[[ARG_V1:.*]]: !s32i {{.*}}, %[[ARG_V2:.*]]: !s32i {{.*}}) { + // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init] + // CHECK-NEXT: %[[V1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var1", init] + // CHECK-NEXT: %[[V2:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var2", init] + // CHECK-NEXT: cir.store %[[ARG_COND]], %[[COND]] : !s32i, !cir.ptr<!s32i> + // CHECK-NEXT: cir.store %[[ARG_V1]], %[[V1]] : !s32i, !cir.ptr<!s32i> + // CHECK-NEXT: cir.store %[[ARG_V2]], %[[V2]] : !s32i, !cir.ptr<!s32i> + +#pragma acc host_data use_device(var1) + {} + // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"} + // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc +#pragma acc host_data use_device(var1, var2) + {} + // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"} + // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"} + // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + +#pragma acc host_data use_device(var1, var2) if_present + {} + // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"} + // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"} + // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } attributes {ifPresent} + +#pragma acc host_data use_device(var1, var2) if(cond) + {} + // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"} + // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"} + // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool + // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1 + // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + +#pragma acc host_data use_device(var1, var2) if(cond) if_present + {} + // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"} + // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"} + // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i + // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool + // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1 + // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } attributes {ifPresent} +} `````````` </details> https://github.com/llvm/llvm-project/pull/143136 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
