diff options
| -rw-r--r-- | clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 6 | ||||
| -rw-r--r-- | clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c | 28 |
2 files changed, 30 insertions, 4 deletions
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 1a07898..5652f03 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -939,7 +939,7 @@ public: 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 @@ public: } 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 7a541a8..068ca4a 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"} } |