diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 8b61d3fae3ad0..f41f776225152 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -14,6 +14,8 @@ #include "CIRGenFunction.h" +#include "clang/AST/ExprCXX.h" + #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/OpenACC/OpenACC.h" #include "llvm/ADT/TypeSwitch.h" @@ -188,7 +190,7 @@ class OpenACCClauseCIREmitter final struct DataOperandInfo { mlir::Location beginLoc; mlir::Value varValue; - llvm::StringRef name; + std::string name; llvm::SmallVector bounds; }; @@ -226,6 +228,10 @@ class OpenACCClauseCIREmitter final mlir::Location exprLoc = cgf.cgm.getLoc(curVarExpr->getBeginLoc()); llvm::SmallVector bounds; + std::string exprString; + llvm::raw_string_ostream os(exprString); + e->printPretty(os, nullptr, cgf.getContext().getPrintingPolicy()); + // Assemble the list of bounds. while (isa(curVarExpr)) { mlir::Location boundLoc = cgf.cgm.getLoc(curVarExpr->getBeginLoc()); @@ -267,20 +273,16 @@ class OpenACCClauseCIREmitter final bounds.push_back(createBound(boundLoc, lowerBound, upperBound, extent)); } - // TODO: OpenACC: if this is a member expr, emit the VarPtrPtr correctly. - if (isa(curVarExpr)) { - cgf.cgm.errorNYI(curVarExpr->getSourceRange(), - "OpenACC Data clause member expr"); - return {exprLoc, {}, {}, std::move(bounds)}; - } + if (const auto *memExpr = dyn_cast(curVarExpr)) + return {exprLoc, cgf.emitMemberExpr(memExpr).getPointer(), exprString, + std::move(bounds)}; // Sema has made sure that only 4 types of things can get here, array // subscript, array section, member expr, or DRE to a var decl (or the // former 3 wrapping a var-decl), so we should be able to assume this is // right. const auto *dre = cast(curVarExpr); - const auto *vd = cast(dre->getFoundDecl()->getCanonicalDecl()); - return {exprLoc, cgf.emitDeclRefLValue(dre).getPointer(), vd->getName(), + return {exprLoc, cgf.emitDeclRefLValue(dre).getPointer(), exprString, std::move(bounds)}; } diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c index 50c0519f0f29d..72471d4ec7874 100644 --- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c @@ -272,14 +272,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[3]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[3]"} loc #pragma acc serial loop copy(localArray[1:3]) for(int i = 0; i < 5; ++i); @@ -290,14 +290,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[1:3]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[1:3]"} loc #pragma acc kernels loop copy(localArray[:3]) for(int i = 0; i < 5; ++i); @@ -307,14 +307,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[:3]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[:3]"} loc #pragma acc parallel loop copy(localArray[1:]) for(int i = 0; i < 5; ++i); @@ -324,14 +324,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[1:]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[1:]"} loc #pragma acc serial loop copy(localArray[localVar1:localVar2]) for(int i = 0; i < 5; ++i); @@ -342,14 +342,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[localVar1:localVar2]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[localVar1:localVar2]"} loc #pragma acc kernels loop copy(localArray[:localVar2]) for(int i = 0; i < 5; ++i); @@ -359,14 +359,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[:localVar2]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[:localVar2]"} loc #pragma acc parallel loop copy(localArray[localVar1:]) for(int i = 0; i < 5; ++i); @@ -376,14 +376,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[localVar1:]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[localVar1:]"} loc #pragma acc serial loop copy(localPointer[3]) for(int i = 0; i < 5; ++i); @@ -393,14 +393,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[3]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[3]"} loc #pragma acc kernels loop copy(localPointer[1:3]) for(int i = 0; i < 5; ++i); @@ -411,14 +411,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[1:3]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[1:3]"} loc #pragma acc parallel loop copy(localPointer[:3]) for(int i = 0; i < 5; ++i); @@ -428,14 +428,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[:3]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[:3]"} loc #pragma acc serial loop copy(localPointer[localVar1:localVar2]) for(int i = 0; i < 5; ++i); @@ -446,14 +446,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[localVar1:localVar2]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[localVar1:localVar2]"} loc #pragma acc kernels loop copy(localPointer[:localVar2]) for(int i = 0; i < 5; ++i); @@ -463,14 +463,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[:localVar2]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[:localVar2]"} loc float *localArrayOfPtrs[5]; #pragma acc parallel loop copy(localArrayOfPtrs[3]) @@ -481,14 +481,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[3]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[3]"} loc #pragma acc serial loop copy(localArrayOfPtrs[3][2]) for(int i = 0; i < 5; ++i); @@ -504,14 +504,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[3][2]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[3][2]"} loc #pragma acc kernels loop copy(localArrayOfPtrs[localVar1:localVar2]) for(int i = 0; i < 5; ++i); @@ -522,14 +522,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2]"} loc #pragma acc parallel loop copy(localArrayOfPtrs[localVar1:]) for(int i = 0; i < 5; ++i); @@ -539,14 +539,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:]"} loc #pragma acc serial loop copy(localArrayOfPtrs[:localVar2]) for(int i = 0; i < 5; ++i); @@ -556,14 +556,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[:localVar2]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[:localVar2]"} loc #pragma acc kernels loop copy(localArrayOfPtrs[localVar1]) for(int i = 0; i < 5; ++i); @@ -573,14 +573,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1]"} loc #pragma acc parallel loop copy(localArrayOfPtrs[localVar1][localVar2]) for(int i = 0; i < 5; ++i); @@ -596,14 +596,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2]"} loc #pragma acc serial loop copy(localArrayOfPtrs[localVar1][localVar2:parmVar]) for(int i = 0; i < 5; ++i); @@ -620,14 +620,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2:parmVar]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2:parmVar]"} loc #pragma acc kernels loop copy(localArrayOfPtrs[localVar1][:parmVar]) for(int i = 0; i < 5; ++i); @@ -643,14 +643,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1][:parmVar]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1][:parmVar]"} loc #pragma acc parallel loop copy(localArrayOfPtrs[localVar1:localVar2][:1]) for(int i = 0; i < 5; ++i); @@ -667,14 +667,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][:1]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][:1]"} loc #pragma acc serial loop copy(localArrayOfPtrs[localVar1:localVar2][1:1]) for(int i = 0; i < 5; ++i); @@ -692,14 +692,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][1:1]"} loc // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.loop combined(serial) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][1:1]"} loc double threeDArray[5][6][7]; #pragma acc kernels loop copy(threeDArray[1][2][3]) @@ -722,14 +722,14 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS3:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray[1][2][3]"} loc // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) { // CHECK-NEXT: acc.loop combined(kernels) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray[1][2][3]"} loc #pragma acc parallel loop copy(threeDArray[1:1][2:1][3:1]) for(int i = 0; i < 5; ++i); @@ -754,12 +754,351 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS3:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray[1:1][2:1][3:1]"} loc // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) { // CHECK-NEXT: acc.loop combined(parallel) { // CHECK: acc.yield // CHECK-NEXT: } // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray[1:1][2:1][3:1]"} loc +} + +typedef struct StructTy { + int scalarMember; + int arrayMember[5]; + short twoDArrayMember[5][3]; + float *ptrArrayMember[5]; + double **ptrPtrMember; +} Struct ; + +void acc_compute_members() { + // CHECK: cir.func @acc_compute_members() + Struct localStruct; + // CHECK-NEXT: %[[LOCALSTRUCT:.*]] = cir.alloca !rec_StructTy, !cir.ptr, ["localStruct"] + +#pragma acc parallel loop copy(localStruct) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALSTRUCT]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "localStruct"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[LOCALSTRUCT]] : !cir.ptr) {dataClause = #acc, name = "localStruct"} + +#pragma acc serial loop copy(localStruct.scalarMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMBER]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "localStruct.scalarMember"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETMEMBER]] : !cir.ptr) {dataClause = #acc, name = "localStruct.scalarMember"} + +#pragma acc kernels loop copy(localStruct.arrayMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember"} loc + +#pragma acc parallel loop copy(localStruct.arrayMember[2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[2]"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[2]"} loc + +#pragma acc serial loop copy(localStruct.arrayMember[1:2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[TWO_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[1:2]"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[1:2]"} loc + +#pragma acc kernels loop copy(localStruct.arrayMember[1:]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[FOUR_CONST:.*]] = arith.constant 4 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[1:]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[1:]"} loc + +#pragma acc parallel loop copy(localStruct.arrayMember[:2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[TWO_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[:2]"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[:2]"} loc + +#pragma acc serial loop copy(localStruct.twoDArrayMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.twoDArrayMember"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) to varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.twoDArrayMember"} + +#pragma acc kernels loop copy(localStruct.twoDArrayMember[3][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.twoDArrayMember[3][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.twoDArrayMember[3][2]"} + +#pragma acc parallel loop copy(localStruct.twoDArrayMember[1:3][1:1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.twoDArrayMember[1:3][1:1]"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.twoDArrayMember[1:3][1:1]"} + +#pragma acc serial loop copy(localStruct.ptrArrayMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.ptrArrayMember"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) to varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.ptrArrayMember"} + +#pragma acc kernels loop copy(localStruct.ptrArrayMember[3][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.ptrArrayMember[3][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.ptrArrayMember[3][2]"} + +#pragma acc parallel loop copy(localStruct.ptrArrayMember[1:3][1:1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.ptrArrayMember[1:3][1:1]"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.ptrArrayMember[1:3][1:1]"} + +#pragma acc serial loop copy(localStruct.ptrPtrMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4] {name = "ptrPtrMember"} : !cir.ptr -> !cir.ptr>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) -> !cir.ptr>> {dataClause = #acc, name = "localStruct.ptrPtrMember"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>>) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) {dataClause = #acc, name = "localStruct.ptrPtrMember"} + +#pragma acc kernels loop copy(localStruct.ptrPtrMember[3][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4] {name = "ptrPtrMember"} : !cir.ptr -> !cir.ptr>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr>> {dataClause = #acc, name = "localStruct.ptrPtrMember[3][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) {dataClause = #acc, name = "localStruct.ptrPtrMember[3][2]"} + +#pragma acc parallel loop copy(localStruct.ptrPtrMember[1:3][1:1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4] {name = "ptrPtrMember"} : !cir.ptr -> !cir.ptr>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr>> {dataClause = #acc, name = "localStruct.ptrPtrMember[1:3][1:1]"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) {dataClause = #acc, name = "localStruct.ptrPtrMember[1:3][1:1]"} +} + +typedef struct InnerTy { + int a; + int b; +} Inner; + +typedef struct OuterTy { + Inner inner[4]; +} Outer; + +void copy_member_of_array_element_member() { + // CHECK: cir.func @copy_member_of_array_element_member() { + Outer outer; + // CHECK-NEXT: %[[OUTER:.*]] = cir.alloca !rec_OuterTy, !cir.ptr, ["outer"] + + #pragma acc parallel loop copy(outer.inner[2].b) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[GETINNER:.*]] = cir.get_member %[[OUTER]][0] {name = "inner"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[INNERDECAY:.*]] = cir.cast(array_to_ptrdecay, %[[GETINNER]] : !cir.ptr>), !cir.ptr + // CHECK-NEXT: %[[STRIDE:.*]] = cir.ptr_stride(%[[INNERDECAY]] : !cir.ptr, %[[TWO]] : !s32i), !cir.ptr + // CHECK-NEXT: %[[GETB:.*]] = cir.get_member %[[STRIDE]][1] {name = "b"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETB]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "outer.inner[2].b"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETB]] : !cir.ptr) {dataClause = #acc, name = "outer.inner[2].b"} } diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.cpp b/clang/test/CIR/CodeGenOpenACC/combined-copy.cpp new file mode 100644 index 0000000000000..c98fccb0f3b82 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.cpp @@ -0,0 +1,413 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +struct InnerStructTy { + int Member[5]; +}; +struct StructTy { + int scalarMember; + int arrayMember[5]; + short twoDArrayMember[5][3]; + InnerStructTy iSTy; + +void InlineFunc() { + // CHECK: cir.func {{.*}}InlineFunc{{.*}} + // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["this", init] + // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr>, !cir.ptr + +#pragma acc parallel loop copy(scalarMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} + +#pragma acc kernels loop copy(arrayMember[2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} + +#pragma acc kernels loop copy(twoDArrayMember[1][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + +#pragma acc kernels loop copy(iSTy) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel loop copy(iSTy.Member) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial loop copy(iSTy.Member[1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} + +#pragma acc parallel loop copy(this->scalarMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} + +#pragma acc kernels loop copy(this->arrayMember[2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} +#pragma acc kernels loop copy(this->twoDArrayMember[1][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + +#pragma acc kernels loop copy(this->iSTy) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel loop copy(this->iSTy.Member) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial loop copy(this->iSTy.Member[1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} +} + +void OutlineFunc(); +}; + +void StructTy::OutlineFunc() { + // CHECK: cir.func {{.*}}OutlineFunc{{.*}} + // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["this", init] + // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr>, !cir.ptr +#pragma acc parallel loop copy(scalarMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} +#pragma acc kernels loop copy(arrayMember[2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} +#pragma acc kernels loop copy(twoDArrayMember[1][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} +#pragma acc kernels loop copy(iSTy) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel loop copy(iSTy.Member) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial loop copy(iSTy.Member[1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} + +#pragma acc parallel loop copy(this->scalarMember) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} +#pragma acc kernels loop copy(this->arrayMember[2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} + +#pragma acc kernels loop copy(this->twoDArrayMember[1][2]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + +#pragma acc kernels loop copy(this->iSTy) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.loop combined(kernels) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel loop copy(this->iSTy.Member) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(parallel) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial loop copy(this->iSTy.Member[1]) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.loop combined(serial) { + // CHECK: acc.yield + // CHECK-NEXT: } + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} +} diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copy.c b/clang/test/CIR/CodeGenOpenACC/compute-copy.c index 549af7802c542..888bad29caa7c 100644 --- a/clang/test/CIR/CodeGenOpenACC/compute-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/compute-copy.c @@ -222,11 +222,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[3]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[3]"} loc #pragma acc serial copy(localArray[1:3]) ; @@ -237,11 +237,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[1:3]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[1:3]"} loc #pragma acc kernels copy(localArray[:3]) ; @@ -251,11 +251,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[:3]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[:3]"} loc #pragma acc parallel copy(localArray[1:]) ; @@ -265,11 +265,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[1:]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[1:]"} loc #pragma acc serial copy(localArray[localVar1:localVar2]) ; // CHECK-NEXT: %[[LV1:.*]] = cir.load{{.*}} %[[LOCAL1]] : !cir.ptr, !s32i @@ -279,11 +279,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[localVar1:localVar2]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[localVar1:localVar2]"} loc #pragma acc kernels copy(localArray[:localVar2]) ; @@ -293,11 +293,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[:localVar2]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[:localVar2]"} loc #pragma acc parallel copy(localArray[localVar1:]) ; @@ -307,11 +307,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAY]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localArray[localVar1:]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAY]] : !cir.ptr>) {dataClause = #acc, name = "localArray[localVar1:]"} loc #pragma acc serial copy(localPointer[3]) ; @@ -321,11 +321,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[3]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[3]"} loc #pragma acc kernels copy(localPointer[1:3]) ; @@ -336,11 +336,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[1:3]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[1:3]"} loc #pragma acc parallel copy(localPointer[:3]) ; @@ -350,11 +350,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[:3]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[:3]"} loc #pragma acc serial copy(localPointer[localVar1:localVar2]) ; @@ -365,11 +365,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[localVar1:localVar2]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[localVar1:localVar2]"} loc #pragma acc kernels copy(localPointer[:localVar2]) ; @@ -379,11 +379,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALPTR]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localPointer[:localVar2]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALPTR]] : !cir.ptr>) {dataClause = #acc, name = "localPointer[:localVar2]"} loc float *localArrayOfPtrs[5]; #pragma acc parallel copy(localArrayOfPtrs[3]) @@ -394,11 +394,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[3]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[3]"} loc #pragma acc serial copy(localArrayOfPtrs[3][2]) ; // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i @@ -413,11 +413,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[3][2]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[3][2]"} loc #pragma acc kernels copy(localArrayOfPtrs[localVar1:localVar2]) ; @@ -428,11 +428,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2]"} loc #pragma acc parallel copy(localArrayOfPtrs[localVar1:]) ; @@ -442,11 +442,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:]"} loc #pragma acc serial copy(localArrayOfPtrs[:localVar2]) ; @@ -456,11 +456,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[:localVar2]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[:localVar2]"} loc #pragma acc kernels copy(localArrayOfPtrs[localVar1]) ; @@ -470,11 +470,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1]"} loc #pragma acc parallel copy(localArrayOfPtrs[localVar1][localVar2]) ; @@ -490,11 +490,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2]"} loc #pragma acc serial copy(localArrayOfPtrs[localVar1][localVar2:parmVar]) ; @@ -511,11 +511,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2:parmVar]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1][localVar2:parmVar]"} loc #pragma acc kernels copy(localArrayOfPtrs[localVar1][:parmVar]) ; @@ -531,11 +531,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1][:parmVar]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1][:parmVar]"} loc #pragma acc parallel copy(localArrayOfPtrs[localVar1:localVar2][:1]) ; @@ -552,11 +552,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][:1]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][:1]"} loc #pragma acc serial copy(localArrayOfPtrs[localVar1:localVar2][1:1]) ; @@ -574,11 +574,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[LV1_CAST]] : si32) extent(%[[LV2_CAST]] : si16) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][1:1]"} loc // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]]) to varPtr(%[[LOCALARRAYOFPTRS]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localArrayOfPtrs[localVar1:localVar2][1:1]"} loc double threeDArray[5][6][7]; #pragma acc kernels copy(threeDArray[1][2][3]) @@ -601,11 +601,11 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS3:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray[1][2][3]"} loc // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) { // CHECK-NEXT: acc.terminator // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray[1][2][3]"} loc #pragma acc parallel copy(threeDArray[1:1][2:1][3:1]) ; @@ -630,9 +630,270 @@ void acc_compute(int parmVar) { // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS3:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) loc - // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) -> !cir.ptr x 6> x 5>> {dataClause = #acc, name = "threeDArray[1:1][2:1][3:1]"} loc // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) { // CHECK-NEXT: acc.yield // CHECK-NEXT: } loc - // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray"} loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 6> x 5>>) bounds(%[[BOUNDS]], %[[BOUNDS2]], %[[BOUNDS3]]) to varPtr(%[[THREEDARRAY]] : !cir.ptr x 6> x 5>>) {dataClause = #acc, name = "threeDArray[1:1][2:1][3:1]"} loc +} + +typedef struct StructTy { + int scalarMember; + int arrayMember[5]; + short twoDArrayMember[5][3]; + float *ptrArrayMember[5]; + double **ptrPtrMember; +} Struct ; + +void acc_compute_members() { + // CHECK: cir.func @acc_compute_members() + Struct localStruct; + // CHECK-NEXT: %[[LOCALSTRUCT:.*]] = cir.alloca !rec_StructTy, !cir.ptr, ["localStruct"] + +#pragma acc parallel copy(localStruct) + ; + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCALSTRUCT]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "localStruct"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[LOCALSTRUCT]] : !cir.ptr) {dataClause = #acc, name = "localStruct"} + +#pragma acc serial copy(localStruct.scalarMember) + ; + // CHECK-NEXT: %[[GETMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMBER]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "localStruct.scalarMember"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETMEMBER]] : !cir.ptr) {dataClause = #acc, name = "localStruct.scalarMember"} + +#pragma acc kernels copy(localStruct.arrayMember) + ; + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember"} loc + +#pragma acc parallel copy(localStruct.arrayMember[2]) + ; + // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[2]"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[2]"} loc + +#pragma acc serial copy(localStruct.arrayMember[1:2]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[TWO_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[1:2]"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[1:2]"} loc + +#pragma acc kernels copy(localStruct.arrayMember[1:]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[FOUR_CONST:.*]] = arith.constant 4 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) upperbound(%[[FOUR_CONST]] : i64) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[1:]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[1:]"} loc + +#pragma acc parallel copy(localStruct.arrayMember[:2]) + ; + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST2:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ZERO_CONST]] : i64) extent(%[[TWO_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST2]] : i64) + // CHECK-NEXT: %[[GETARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "localStruct.arrayMember[:2]"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS]]) to varPtr(%[[GETARRAYMEMBER]] : !cir.ptr>) {dataClause = #acc, name = "localStruct.arrayMember[:2]"} loc + +#pragma acc serial copy(localStruct.twoDArrayMember) + ; + // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.twoDArrayMember"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) to varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.twoDArrayMember"} + +#pragma acc kernels copy(localStruct.twoDArrayMember[3][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.twoDArrayMember[3][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.twoDArrayMember[3][2]"} + +#pragma acc parallel copy(localStruct.twoDArrayMember[1:3][1:1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.twoDArrayMember[1:3][1:1]"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.twoDArrayMember[1:3][1:1]"} + +#pragma acc serial copy(localStruct.ptrArrayMember) + ; + // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.ptrArrayMember"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) to varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.ptrArrayMember"} + +#pragma acc kernels copy(localStruct.ptrArrayMember[3][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.ptrArrayMember[3][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.ptrArrayMember[3][2]"} + +#pragma acc parallel copy(localStruct.ptrArrayMember[1:3][1:1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRARRAYMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][3] {name = "ptrArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "localStruct.ptrArrayMember[1:3][1:1]"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRARRAYMEMBER]] : !cir.ptr x 5>>) {dataClause = #acc, name = "localStruct.ptrArrayMember[1:3][1:1]"} + +#pragma acc serial copy(localStruct.ptrPtrMember) + ; + // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4] {name = "ptrPtrMember"} : !cir.ptr -> !cir.ptr>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) -> !cir.ptr>> {dataClause = #acc, name = "localStruct.ptrPtrMember"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>>) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) {dataClause = #acc, name = "localStruct.ptrPtrMember"} + +#pragma acc kernels copy(localStruct.ptrPtrMember[3][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[THREE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4] {name = "ptrPtrMember"} : !cir.ptr -> !cir.ptr>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr>> {dataClause = #acc, name = "localStruct.ptrPtrMember[3][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) {dataClause = #acc, name = "localStruct.ptrPtrMember[3][2]"} + +#pragma acc parallel copy(localStruct.ptrPtrMember[1:3][1:1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[THREE:.*]] = cir.const #cir.int<3> : !s32i + // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE]] : !s32i to si32 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[THREE_CAST]] : si32) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETPTRPTRMEMBER:.*]] = cir.get_member %[[LOCALSTRUCT]][4] {name = "ptrPtrMember"} : !cir.ptr -> !cir.ptr>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr>> {dataClause = #acc, name = "localStruct.ptrPtrMember[1:3][1:1]"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GETPTRPTRMEMBER]] : !cir.ptr>>) {dataClause = #acc, name = "localStruct.ptrPtrMember[1:3][1:1]"} + } diff --git a/clang/test/CIR/CodeGenOpenACC/compute-copy.cpp b/clang/test/CIR/CodeGenOpenACC/compute-copy.cpp new file mode 100644 index 0000000000000..6c5d6a7a617b4 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/compute-copy.cpp @@ -0,0 +1,341 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +struct InnerStructTy { + int Member[5]; +}; +struct StructTy { + int scalarMember; + int arrayMember[5]; + short twoDArrayMember[5][3]; + InnerStructTy iSTy; + +void InlineFunc() { + // CHECK: cir.func {{.*}}InlineFunc{{.*}} + // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["this", init] + // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr>, !cir.ptr + +#pragma acc parallel copy(scalarMember) + ; + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} + +#pragma acc kernels copy(arrayMember[2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} + +#pragma acc kernels copy(twoDArrayMember[1][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + +#pragma acc kernels copy(iSTy) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel copy(iSTy.Member) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial copy(iSTy.Member[1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} + +#pragma acc parallel copy(this->scalarMember) + ; + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} + +#pragma acc kernels copy(this->arrayMember[2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} +#pragma acc kernels copy(this->twoDArrayMember[1][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + +#pragma acc kernels copy(this->iSTy) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel copy(this->iSTy.Member) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial copy(this->iSTy.Member[1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} +} + +void OutlineFunc(); +}; + +void StructTy::OutlineFunc() { + // CHECK: cir.func {{.*}}OutlineFunc{{.*}} + // CHECK-NEXT: %[[THIS:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["this", init] + // CHECK-NEXT: cir.store %[[THIS_ARG:.*]], %[[THIS]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: %[[THIS_LOAD:.*]] = cir.load %[[THIS]] : !cir.ptr>, !cir.ptr +#pragma acc parallel copy(scalarMember) + ; + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} +#pragma acc kernels copy(arrayMember[2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} +#pragma acc kernels copy(twoDArrayMember[1][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} +#pragma acc kernels copy(iSTy) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel copy(iSTy.Member) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial copy(iSTy.Member[1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} + +#pragma acc parallel copy(this->scalarMember) + ; + // CHECK-NEXT: %[[GETSCALARMEM:.*]] = cir.get_member %[[THIS_LOAD]][0] {name = "scalarMember"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSCALARMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->scalarMember"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSCALARMEM]] : !cir.ptr) {dataClause = #acc, name = "this->scalarMember"} +#pragma acc kernels copy(this->arrayMember[2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][1] {name = "arrayMember"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETARRAYMEM]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->arrayMember[2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETARRAYMEM]] : !cir.ptr>) {dataClause = #acc, name = "this->arrayMember[2]"} + +#pragma acc kernels copy(this->twoDArrayMember[1][2]) + ; + // CHECK-NEXT: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i + // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[TWO_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS2:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GET2DARRAYMEM:.*]] = cir.get_member %[[THIS_LOAD]][2] {name = "twoDArrayMember"} : !cir.ptr -> !cir.ptr x 5>> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) -> !cir.ptr x 5>> {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr x 5>>) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr x 5>>) bounds(%[[BOUNDS1]], %[[BOUNDS2]]) to varPtr(%[[GET2DARRAYMEM]] : !cir.ptr x 5>>) {dataClause = #acc, name = "this->twoDArrayMember[1][2]"} + +#pragma acc kernels copy(this->iSTy) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "this->iSTy"} + // CHECK-NEXT: acc.kernels dataOperands(%[[COPYIN1]] : !cir.ptr) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr) to varPtr(%[[GETSTRUCTMEM]] : !cir.ptr) {dataClause = #acc, name = "this->iSTy"} + +#pragma acc parallel copy(this->iSTy.Member) + ; + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member"} + // CHECK-NEXT: acc.parallel dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member"} + +#pragma acc serial copy(this->iSTy.Member[1]) + ; + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUNDS1:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) startIdx(%[[ZERO_CONST]] : i64) + // CHECK-NEXT: %[[GETSTRUCTMEM:.*]] = cir.get_member %[[THIS_LOAD]][3] {name = "iSTy"} : !cir.ptr -> !cir.ptr + // CHECK-NEXT: %[[GETMEMOFSTRUCT:.*]] = cir.get_member %[[GETSTRUCTMEM]][0] {name = "Member"} : !cir.ptr -> !cir.ptr> + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) bounds(%[[BOUNDS1]]) -> !cir.ptr> {dataClause = #acc, name = "this->iSTy.Member[1]"} + // CHECK-NEXT: acc.serial dataOperands(%[[COPYIN1]] : !cir.ptr>) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr>) bounds(%[[BOUNDS1]]) to varPtr(%[[GETMEMOFSTRUCT]] : !cir.ptr>) {dataClause = #acc, name = "this->iSTy.Member[1]"} +}