llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) <details> <summary>Changes</summary> The 'update' construct has 3 'var-list' clauses, device, self, and host. Each has a pretty simple data-operand type syntax in the IR, so this patch implements them as well. At least one of those is required to be present on an 'update', so we cannot do any lowering without them. Note that 'self' and 'host' are aliases. --- Full diff: https://github.com/llvm/llvm-project/pull/146378.diff 3 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+30-3) - (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+4-2) - (added) clang/test/CIR/CodeGenOpenACC/update.c (+67) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index cc0f3b77c1a65..b7a73e2f62945 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -613,12 +613,39 @@ class OpenACCClauseCIREmitter final } else { llvm_unreachable("var-list version of self shouldn't get here"); } + } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) { + assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() && + "var-list version of self required for update"); + for (const Expr *var : clause.getVarList()) + addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>( + var, mlir::acc::DataClause::acc_update_self, {}, + /*structured=*/false, /*implicit=*/false); } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); } else { - // TODO: When we've implemented this for everything, switch this to an - // unreachable. update construct remains. - return clauseNotImplemented(clause); + llvm_unreachable("Unknown construct kind in VisitSelfClause"); + } + } + + void VisitHostClause(const OpenACCHostClause &clause) { + if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) { + for (const Expr *var : clause.getVarList()) + addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>( + var, mlir::acc::DataClause::acc_update_host, {}, + /*structured=*/false, /*implicit=*/false); + } else { + llvm_unreachable("Unknown construct kind in VisitHostClause"); + } + } + + void VisitDeviceClause(const OpenACCDeviceClause &clause) { + if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) { + for (const Expr *var : clause.getVarList()) + addDataOperand<mlir::acc::UpdateDeviceOp>( + var, mlir::acc::DataClause::acc_update_device, {}, + /*structured=*/false, /*implicit=*/false); + } else { + llvm_unreachable("Unknown construct kind in VisitDeviceClause"); } } @@ -1095,6 +1122,7 @@ EXPL_SPEC(mlir::acc::WaitOp) EXPL_SPEC(mlir::acc::HostDataOp) EXPL_SPEC(mlir::acc::EnterDataOp) EXPL_SPEC(mlir::acc::ExitDataOp) +EXPL_SPEC(mlir::acc::UpdateOp) #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 f3a635b7c83eb..5993056bf06ba 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -266,8 +266,10 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct( mlir::LogicalResult CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) { - cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct"); - return mlir::failure(); + mlir::Location start = getLoc(s.getSourceRange().getBegin()); + emitOpenACCOp<UpdateOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(), + s.clauses()); + return mlir::success(); } mlir::LogicalResult CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) { diff --git a/clang/test/CIR/CodeGenOpenACC/update.c b/clang/test/CIR/CodeGenOpenACC/update.c new file mode 100644 index 0000000000000..4e25a1df2a42b --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/update.c @@ -0,0 +1,67 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +void acc_update(int parmVar, int *ptrParmVar) { + // CHECK: cir.func{{.*}} @acc_update(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) { + // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init] + // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i> + // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>> + +#pragma acc update device(parmVar) + // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]] : !cir.ptr<!s32i>) + +#pragma acc update device(parmVar, ptrParmVar) + // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false} + // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + +#pragma acc update device(parmVar) device(ptrParmVar) + // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false} + // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + +#pragma acc update host(parmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false} + +#pragma acc update host(parmVar, ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false} + +#pragma acc update host(parmVar) host(ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false} + +#pragma acc update self(parmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + +#pragma acc update self(parmVar, ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + +#pragma acc update self(parmVar) self(ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + +#pragma acc update self(parmVar) device(ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} + // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false} +} `````````` </details> https://github.com/llvm/llvm-project/pull/146378 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
