Author: erichkeane Date: 2025-06-27T10:24:38-07:00 New Revision: e4d8e06f83f866786cb5af602857f6d161501593
URL: https://github.com/llvm/llvm-project/commit/e4d8e06f83f866786cb5af602857f6d161501593 DIFF: https://github.com/llvm/llvm-project/commit/e4d8e06f83f866786cb5af602857f6d161501593.diff LOG: [OpenACC][CIR] Implement present/deviceptr/attach lowering for data These three are once again are IR clones of what the compute IR looks like, so this patch is just adding the implementation and writing sufficient tests. Added: Modified: clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp clang/test/CIR/CodeGenOpenACC/data.c Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 5652f03c92b13..b17994ee8771e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -922,7 +922,7 @@ class OpenACCClauseCIREmitter final void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) { if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp, - mlir::acc::KernelsOp>) { + mlir::acc::KernelsOp, mlir::acc::DataOp>) { for (const Expr *var : clause.getVarList()) addDataOperand<mlir::acc::DevicePtrOp>( var, mlir::acc::DataClause::acc_deviceptr, {}, @@ -932,7 +932,7 @@ class OpenACCClauseCIREmitter final applyToComputeOp(clause); } else { // TODO: When we've implemented this for everything, switch this to an - // unreachable. data, declare remain. + // unreachable. declare remains. return clauseNotImplemented(clause); } } @@ -953,7 +953,7 @@ class OpenACCClauseCIREmitter final void VisitPresentClause(const OpenACCPresentClause &clause) { if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp, - mlir::acc::KernelsOp>) { + mlir::acc::KernelsOp, mlir::acc::DataOp>) { for (const Expr *var : clause.getVarList()) addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>( var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true, @@ -962,14 +962,14 @@ class OpenACCClauseCIREmitter final applyToComputeOp(clause); } else { // TODO: When we've implemented this for everything, switch this to an - // unreachable. data & declare remain. + // unreachable. declare remains. return clauseNotImplemented(clause); } } void VisitAttachClause(const OpenACCAttachClause &clause) { if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp, - mlir::acc::KernelsOp>) { + mlir::acc::KernelsOp, mlir::acc::DataOp>) { for (const Expr *var : clause.getVarList()) addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>( var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true, @@ -978,7 +978,7 @@ class OpenACCClauseCIREmitter final applyToComputeOp(clause); } else { // TODO: When we've implemented this for everything, switch this to an - // unreachable. data, enter data remain. + // unreachable. enter data remains. return clauseNotImplemented(clause); } } diff --git a/clang/test/CIR/CodeGenOpenACC/data.c b/clang/test/CIR/CodeGenOpenACC/data.c index 77a373f9c049f..1f6a76ce1ea7c 100644 --- a/clang/test/CIR/CodeGenOpenACC/data.c +++ b/clang/test/CIR/CodeGenOpenACC/data.c @@ -3,6 +3,9 @@ void acc_data(int cond) { // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}) { // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init] + + int *ptr; + // CHECK-NEXT: %[[PTR:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptr"] // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i> #pragma acc data default(none) @@ -221,5 +224,50 @@ void acc_data(int cond) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>} +#pragma acc data deviceptr(ptr) + {} + // CHECK-NEXT: %[[DEV_PTR:.*]] = acc.deviceptr varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"} + // CHECK-NEXT: acc.data dataOperands(%[[DEV_PTR]] : !cir.ptr<!cir.ptr<!s32i>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc +#pragma acc data deviceptr(ptr) device_type(radeon) async + {} + // CHECK-NEXT: %[[DEV_PTR:.*]] = acc.deviceptr varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"} + // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[DEV_PTR]] : !cir.ptr<!cir.ptr<!s32i>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + +#pragma acc data present(cond) + {} + // CHECK-NEXT: %[[PRESENT:.*]] = acc.present varPtr(%[[COND]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "cond"} + // CHECK-NEXT: acc.data dataOperands(%[[PRESENT]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.delete accPtr(%[[PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "cond"} + +#pragma acc data present(cond) device_type(radeon) async + {} + // CHECK-NEXT: %[[PRESENT:.*]] = acc.present varPtr(%[[COND]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) -> !cir.ptr<!s32i> {name = "cond"} + // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[PRESENT]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.delete accPtr(%[[PRESENT]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_present>, name = "cond"} + +#pragma acc data attach(ptr) + {} + // CHECK-NEXT: %[[ATTACH:.*]] = acc.attach varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"} + // CHECK-NEXT: acc.data dataOperands(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.detach accPtr(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_attach>, name = "ptr"} + +#pragma acc data attach(ptr) device_type(radeon) async + {} + // CHECK-NEXT: %[[ATTACH:.*]] = acc.attach varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"} + // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.detach accPtr(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_attach>, name = "ptr"} + // CHECK-NEXT: cir.return } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
