From 125dbe103e4fdff8b23bfb4b76ad960f0069f63e Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 30 Jun 2025 11:24:17 -0700 Subject: [PATCH] [OpenACC][CIR] 'update' construct lowering + a few clauses (#146378) 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. --- clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 34 +++++++++- clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 6 +- clang/test/CIR/CodeGenOpenACC/update.c | 67 +++++++++++++++++++ 3 files changed, 102 insertions(+), 5 deletions(-) create mode 100644 clang/test/CIR/CodeGenOpenACC/update.c diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index cc0f3b77c1a6..b7a73e2f6294 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -613,12 +613,39 @@ public: } else { llvm_unreachable("var-list version of self shouldn't get here"); } + } else if constexpr (isOneOfTypes) { + assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() && + "var-list version of self required for update"); + for (const Expr *var : clause.getVarList()) + addDataOperand( + var, mlir::acc::DataClause::acc_update_self, {}, + /*structured=*/false, /*implicit=*/false); } else if constexpr (isCombinedType) { 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) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + 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) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + 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 diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index f3a635b7c83e..5993056bf06b 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(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 000000000000..4e25a1df2a42 --- /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{{.*}}) { + // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr, ["parmVar", init] + // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["ptrParmVar", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr + // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr, !cir.ptr> + +#pragma acc update device(parmVar) + // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]] : !cir.ptr) + +#pragma acc update device(parmVar, ptrParmVar) + // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr, !cir.ptr>) + +#pragma acc update device(parmVar) device(ptrParmVar) + // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr, !cir.ptr>) + +#pragma acc update host(parmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {name = "parmVar", structured = false} + +#pragma acc update host(parmVar, ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr>) to varPtr(%[[PTRPARM]] : !cir.ptr>) {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {name = "parmVar", structured = false} + +#pragma acc update host(parmVar) host(ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr>) to varPtr(%[[PTRPARM]] : !cir.ptr>) {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {name = "parmVar", structured = false} + +#pragma acc update self(parmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {dataClause = #acc, name = "parmVar", structured = false} + +#pragma acc update self(parmVar, ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr>) to varPtr(%[[PTRPARM]] : !cir.ptr>) {dataClause = #acc, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {dataClause = #acc, name = "parmVar", structured = false} + +#pragma acc update self(parmVar) self(ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr>) to varPtr(%[[PTRPARM]] : !cir.ptr>) {dataClause = #acc, name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {dataClause = #acc, name = "parmVar", structured = false} + +#pragma acc update self(parmVar) device(ptrParmVar) + // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[UPD_DEV2]] : !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr) to varPtr(%[[PARM]] : !cir.ptr) {dataClause = #acc, name = "parmVar", structured = false} +}