Author: erichkeane Date: 2025-06-27T09:27:59-07:00 New Revision: 2557f9946373148664ce90dc6da5a6677424ac2c
URL: https://github.com/llvm/llvm-project/commit/2557f9946373148664ce90dc6da5a6677424ac2c DIFF: https://github.com/llvm/llvm-project/commit/2557f9946373148664ce90dc6da5a6677424ac2c.diff LOG: [OpenACC][CIR] Implement 'no_create' lowering for data This lowering ends up being identical to 'create', except it is a acc.nocreate for the start operation, and it doesn't permit modifier list. This patch implements this by adding it to the list of permitted handlers (along with compute), plus adds tests. Added: Modified: clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 1a078981a3a05..5652f03c92b13 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -939,7 +939,7 @@ class OpenACCClauseCIREmitter final void VisitNoCreateClause(const OpenACCNoCreateClause &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::NoCreateOp, mlir::acc::DeleteOp>( var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true, @@ -947,9 +947,7 @@ class OpenACCClauseCIREmitter final } else if constexpr (isCombinedType<OpTy>) { applyToComputeOp(clause); } else { - // TODO: When we've implemented this for everything, switch this to an - // unreachable. data remains. - return clauseNotImplemented(clause); + llvm_unreachable("Unknown construct kind in VisitNoCreateClause"); } } diff --git a/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c b/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c index 7a541a8c62ea6..068ca4a63e331 100644 --- a/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c +++ b/clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c @@ -187,4 +187,32 @@ void acc_data(int parmVar) { // CHECK-NEXT: } loc // CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "parmVar"} // CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"} + +#pragma acc data no_create(parmVar) + ; + // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"} + // CHECK-NEXT: acc.data dataOperands(%[[NOCREATE1]] : !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "parmVar"} + +#pragma acc data no_create(parmVar) no_create(localVar1) + ; + // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"} + // CHECK-NEXT: %[[NOCREATE2:.*]] = acc.nocreate varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"} + // CHECK-NEXT: acc.data dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "localVar1"} + // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "parmVar"} + +#pragma acc data no_create(parmVar, localVar1) + ; + // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"} + // CHECK-NEXT: %[[NOCREATE2:.*]] = acc.nocreate varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"} + // CHECK-NEXT: acc.data dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "localVar1"} + // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_no_create>, name = "parmVar"} } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
