https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/176699
None >From f2de4e4ce4b6127451fda9a67bd5e90608cf48c9 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 19 Jan 2026 02:18:34 -0500 Subject: [PATCH] test: minimal formatting - only logical changes formatted --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 16 + ...d_ptr_variable_count_and_stride_messages.c | 62 +++ ...date_strided_ptr_variable_count_messages.c | 57 ++ ...ate_strided_ptr_variable_stride_messages.c | 64 +++ ...truct_variable_count_and_stride_messages.c | 72 +++ ...pdate_variable_count_and_stride_messages.c | 85 +++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 495 ++---------------- .../strided_update_count_expression.c | 133 +++++ .../strided_update_count_expression_complex.c | 289 ++++++++++ .../strided_update_count_expression_misc.c | 99 ++++ ..._update_multiple_arrays_count_expression.c | 161 ++++++ ...d_update_multiple_arrays_variable_stride.c | 145 +++++ ...strided_update_variable_count_and_stride.c | 136 +++++ .../strided_update_variable_stride.c | 135 +++++ .../strided_update_variable_stride_complex.c | 293 +++++++++++ .../strided_update_variable_stride_misc.c | 94 ++++ .../target_update_ptr_count_expression.c | 99 ++++ ...get_update_ptr_variable_count_and_stride.c | 94 ++++ .../target_update_ptr_variable_stride.c | 95 ++++ ...t_update_strided_struct_count_expression.c | 97 ++++ ...strided_struct_variable_count_and_stride.c | 96 ++++ ...et_update_strided_struct_variable_stride.c | 95 ++++ 22 files changed, 2460 insertions(+), 452 deletions(-) create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c create mode 100644 clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c create mode 100644 clang/test/OpenMP/target_update_variable_count_and_stride_messages.c create mode 100644 offload/test/offloading/strided_update_count_expression.c create mode 100644 offload/test/offloading/strided_update_count_expression_complex.c create mode 100644 offload/test/offloading/strided_update_count_expression_misc.c create mode 100644 offload/test/offloading/strided_update_multiple_arrays_count_expression.c create mode 100644 offload/test/offloading/strided_update_multiple_arrays_variable_stride.c create mode 100644 offload/test/offloading/strided_update_variable_count_and_stride.c create mode 100644 offload/test/offloading/strided_update_variable_stride.c create mode 100644 offload/test/offloading/strided_update_variable_stride_complex.c create mode 100644 offload/test/offloading/strided_update_variable_stride_misc.c create mode 100644 offload/test/offloading/target_update_ptr_count_expression.c create mode 100644 offload/test/offloading/target_update_ptr_variable_count_and_stride.c create mode 100644 offload/test/offloading/target_update_ptr_variable_stride.c create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression.c create mode 100644 offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c create mode 100644 offload/test/offloading/target_update_strided_struct_variable_stride.c diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8981a0de6d0e4..e6f5f00a86922 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8034,11 +8034,27 @@ class MappableExprsHandler { if (!StrideExpr) return false; + assert(StrideExpr->getType()->isIntegerType() && + "Stride expression must be of integer type"); + + // If the stride is a variable (not a constant), it's non-contiguous. + const Expr *S = StrideExpr->IgnoreParenImpCasts(); + if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) { + if (isa<VarDecl>(DRE->getDecl()) || + isa<ParmVarDecl>(DRE->getDecl())) + return true; + } + if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S)) + return true; + + // If stride is not evaluatable as a constant, treat as + // non-contiguous. const auto Constant = StrideExpr->getIntegerConstantExpr(CGF.getContext()); if (!Constant) return false; + // Treat non-unitary strides as non-contiguous. return !Constant->isOne(); }); diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c new file mode 100644 index 0000000000000..932cd6b1c97bb --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c @@ -0,0 +1,62 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int count = 8; + int stride = 2; + int stride_large = 5; + double *data; + + // Valid strided array sections with both variable count and variable stride (FROM) + #pragma omp target update from(data[0:count:stride]) // OK - both variable + {} + + #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride + {} + + #pragma omp target update from(data[0:count:stride_large]) // OK - variable count, different stride + {} + + #pragma omp target update from(data[1:len-2:stride]) // OK - with offset, count expression + {} + + #pragma omp target update from(data[0:count/2:stride*2]) // OK - both expressions + {} + + #pragma omp target update from(data[0:(len+1)/2:stride+1]) // OK - complex expressions + {} + + #pragma omp target update from(data[2:count-2:len/4]) // OK - all expressions + {} + + // Edge cases + int stride_one = 1; + #pragma omp target update from(data[0:count:stride_one]) // OK - variable count, stride=1 + {} + + #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride + {} + + // Invalid compile-time constant strides with variable count + #pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:len/2:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable count and stride (TO) + #pragma omp target update to(data[0:count:stride]) // OK + {} + + #pragma omp target update to(data[0:len/2:stride]) // OK + {} + + #pragma omp target update to(data[0:count:stride*2]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c new file mode 100644 index 0000000000000..23fba9c8bc84f --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c @@ -0,0 +1,57 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int count = 8; + int divisor = 2; + double *data; + + // Valid strided array sections with variable count expressions (FROM) + #pragma omp target update from(data[0:count:2]) // OK - variable count + {} + + #pragma omp target update from(data[0:len/2:2]) // OK - count expression + {} + + #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction + {} + + #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression + {} + + #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication + {} + + #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo + {} + + // Variable count with stride = 1 (contiguous) + #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride + {} + + #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride + {} + + // Invalid stride expressions with variable count + #pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:len/2:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable count expressions (TO) + #pragma omp target update to(data[0:count:2]) // OK + {} + + #pragma omp target update to(data[0:len/2:2]) // OK + {} + + #pragma omp target update to(data[0:len-4:3]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c new file mode 100644 index 0000000000000..3f85ed0c48d66 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c @@ -0,0 +1,64 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int stride = 2; + int stride_large = 5; + double *data; + + // Valid strided array sections with variable stride (FROM) + #pragma omp target update from(data[0:8:stride]) // OK - variable stride + {} + + #pragma omp target update from(data[0:4:stride_large]) // OK - different variable stride + {} + + #pragma omp target update from(data[1:6:stride]) // OK - with offset + {} + + #pragma omp target update from(data[0:5:stride+1]) // OK - stride expression + {} + + #pragma omp target update from(data[0:4:stride*2]) // OK - stride multiplication + {} + + #pragma omp target update from(data[2:3:len/4]) // OK - stride from expression + {} + + // Edge case: stride = 1 (should be contiguous, not non-contiguous) + int stride_one = 1; + #pragma omp target update from(data[0:8:stride_one]) // OK - stride=1 is contiguous + {} + + // Invalid variable stride expressions + int zero_stride = 0; + int neg_stride = -1; + + // Note: These are runtime checks, so no compile-time error + #pragma omp target update from(data[0:8:zero_stride]) // OK at compile-time (runtime will fail) + {} + + #pragma omp target update from(data[0:4:neg_stride]) // OK at compile-time (runtime will fail) + {} + + // Compile-time constant invalid strides + #pragma omp target update from(data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable stride (TO) + #pragma omp target update to(data[0:8:stride]) // OK + {} + + #pragma omp target update to(data[0:5:stride+1]) // OK + {} + + #pragma omp target update to(data[0:4:stride*2]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c new file mode 100644 index 0000000000000..70775d5c8322c --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c @@ -0,0 +1,72 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 20 +typedef struct { + double data[N]; + int len; + int stride; +} T; + +int main(int argc, char **argv) { + T s; + s.len = 16; + s.stride = 2; + int count = 8; + int ext_stride = 3; + + // Valid strided struct member array sections with variable count/stride (FROM) + #pragma omp target update from(s.data[0:s.len/2:2]) // OK - member count expression + {} + + #pragma omp target update from(s.data[0:count:s.stride]) // OK - external count, member stride + {} + + #pragma omp target update from(s.data[0:s.len:ext_stride]) // OK - member count, external stride + {} + + #pragma omp target update from(s.data[0:count:ext_stride]) // OK - both external + {} + + #pragma omp target update from(s.data[0:s.len/2:s.stride]) // OK - both from struct + {} + + #pragma omp target update from(s.data[1:(s.len-2)/2:s.stride]) // OK - complex count expression + {} + + #pragma omp target update from(s.data[0:count*2:s.stride+1]) // OK - expressions for both + {} + + // Edge cases + int stride_one = 1; + #pragma omp target update from(s.data[0:s.len:stride_one]) // OK - stride=1 + {} + + #pragma omp target update from(s.data[0:s.len/s.stride:s.stride]) // OK - count depends on stride + {} + + // Invalid compile-time constant strides with variable count + #pragma omp target update from(s.data[0:s.len:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(s.data[0:count:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(s.data[1:s.len/2:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided struct member array sections with variable count and stride (TO) + #pragma omp target update to(s.data[0:s.len/2:2]) // OK + {} + + #pragma omp target update to(s.data[0:count:s.stride]) // OK + {} + + #pragma omp target update to(s.data[0:s.len:ext_stride]) // OK + {} + + #pragma omp target update to(s.data[0:count*2:s.stride+1]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(s.data[0:s.len:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c new file mode 100644 index 0000000000000..0082539538a32 --- /dev/null +++ b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c @@ -0,0 +1,85 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int count = 8; + int stride = 2; + int divisor = 2; + double data[100]; + + // Valid strided array sections with variable count expressions (FROM) + #pragma omp target update from(data[0:count:2]) // OK - variable count + {} + + #pragma omp target update from(data[0:len/2:2]) // OK - count expression + {} + + #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction + {} + + #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression + {} + + #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication + {} + + #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo + {} + + // Variable stride with constant/variable count + #pragma omp target update from(data[0:10:stride]) // OK - constant count, variable stride + {} + + #pragma omp target update from(data[0:count:stride]) // OK - both variable + {} + + #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride + {} + + #pragma omp target update from(data[0:count:stride*2]) // OK - variable count, stride expression + {} + + #pragma omp target update from(data[0:len/divisor:stride+1]) // OK - both expressions + {} + + // Variable count with stride = 1 (contiguous) + #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride + {} + + #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride + {} + + // Edge cases + int stride_one = 1; + #pragma omp target update from(data[0:len:stride_one]) // OK - stride=1 variable + {} + + #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride + {} + + // Invalid stride expressions with variable count + #pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:len/2:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable count expressions (TO) + #pragma omp target update to(data[0:count:2]) // OK + {} + + #pragma omp target update to(data[0:len/2:stride]) // OK + {} + + #pragma omp target update to(data[0:count:stride]) // OK + {} + + #pragma omp target update to(data[0:len/divisor:stride+1]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8d7a207a91f5a..418c6142380eb 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -21,7 +21,6 @@ #include "llvm/Analysis/CodeMetrics.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/OptimizationRemarkEmitter.h" -#include "llvm/Analysis/PostDominators.h" #include "llvm/Analysis/ScalarEvolution.h" #include "llvm/Analysis/TargetLibraryInfo.h" #include "llvm/Bitcode/BitcodeReader.h" @@ -50,7 +49,6 @@ #include "llvm/IR/Value.h" #include "llvm/MC/TargetRegistry.h" #include "llvm/Support/CommandLine.h" -#include "llvm/Support/Error.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/VirtualFileSystem.h" @@ -404,19 +402,18 @@ Value *createFakeIntVal(IRBuilderBase &Builder, OpenMPIRBuilder::InsertPointTy OuterAllocaIP, llvm::SmallVectorImpl<Instruction *> &ToBeDeleted, OpenMPIRBuilder::InsertPointTy InnerAllocaIP, - const Twine &Name = "", bool AsPtr = true, - bool Is64Bit = false) { + const Twine &Name = "", bool AsPtr = true) { Builder.restoreIP(OuterAllocaIP); - IntegerType *IntTy = Is64Bit ? Builder.getInt64Ty() : Builder.getInt32Ty(); Instruction *FakeVal; AllocaInst *FakeValAddr = - Builder.CreateAlloca(IntTy, nullptr, Name + ".addr"); + Builder.CreateAlloca(Builder.getInt32Ty(), nullptr, Name + ".addr"); ToBeDeleted.push_back(FakeValAddr); if (AsPtr) { FakeVal = FakeValAddr; } else { - FakeVal = Builder.CreateLoad(IntTy, FakeValAddr, Name + ".val"); + FakeVal = + Builder.CreateLoad(Builder.getInt32Ty(), FakeValAddr, Name + ".val"); ToBeDeleted.push_back(FakeVal); } @@ -424,10 +421,11 @@ Value *createFakeIntVal(IRBuilderBase &Builder, Builder.restoreIP(InnerAllocaIP); Instruction *UseFakeVal; if (AsPtr) { - UseFakeVal = Builder.CreateLoad(IntTy, FakeVal, Name + ".use"); + UseFakeVal = + Builder.CreateLoad(Builder.getInt32Ty(), FakeVal, Name + ".use"); } else { - UseFakeVal = cast<BinaryOperator>(Builder.CreateAdd( - FakeVal, Is64Bit ? Builder.getInt64(10) : Builder.getInt32(10))); + UseFakeVal = + cast<BinaryOperator>(Builder.CreateAdd(FakeVal, Builder.getInt32(10))); } ToBeDeleted.push_back(UseFakeVal); return FakeVal; @@ -765,28 +763,6 @@ static void raiseUserConstantDataAllocasToEntryBlock(IRBuilderBase &Builder, } } -static void hoistNonEntryAllocasToEntryBlock(llvm::BasicBlock &Block) { - llvm::SmallVector<llvm::Instruction *> AllocasToMove; - - auto ShouldHoistAlloca = [](const llvm::AllocaInst &AllocaInst) { - // TODO: For now, we support simple static allocations, we might need to - // move non-static ones as well. However, this will need further analysis to - // move the lenght arguments as well. - return !AllocaInst.isArrayAllocation(); - }; - - for (llvm::Instruction &Inst : Block) - if (auto *AllocaInst = llvm::dyn_cast<llvm::AllocaInst>(&Inst)) - if (ShouldHoistAlloca(*AllocaInst)) - AllocasToMove.push_back(AllocaInst); - - auto InsertPoint = - Block.getParent()->getEntryBlock().getTerminator()->getIterator(); - - for (llvm::Instruction *AllocaInst : AllocasToMove) - AllocaInst->moveBefore(InsertPoint); -} - void OpenMPIRBuilder::finalize(Function *Fn) { SmallPtrSet<BasicBlock *, 32> ParallelRegionBlockSet; SmallVector<BasicBlock *, 32> Blocks; @@ -831,8 +807,7 @@ void OpenMPIRBuilder::finalize(Function *Fn) { for (auto *V : OI.ExcludeArgsFromAggregate) Extractor.excludeArgFromAggregate(V); - Function *OutlinedFn = - Extractor.extractCodeRegion(CEAC, OI.Inputs, OI.Outputs); + Function *OutlinedFn = Extractor.extractCodeRegion(CEAC); // Forward target-cpu, target-features attributes to the outlined function. auto TargetCpuAttr = OuterFn->getFnAttribute("target-cpu"); @@ -892,13 +867,6 @@ void OpenMPIRBuilder::finalize(Function *Fn) { // Run a user callback, e.g. to add attributes. if (OI.PostOutlineCB) OI.PostOutlineCB(*OutlinedFn); - - if (OI.FixUpNonEntryAllocas) { - PostDominatorTree PostDomTree(*OutlinedFn); - for (llvm::BasicBlock &BB : *OutlinedFn) - if (PostDomTree.properlyDominates(&BB, &OutlinedFn->getEntryBlock())) - hoistNonEntryAllocasToEntryBlock(BB); - } } // Remove work items that have been completed. @@ -1726,7 +1694,6 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createParallel( IfCondition, NumThreads, PrivTID, PrivTIDAddr, ThreadID, ToBeDeletedVec); }; - OI.FixUpNonEntryAllocas = true; } else { // Generate OpenMP host runtime call OI.PostOutlineCB = [=, ToBeDeletedVec = @@ -1734,7 +1701,6 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createParallel( hostParallelCallback(this, OutlinedFn, OuterFn, Ident, IfCondition, PrivTID, PrivTIDAddr, ToBeDeletedVec); }; - OI.FixUpNonEntryAllocas = true; } OI.OuterAllocaBB = OuterAllocaBlock; @@ -2027,382 +1993,6 @@ static Value *emitTaskDependencies( return DepArray; } -/// Create the task duplication function passed to kmpc_taskloop. -Expected<Value *> OpenMPIRBuilder::createTaskDuplicationFunction( - Type *PrivatesTy, int32_t PrivatesIndex, TaskDupCallbackTy DupCB) { - unsigned ProgramAddressSpace = M.getDataLayout().getProgramAddressSpace(); - if (!DupCB) - return Constant::getNullValue( - PointerType::get(Builder.getContext(), ProgramAddressSpace)); - - // From OpenMP Runtime p_task_dup_t: - // Routine optionally generated by the compiler for setting the lastprivate - // flag and calling needed constructors for private/firstprivate objects (used - // to form taskloop tasks from pattern task) Parameters: dest task, src task, - // lastprivate flag. - // typedef void (*p_task_dup_t)(kmp_task_t *, kmp_task_t *, kmp_int32); - - auto *VoidPtrTy = PointerType::get(Builder.getContext(), ProgramAddressSpace); - - FunctionType *DupFuncTy = FunctionType::get( - Builder.getVoidTy(), {VoidPtrTy, VoidPtrTy, Builder.getInt32Ty()}, - /*isVarArg=*/false); - - Function *DupFunction = Function::Create(DupFuncTy, Function::InternalLinkage, - "omp_taskloop_dup", M); - Value *DestTaskArg = DupFunction->getArg(0); - Value *SrcTaskArg = DupFunction->getArg(1); - Value *LastprivateFlagArg = DupFunction->getArg(2); - DestTaskArg->setName("dest_task"); - SrcTaskArg->setName("src_task"); - LastprivateFlagArg->setName("lastprivate_flag"); - - IRBuilderBase::InsertPointGuard Guard(Builder); - Builder.SetInsertPoint( - BasicBlock::Create(Builder.getContext(), "entry", DupFunction)); - - auto GetTaskContextPtrFromArg = [&](Value *Arg) -> Value * { - Type *TaskWithPrivatesTy = - StructType::get(Builder.getContext(), {Task, PrivatesTy}); - Value *TaskPrivates = Builder.CreateGEP( - TaskWithPrivatesTy, Arg, {Builder.getInt32(0), Builder.getInt32(1)}); - Value *ContextPtr = Builder.CreateGEP( - PrivatesTy, TaskPrivates, - {Builder.getInt32(0), Builder.getInt32(PrivatesIndex)}); - return ContextPtr; - }; - - Value *DestTaskContextPtr = GetTaskContextPtrFromArg(DestTaskArg); - Value *SrcTaskContextPtr = GetTaskContextPtrFromArg(SrcTaskArg); - - DestTaskContextPtr->setName("destPtr"); - SrcTaskContextPtr->setName("srcPtr"); - - InsertPointTy AllocaIP(&DupFunction->getEntryBlock(), - DupFunction->getEntryBlock().begin()); - InsertPointTy CodeGenIP = Builder.saveIP(); - Expected<IRBuilderBase::InsertPoint> AfterIPOrError = - DupCB(AllocaIP, CodeGenIP, DestTaskContextPtr, SrcTaskContextPtr); - if (!AfterIPOrError) - return AfterIPOrError.takeError(); - Builder.restoreIP(*AfterIPOrError); - - Builder.CreateRetVoid(); - - return DupFunction; -} - -OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createTaskloop( - const LocationDescription &Loc, InsertPointTy AllocaIP, - BodyGenCallbackTy BodyGenCB, - llvm::function_ref<llvm::Expected<llvm::CanonicalLoopInfo *>()> LoopInfo, - Value *LBVal, Value *UBVal, Value *StepVal, bool Untied, Value *IfCond, - Value *GrainSize, bool NoGroup, int Sched, Value *Final, bool Mergeable, - Value *Priority, TaskDupCallbackTy DupCB, Value *TaskContextStructPtrVal) { - - if (!updateToLocation(Loc)) - return InsertPointTy(); - - uint32_t SrcLocStrSize; - Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize); - Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize); - - BasicBlock *TaskloopExitBB = - splitBB(Builder, /*CreateBranch=*/true, "taskloop.exit"); - BasicBlock *TaskloopBodyBB = - splitBB(Builder, /*CreateBranch=*/true, "taskloop.body"); - BasicBlock *TaskloopAllocaBB = - splitBB(Builder, /*CreateBranch=*/true, "taskloop.alloca"); - - InsertPointTy TaskloopAllocaIP = - InsertPointTy(TaskloopAllocaBB, TaskloopAllocaBB->begin()); - InsertPointTy TaskloopBodyIP = - InsertPointTy(TaskloopBodyBB, TaskloopBodyBB->begin()); - - if (Error Err = BodyGenCB(TaskloopAllocaIP, TaskloopBodyIP)) - return Err; - - llvm::Expected<llvm::CanonicalLoopInfo *> result = LoopInfo(); - if (!result) { - return result.takeError(); - } - - llvm::CanonicalLoopInfo *CLI = result.get(); - OutlineInfo OI; - OI.EntryBB = TaskloopAllocaBB; - OI.OuterAllocaBB = AllocaIP.getBlock(); - OI.ExitBB = TaskloopExitBB; - - // Add the thread ID argument. - SmallVector<Instruction *> ToBeDeleted; - // dummy instruction to be used as a fake argument - OI.ExcludeArgsFromAggregate.push_back(createFakeIntVal( - Builder, AllocaIP, ToBeDeleted, TaskloopAllocaIP, "global.tid", false)); - Value *FakeLB = createFakeIntVal(Builder, AllocaIP, ToBeDeleted, - TaskloopAllocaIP, "lb", false, true); - Value *FakeUB = createFakeIntVal(Builder, AllocaIP, ToBeDeleted, - TaskloopAllocaIP, "ub", false, true); - Value *FakeStep = createFakeIntVal(Builder, AllocaIP, ToBeDeleted, - TaskloopAllocaIP, "step", false, true); - // For Taskloop, we want to force the bounds being the first 3 inputs in the - // aggregate struct - OI.Inputs.insert(FakeLB); - OI.Inputs.insert(FakeUB); - OI.Inputs.insert(FakeStep); - if (TaskContextStructPtrVal) - OI.Inputs.insert(TaskContextStructPtrVal); - assert( - (TaskContextStructPtrVal && DupCB) || - (!TaskContextStructPtrVal && !DupCB) && - "Task context struct ptr and duplication callback must be both set " - "or both null"); - - // It isn't safe to run the duplication bodygen callback inside the post - // outlining callback so this has to be run now before we know the real task - // shareds structure type. - unsigned ProgramAddressSpace = M.getDataLayout().getProgramAddressSpace(); - Type *PointerTy = PointerType::get(Builder.getContext(), ProgramAddressSpace); - Type *FakeSharedsTy = StructType::get( - Builder.getContext(), - {FakeLB->getType(), FakeUB->getType(), FakeStep->getType(), PointerTy}); - Expected<Value *> TaskDupFnOrErr = createTaskDuplicationFunction( - FakeSharedsTy, - /*PrivatesIndex: the pointer after the three indices above*/ 3, DupCB); - if (!TaskDupFnOrErr) { - return TaskDupFnOrErr.takeError(); - } - Value *TaskDupFn = *TaskDupFnOrErr; - - OI.PostOutlineCB = [this, Ident, LBVal, UBVal, StepVal, Untied, - TaskloopAllocaBB, CLI, Loc, TaskDupFn, ToBeDeleted, - IfCond, GrainSize, NoGroup, Sched, FakeLB, FakeUB, - FakeStep, Final, Mergeable, - Priority](Function &OutlinedFn) mutable { - // Replace the Stale CI by appropriate RTL function call. - assert(OutlinedFn.hasOneUse() && - "there must be a single user for the outlined function"); - CallInst *StaleCI = cast<CallInst>(OutlinedFn.user_back()); - - /* Create the casting for the Bounds Values that can be used when outlining - * to replace the uses of the fakes with real values */ - BasicBlock *CodeReplBB = StaleCI->getParent(); - IRBuilderBase::InsertPoint CurrentIp = Builder.saveIP(); - Builder.SetInsertPoint(CodeReplBB->getFirstInsertionPt()); - Value *CastedLBVal = - Builder.CreateIntCast(LBVal, Builder.getInt64Ty(), true, "lb64"); - Value *CastedUBVal = - Builder.CreateIntCast(UBVal, Builder.getInt64Ty(), true, "ub64"); - Value *CastedStepVal = - Builder.CreateIntCast(StepVal, Builder.getInt64Ty(), true, "step64"); - Builder.restoreIP(CurrentIp); - - Builder.SetInsertPoint(StaleCI); - - // Gather the arguments for emitting the runtime call for - // @__kmpc_omp_task_alloc - Function *TaskAllocFn = - getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_task_alloc); - - Value *ThreadID = getOrCreateThreadID(Ident); - - if (!NoGroup) { - // Emit runtime call for @__kmpc_taskgroup - Function *TaskgroupFn = - getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_taskgroup); - Builder.CreateCall(TaskgroupFn, {Ident, ThreadID}); - } - - // `flags` Argument Configuration - // Task is tied if (Flags & 1) == 1. - // Task is untied if (Flags & 1) == 0. - // Task is final if (Flags & 2) == 2. - // Task is not final if (Flags & 2) == 0. - // Task is mergeable if (Flags & 4) == 4. - // Task is not mergeable if (Flags & 4) == 0. - // Task is priority if (Flags & 32) == 32. - // Task is not priority if (Flags & 32) == 0. - Value *Flags = Builder.getInt32(Untied ? 0 : 1); - if (Final) - Flags = Builder.CreateOr(Builder.getInt32(2), Flags); - if (Mergeable) - Flags = Builder.CreateOr(Builder.getInt32(4), Flags); - if (Priority) - Flags = Builder.CreateOr(Builder.getInt32(32), Flags); - - Value *TaskSize = Builder.getInt64( - divideCeil(M.getDataLayout().getTypeSizeInBits(Task), 8)); - - AllocaInst *ArgStructAlloca = - dyn_cast<AllocaInst>(StaleCI->getArgOperand(1)); - assert(ArgStructAlloca && - "Unable to find the alloca instruction corresponding to arguments " - "for extracted function"); - StructType *ArgStructType = - dyn_cast<StructType>(ArgStructAlloca->getAllocatedType()); - assert(ArgStructType && "Unable to find struct type corresponding to " - "arguments for extracted function"); - Value *SharedsSize = - Builder.getInt64(M.getDataLayout().getTypeStoreSize(ArgStructType)); - - // Emit the @__kmpc_omp_task_alloc runtime call - // The runtime call returns a pointer to an area where the task captured - // variables must be copied before the task is run (TaskData) - CallInst *TaskData = Builder.CreateCall( - TaskAllocFn, {/*loc_ref=*/Ident, /*gtid=*/ThreadID, /*flags=*/Flags, - /*sizeof_task=*/TaskSize, /*sizeof_shared=*/SharedsSize, - /*task_func=*/&OutlinedFn}); - - Value *Shareds = StaleCI->getArgOperand(1); - Align Alignment = TaskData->getPointerAlignment(M.getDataLayout()); - Value *TaskShareds = Builder.CreateLoad(VoidPtr, TaskData); - Builder.CreateMemCpy(TaskShareds, Alignment, Shareds, Alignment, - SharedsSize); - // Get the pointer to loop lb, ub, step from task ptr - // and set up the lowerbound,upperbound and step values - llvm::Value *Lb = Builder.CreateGEP( - ArgStructType, TaskShareds, {Builder.getInt32(0), Builder.getInt32(0)}); - - llvm::Value *Ub = Builder.CreateGEP( - ArgStructType, TaskShareds, {Builder.getInt32(0), Builder.getInt32(1)}); - - llvm::Value *Step = Builder.CreateGEP( - ArgStructType, TaskShareds, {Builder.getInt32(0), Builder.getInt32(2)}); - llvm::Value *Loadstep = Builder.CreateLoad(Builder.getInt64Ty(), Step); - - // set up the arguments for emitting kmpc_taskloop runtime call - // setting values for ifval, nogroup, sched, grainsize, task_dup - Value *IfCondVal = - IfCond ? Builder.CreateIntCast(IfCond, Builder.getInt32Ty(), true) - : Builder.getInt32(1); - // As __kmpc_taskgroup is called manually in OMPIRBuilder, NoGroupVal should - // always be 1 when calling __kmpc_taskloop to ensure it is not called again - Value *NoGroupVal = Builder.getInt32(1); - Value *SchedVal = Builder.getInt32(Sched); - Value *GrainSizeVal = - GrainSize ? Builder.CreateIntCast(GrainSize, Builder.getInt64Ty(), true) - : Builder.getInt64(0); - Value *TaskDup = TaskDupFn; - - Value *Args[] = {Ident, ThreadID, TaskData, IfCondVal, Lb, Ub, - Loadstep, NoGroupVal, SchedVal, GrainSizeVal, TaskDup}; - - // taskloop runtime call - Function *TaskloopFn = - getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_taskloop); - Builder.CreateCall(TaskloopFn, Args); - - // Emit the @__kmpc_end_taskgroup runtime call to end the taskgroup if - // nogroup is not defined - if (!NoGroup) { - Function *EndTaskgroupFn = - getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_end_taskgroup); - Builder.CreateCall(EndTaskgroupFn, {Ident, ThreadID}); - } - - StaleCI->eraseFromParent(); - - Builder.SetInsertPoint(TaskloopAllocaBB, TaskloopAllocaBB->begin()); - - LoadInst *SharedsOutlined = - Builder.CreateLoad(VoidPtr, OutlinedFn.getArg(1)); - OutlinedFn.getArg(1)->replaceUsesWithIf( - SharedsOutlined, - [SharedsOutlined](Use &U) { return U.getUser() != SharedsOutlined; }); - - Value *IV = CLI->getIndVar(); - Type *IVTy = IV->getType(); - Constant *One = ConstantInt::get(Builder.getInt64Ty(), 1); - - // When outlining, CodeExtractor will create GEP's to the LowerBound and - // UpperBound. These GEP's can be reused for loading the tasks respective - // bounds. - Value *TaskLB = nullptr; - Value *TaskUB = nullptr; - Value *LoadTaskLB = nullptr; - Value *LoadTaskUB = nullptr; - for (Instruction &I : *TaskloopAllocaBB) { - if (I.getOpcode() == Instruction::GetElementPtr) { - GetElementPtrInst &Gep = cast<GetElementPtrInst>(I); - if (ConstantInt *CI = dyn_cast<ConstantInt>(Gep.getOperand(2))) { - switch (CI->getZExtValue()) { - case 0: - TaskLB = &I; - break; - case 1: - TaskUB = &I; - break; - } - } - } else if (I.getOpcode() == Instruction::Load) { - LoadInst &Load = cast<LoadInst>(I); - if (Load.getPointerOperand() == TaskLB) { - assert(TaskLB != nullptr && "Expected value for TaskLB"); - LoadTaskLB = &I; - } else if (Load.getPointerOperand() == TaskUB) { - assert(TaskUB != nullptr && "Expected value for TaskUB"); - LoadTaskUB = &I; - } - } - } - - Builder.SetInsertPoint(CLI->getPreheader()->getTerminator()); - - assert(LoadTaskLB != nullptr && "Expected value for LoadTaskLB"); - assert(LoadTaskUB != nullptr && "Expected value for LoadTaskUB"); - Value *TripCountMinusOne = - Builder.CreateSDiv(Builder.CreateSub(LoadTaskUB, LoadTaskLB), FakeStep); - Value *TripCount = Builder.CreateAdd(TripCountMinusOne, One, "trip_cnt"); - Value *CastedTripCount = Builder.CreateIntCast(TripCount, IVTy, true); - Value *CastedTaskLB = Builder.CreateIntCast(LoadTaskLB, IVTy, true); - // set the trip count in the CLI - CLI->setTripCount(CastedTripCount); - - Builder.SetInsertPoint(CLI->getBody(), - CLI->getBody()->getFirstInsertionPt()); - - // The canonical loop is generated with a fixed lower bound. We need to - // update the index calculation code to use the task's lower bound. The - // generated code looks like this: - // %omp_loop.iv = phi ... - // ... - // %tmp = mul [type] %omp_loop.iv, step - // %user_index = add [type] tmp, lb - // OpenMPIRBuilder constructs canonical loops to have exactly three uses of - // the normalised induction variable: - // 1. This one: converting the normalised IV to the user IV - // 2. The increment (add) - // 3. The comparison against the trip count (icmp) - // (1) is the only use that is a mul followed by an add so this cannot match - // other IR. - assert(CLI->getIndVar()->getNumUses() == 3 && - "Canonical loop should have exactly three uses of the ind var"); - for (User *IVUser : CLI->getIndVar()->users()) { - if (auto *Mul = dyn_cast<BinaryOperator>(IVUser)) { - if (Mul->getOpcode() == Instruction::Mul) { - for (User *MulUser : Mul->users()) { - if (auto *Add = dyn_cast<BinaryOperator>(MulUser)) { - if (Add->getOpcode() == Instruction::Add) { - Add->setOperand(1, CastedTaskLB); - } - } - } - } - } - } - - FakeLB->replaceAllUsesWith(CastedLBVal); - FakeUB->replaceAllUsesWith(CastedUBVal); - FakeStep->replaceAllUsesWith(CastedStepVal); - for (Instruction *I : llvm::reverse(ToBeDeleted)) { - I->eraseFromParent(); - } - }; - - addOutlineInfo(std::move(OI)); - Builder.SetInsertPoint(TaskloopExitBB, TaskloopExitBB->begin()); - return Builder.saveIP(); -} - OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createTask( const LocationDescription &Loc, InsertPointTy AllocaIP, BodyGenCallbackTy BodyGenCB, bool Tied, Value *Final, Value *IfCondition, @@ -8320,7 +7910,7 @@ static Expected<Function *> createOutlinedFunction( // multiple mappings (technically not legal in OpenMP, but there is a case // in Fortran for Common Blocks where this is neccesary), we will end up // with GEP's into this array inside the kernel, that refer to the Global - // but are technically separate arguments to the kernel for all intents and + // but are technically seperate arguments to the kernel for all intents and // purposes. If we have mapped a segment that requires a GEP into the 0-th // index, it will fold into an referal to the Global, if we then encounter // this folded GEP during replacement all of the references to the @@ -8328,7 +7918,7 @@ static Expected<Function *> createOutlinedFunction( // that corresponds to it, including any other GEP's that refer to the // Global that may be other arguments. This will invalidate all of the other // preceding mapped arguments that refer to the same global that may be - // separate segments. To prevent this, we defer global processing until all + // seperate segments. To prevent this, we defer global processing until all // other processing has been performed. if (llvm::isa<llvm::GlobalValue, llvm::GlobalObject, llvm::GlobalVariable>( removeASCastIfPresent(Input))) { @@ -9090,6 +8680,8 @@ static void emitTargetCall( } unsigned NumTargetItems = Info.NumberOfPtrs; + // TODO: Use correct device ID + Value *DeviceID = Builder.getInt64(OMP_DEVICEID_UNDEF); uint32_t SrcLocStrSize; Constant *SrcLocStr = OMPBuilder.getOrCreateDefaultSrcLocStr(SrcLocStrSize); Value *RTLoc = OMPBuilder.getOrCreateIdent(SrcLocStr, SrcLocStrSize, @@ -9115,13 +8707,13 @@ static void emitTargetCall( // The presence of certain clauses on the target directive require the // explicit generation of the target task. if (RequiresOuterTargetTask) - return OMPBuilder.emitTargetTask(TaskBodyCB, RuntimeAttrs.DeviceID, - RTLoc, AllocaIP, Dependencies, - KArgs.RTArgs, Info.HasNoWait); + return OMPBuilder.emitTargetTask(TaskBodyCB, DeviceID, RTLoc, AllocaIP, + Dependencies, KArgs.RTArgs, + Info.HasNoWait); - return OMPBuilder.emitKernelLaunch( - Builder, OutlinedFnID, EmitTargetCallFallbackCB, KArgs, - RuntimeAttrs.DeviceID, RTLoc, AllocaIP); + return OMPBuilder.emitKernelLaunch(Builder, OutlinedFnID, + EmitTargetCallFallbackCB, KArgs, + DeviceID, RTLoc, AllocaIP); }()); Builder.restoreIP(AfterIP); @@ -9773,16 +9365,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays( ConstantInt::get(Int64Ty, 0)); SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size()); for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) { + bool IsNonContigEntry = + IsNonContiguous && + (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + CombinedInfo.Types[I] & + OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0); + // For NON_CONTIG entries ArgSizes must carry the dimension count + // (number of descriptor_dim records) – NOT the byte size expression. + // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a + // non-constant size so we marked them runtime and stored the byte size, + // leading the runtime to treat it as DimSize and overrun descriptors. + if (IsNonContigEntry) { + // Dims must be long enough and positive. + assert(I < CombinedInfo.NonContigInfo.Dims.size() && + "Induction variable is in-bounds with the NON_CONTIG Dims array"); + const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I]; + assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0"); + ConstSizes[I] = + ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); + continue; + } if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) { if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) { - if (IsNonContiguous && - static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - CombinedInfo.Types[I] & - OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG)) - ConstSizes[I] = - ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); - else - ConstSizes[I] = CI; + ConstSizes[I] = CI; continue; } } @@ -11010,13 +10615,6 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata( continue; } break; - case OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect: - case OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable: - if (!CE->getAddress()) { - ErrorFn(EMIT_MD_GLOBAL_VAR_INDIRECT_ERROR, E.second); - continue; - } - break; default: break; } @@ -11026,17 +10624,12 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata( // entry. Indirect variables are handled separately on the device. if (auto *GV = dyn_cast<GlobalValue>(CE->getAddress())) if ((GV->hasLocalLinkage() || GV->hasHiddenVisibility()) && - (Flags != - OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect && - Flags != OffloadEntriesInfoManager:: - OMPTargetGlobalVarEntryIndirectVTable)) + Flags != OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect) continue; // Indirect globals need to use a special name that doesn't match the name // of the associated host global. - if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect || - Flags == - OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable) + if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect) createOffloadEntry(CE->getAddress(), CE->getAddress(), CE->getVarSize(), Flags, CE->getLinkage(), CE->getVarName()); else @@ -11473,9 +11066,7 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo( } return; } - if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect || - Flags == - OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirectVTable) + if (Flags == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryIndirect) OffloadEntriesDeviceGlobalVar.try_emplace(VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, VarName.str()); diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c new file mode 100644 index 0000000000000..a87da289a9154 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression.c @@ -0,0 +1,133 @@ +// This test checks that "update from" and "update to" clauses in OpenMP are +// supported when elements are updated in a non-contiguous manner with variable +// count expression. Tests #pragma omp target update from/to(data[0:len/2:2]) +// where the count (len/2) is a variable expression, not a constant. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 10; + double data[len]; + + // ==================================================================== + // TEST 1: Update FROM device (device -> host) + // ==================================================================== + +#pragma omp target map(tofrom : len, data[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + } + } + + printf("Test 1: Update FROM device\n"); + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(to : len, data[0 : len]) + { +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += i; + } + +#pragma omp target update from(data[0 : len / 2 : 2]) + } + + printf("from target array results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + + // ==================================================================== + // TEST 2: Update TO device (host -> device) + // ==================================================================== + + for (int i = 0; i < len; i++) { + data[i] = i; + } + + printf("\nTest 2: Update TO device\n"); + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(tofrom : len, data[0 : len]) + { +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] = 20.0; + } + + data[0] = 10.0; + data[2] = 10.0; + data[4] = 10.0; + data[6] = 10.0; + data[8] = 10.0; + +#pragma omp target update to(data[0 : len / 2 : 2]) + +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += 5.0; + } + } + + printf("device array values after update to:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + + return 0; +} + +// CHECK: Test 1: Update FROM device +// CHECK: original host array values: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: from target array results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 16.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: Test 2: Update TO device +// CHECK: original host array values: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: device array values after update to: +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 25.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 25.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 25.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 25.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 25.000000 diff --git a/offload/test/offloading/strided_update_count_expression_complex.c b/offload/test/offloading/strided_update_count_expression_complex.c new file mode 100644 index 0000000000000..f9beef513da24 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression_complex.c @@ -0,0 +1,289 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests non-contiguous array sections with complex expression-based count +// scenarios including multiple struct arrays and non-zero offset. + +#include <omp.h> +#include <stdio.h> + +struct Data { + int offset; + int len; + double arr[20]; +}; + +int main() { + struct Data s1, s2; + + // Test 1: Multiple arrays with different count expressions + s1.len = 10; + s2.len = 10; + + // Initialize on device +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] = i; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] = i * 10; + } + } + + // Test FROM: Update multiple struct arrays with complex count expressions +#pragma omp target data map(to : s1, s2) + { +#pragma omp target + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] += i; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] += i * 10; + } + } + + // Complex count: (len-2)/2 and len*2/5 +#pragma omp target update from(s1.arr[0 : (s1.len - 2) / 2 : 2], \ + s2.arr[0 : s2.len * 2 / 5 : 2]) + } + + printf("Test 1 - complex count expressions (from):\n"); + printf("s1 results:\n"); + for (int i = 0; i < s1.len; i++) + printf("%f\n", s1.arr[i]); + + printf("s2 results:\n"); + for (int i = 0; i < s2.len; i++) + printf("%f\n", s2.arr[i]); + + // Reset for TO test +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] = i * 2; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] = i * 20; + } + } + + // Modify host data + for (int i = 0; i < (s1.len - 2) / 2; i++) { + s1.arr[i * 2] = i + 100; + } + for (int i = 0; i < s2.len * 2 / 5; i++) { + s2.arr[i * 2] = i + 50; + } + + // Test TO: Update with complex count expressions +#pragma omp target data map(to : s1, s2) + { +#pragma omp target update to(s1.arr[0 : (s1.len - 2) / 2 : 2], \ + s2.arr[0 : s2.len * 2 / 5 : 2]) + +#pragma omp target + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] += 100; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] += 100; + } + } + } + + printf("Test 1 - complex count expressions (to):\n"); + printf("s1 results:\n"); + for (int i = 0; i < s1.len; i++) + printf("%f\n", s1.arr[i]); + + printf("s2 results:\n"); + for (int i = 0; i < s2.len; i++) + printf("%f\n", s2.arr[i]); + + // Test 2: Complex count with non-zero offset + s1.offset = 2; + s1.len = 10; + s2.offset = 1; + s2.len = 10; + + // Initialize on device +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] = i; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] = i * 10; + } + } + + // Test FROM: Complex count with offset +#pragma omp target data map(to : s1, s2) + { +#pragma omp target + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] += i; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] += i * 10; + } + } + + // Count: (len-offset)/2 with stride 2 +#pragma omp target update from( \ + s1.arr[s1.offset : (s1.len - s1.offset) / 2 : 2], \ + s2.arr[s2.offset : (s2.len - s2.offset) / 2 : 2]) + } + + printf("Test 2 - complex count with offset (from):\n"); + printf("s1 results:\n"); + for (int i = 0; i < s1.len; i++) + printf("%f\n", s1.arr[i]); + + printf("s2 results:\n"); + for (int i = 0; i < s2.len; i++) + printf("%f\n", s2.arr[i]); + + // Reset for TO test +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] = i * 2; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] = i * 20; + } + } + + // Modify host data + for (int i = 0; i < (s1.len - s1.offset) / 2; i++) { + s1.arr[s1.offset + i * 2] = i + 100; + } + for (int i = 0; i < (s2.len - s2.offset) / 2; i++) { + s2.arr[s2.offset + i * 2] = i + 50; + } + + // Test TO: Update with complex count and offset +#pragma omp target data map(to : s1, s2) + { +#pragma omp target update to( \ + s1.arr[s1.offset : (s1.len - s1.offset) / 2 : 2], \ + s2.arr[s2.offset : (s2.len - s2.offset) / 2 : 2]) + +#pragma omp target + { + for (int i = 0; i < s1.len; i++) { + s1.arr[i] += 100; + } + for (int i = 0; i < s2.len; i++) { + s2.arr[i] += 100; + } + } + } + + printf("Test 2 - complex count with offset (to):\n"); + printf("s1 results:\n"); + for (int i = 0; i < s1.len; i++) + printf("%f\n", s1.arr[i]); + + printf("s2 results:\n"); + for (int i = 0; i < s2.len; i++) + printf("%f\n", s2.arr[i]); + + return 0; +} + +// CHECK: Test 1 - complex count expressions (from): +// CHECK: s1 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 +// CHECK: s2 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 70.000000 +// CHECK-NEXT: 80.000000 +// CHECK-NEXT: 90.000000 +// CHECK: Test 1 - complex count expressions (to): +// CHECK: s1 results: +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 101.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 103.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 16.000000 +// CHECK-NEXT: 18.000000 +// CHECK: s2 results: +// CHECK-NEXT: 50.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 51.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 52.000000 +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 53.000000 +// CHECK-NEXT: 140.000000 +// CHECK-NEXT: 160.000000 +// CHECK-NEXT: 180.000000 +// CHECK: Test 2 - complex count with offset (from): +// CHECK: s1 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 18.000000 +// CHECK: s2 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 140.000000 +// CHECK-NEXT: 80.000000 +// CHECK-NEXT: 90.000000 +// CHECK: Test 2 - complex count with offset (to): +// CHECK: s1 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 101.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 103.000000 +// CHECK-NEXT: 18.000000 +// CHECK: s2 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 50.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 51.000000 +// CHECK-NEXT: 80.000000 +// CHECK-NEXT: 52.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 53.000000 +// CHECK-NEXT: 160.000000 +// CHECK-NEXT: 180.000000 diff --git a/offload/test/offloading/strided_update_count_expression_misc.c b/offload/test/offloading/strided_update_count_expression_misc.c new file mode 100644 index 0000000000000..0e93a6d7df2cb --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression_misc.c @@ -0,0 +1,99 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Miscellaneous tests for count expressions: tests modulo, large stride with +// computed count, and boundary calculations to ensure expression semantics work +// correctly. + +#include <omp.h> +#include <stdio.h> + +int main() { + // ==================================================================== + // TEST 1: Modulo operation in count expression + // ==================================================================== + + int len1 = 10; + int divisor = 5; + double data1[len1]; + +#pragma omp target map(tofrom : len1, divisor, data1[0 : len1]) + { + for (int i = 0; i < len1; i++) { + data1[i] = i; + } + } + +#pragma omp target data map(to : len1, divisor, data1[0 : len1]) + { +#pragma omp target + { + for (int i = 0; i < len1; i++) { + data1[i] += i; + } + } + + // data[0:10%5:2] = data[0:0:2] updates no indices (count=0) +#pragma omp target update from(data1[0 : len1 % divisor : 2]) + } + + printf("Test 1: Modulo count expression\n"); + for (int i = 0; i < len1; i++) + printf("%f\n", data1[i]); + + // ==================================================================== + // TEST 2: Large stride with computed count for boundary coverage + // ==================================================================== + + int len2 = 10; + int stride = 5; + double data2[len2]; + +#pragma omp target map(tofrom : len2, stride, data2[0 : len2]) + { + for (int i = 0; i < len2; i++) { + data2[i] = i; + } + } + +#pragma omp target data map(to : len2, stride, data2[0 : len2]) + { +#pragma omp target + { + for (int i = 0; i < len2; i++) { + data2[i] += i; + } + } + + // data[0:(10+5-1)/5:5] = data[0:2:5] updates indices: 0, 5 +#pragma omp target update from(data2[0 : (len2 + stride - 1) / stride : stride]) + } + + printf("\nTest 2: Large stride count expression\n"); + for (int i = 0; i < len2; i++) + printf("%f\n", data2[i]); + + return 0; +} + +// CHECK: Test 1: Modulo count expression +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: Test 2: Large stride count expression +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 diff --git a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c new file mode 100644 index 0000000000000..9449baa663f67 --- /dev/null +++ b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c @@ -0,0 +1,161 @@ +// This test checks "update from" and "update to" with multiple arrays and +// variable count expressions. Tests both: (1) multiple arrays in single update +// clause with different count expressions, and (2) overlapping updates to the +// same array with various count expressions. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int n1 = 10, n2 = 10; + double arr1[n1], arr2[n2]; + + // ==================================================================== + // TEST 1: Update FROM - Multiple arrays in single update clause + // ==================================================================== + +#pragma omp target map(tofrom : n1, n2, arr1[0 : n1], arr2[0 : n2]) + { + for (int i = 0; i < n1; i++) { + arr1[i] = i; + } + for (int i = 0; i < n2; i++) { + arr2[i] = i * 10; + } + } + + printf("Test 1: Update FROM - Multiple arrays\n"); + +#pragma omp target data map(to : n1, n2, arr1[0 : n1], arr2[0 : n2]) + { +#pragma omp target + { + for (int i = 0; i < n1; i++) { + arr1[i] += i; + } + for (int i = 0; i < n2; i++) { + arr2[i] += 100; + } + } + + // Update with different count expressions in single clause: + // arr1[0:n1/2:2] = arr1[0:5:2] updates indices 0,2,4,6,8 + // arr2[0:n2/5:2] = arr2[0:2:2] updates indices 0,2 +#pragma omp target update from(arr1[0 : n1 / 2 : 2], arr2[0 : n2 / 5 : 2]) + } + + printf("from target arr1 results:\n"); + for (int i = 0; i < n1; i++) + printf("%f\n", arr1[i]); + + printf("\nfrom target arr2 results:\n"); + for (int i = 0; i < n2; i++) + printf("%f\n", arr2[i]); + + // ==================================================================== + // TEST 2: Update TO - Multiple arrays in single update clause + // ==================================================================== + + for (int i = 0; i < n1; i++) { + arr1[i] = i; + } + for (int i = 0; i < n2; i++) { + arr2[i] = i * 10; + } + + printf("\nTest 2: Update TO - Multiple arrays\n"); + +#pragma omp target data map(tofrom : n1, n2, arr1[0 : n1], arr2[0 : n2]) + { +#pragma omp target + { + for (int i = 0; i < n1; i++) { + arr1[i] = 100.0; + } + for (int i = 0; i < n2; i++) { + arr2[i] = 20.0; + } + } + + // Modify host + for (int i = 0; i < n1; i += 2) { + arr1[i] = 10.0; + } + for (int i = 0; i < n2; i += 2) { + arr2[i] = 5.0; + } + +#pragma omp target update to(arr1[0 : n1 / 2 : 2], arr2[0 : n2 / 5 : 2]) + +#pragma omp target + { + for (int i = 0; i < n1; i++) { + arr1[i] += 2.0; + } + for (int i = 0; i < n2; i++) { + arr2[i] += 2.0; + } + } + } + + printf("device arr1 values after update to:\n"); + for (int i = 0; i < n1; i++) + printf("%f\n", arr1[i]); + + printf("\ndevice arr2 values after update to:\n"); + for (int i = 0; i < n2; i++) + printf("%f\n", arr2[i]); + + return 0; +} + +// CHECK: Test 1: Update FROM - Multiple arrays +// CHECK: from target arr1 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 16.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: from target arr2 results: +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 30.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 50.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 70.000000 +// CHECK-NEXT: 80.000000 +// CHECK-NEXT: 90.000000 + +// CHECK: Test 2: Update TO - Multiple arrays +// CHECK: device arr1 values after update to: +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 + +// CHECK: device arr2 values after update to: +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 22.000000 diff --git a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c new file mode 100644 index 0000000000000..68c3eca4ccc56 --- /dev/null +++ b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c @@ -0,0 +1,145 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests multiple arrays with different variable strides in single update +// clause. + +#include <omp.h> +#include <stdio.h> + +int main() { + int stride1 = 2; + int stride2 = 2; + double data1[10], data2[10]; + + // ==================================================================== + // TEST 1: Update FROM - Multiple arrays with variable strides + // ==================================================================== + +#pragma omp target map(tofrom : stride1, stride2, data1[0 : 10], data2[0 : 10]) + { + for (int i = 0; i < 10; i++) { + data1[i] = i; + data2[i] = i * 10; + } + } + + printf("Test 1: Update FROM - Multiple arrays\n"); + +#pragma omp target data map(to : stride1, stride2, data1[0 : 10], data2[0 : 10]) + { +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data1[i] += i; + data2[i] += 100; + } + } + +#pragma omp target update from(data1[0 : 5 : stride1], data2[0 : 5 : stride2]) + } + + printf("from target data1:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data1[i]); + + printf("\nfrom target data2:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data2[i]); + + // ==================================================================== + // TEST 2: Update TO - Multiple arrays with variable strides + // ==================================================================== + + for (int i = 0; i < 10; i++) { + data1[i] = i; + data2[i] = i * 10; + } + + printf("\nTest 2: Update TO - Multiple arrays\n"); + +#pragma omp target data map(tofrom : stride1, stride2, data1[0 : 10], \ + data2[0 : 10]) + { +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data1[i] = 100.0; + data2[i] = 20.0; + } + } + + for (int i = 0; i < 10; i += 2) { + data1[i] = 10.0; + data2[i] = 5.0; + } + +#pragma omp target update to(data1[0 : 5 : stride1], data2[0 : 5 : stride2]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data1[i] += 2.0; + data2[i] += 2.0; + } + } + } + + printf("device data1 after update to:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data1[i]); + + printf("\ndevice data2 after update to:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data2[i]); + + return 0; +} + +// CHECK: Test 1: Update FROM - Multiple arrays +// CHECK: from target data1: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 16.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: from target data2: +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 30.000000 +// CHECK-NEXT: 140.000000 +// CHECK-NEXT: 50.000000 +// CHECK-NEXT: 160.000000 +// CHECK-NEXT: 70.000000 +// CHECK-NEXT: 180.000000 +// CHECK-NEXT: 90.000000 + +// CHECK: Test 2: Update TO - Multiple arrays +// CHECK: device data1 after update to: +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 102.000000 + +// CHECK: device data2 after update to: +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 22.000000 diff --git a/offload/test/offloading/strided_update_variable_count_and_stride.c b/offload/test/offloading/strided_update_variable_count_and_stride.c new file mode 100644 index 0000000000000..36056ab64250a --- /dev/null +++ b/offload/test/offloading/strided_update_variable_count_and_stride.c @@ -0,0 +1,136 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests combining variable count expression AND variable stride in array +// sections. + +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 10; + int stride = 2; + double data[len]; + + // ==================================================================== + // TEST 1: Update FROM - Variable count and stride + // ==================================================================== + +#pragma omp target map(tofrom : len, stride, data[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + } + } + + printf("Test 1: Update FROM - Variable count and stride\n"); + printf("original values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(to : len, stride, data[0 : len]) + { +#pragma omp target + { + for (int i = 0; i < len; i++) { + data[i] += i; + } + } + +#pragma omp target update from(data[0 : len / 2 : stride]) + } + + printf("from target results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + + // ==================================================================== + // TEST 2: Update TO - Variable count and stride + // ==================================================================== + + for (int i = 0; i < len; i++) { + data[i] = i; + } + + printf("\nTest 2: Update TO - Variable count and stride\n"); + printf("original values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(tofrom : len, stride, data[0 : len]) + { +#pragma omp target + { + for (int i = 0; i < len; i++) { + data[i] = 50.0; + } + } + + for (int i = 0; i < len / 2; i++) { + data[i * stride] = 10.0; + } + +#pragma omp target update to(data[0 : len / 2 : stride]) + +#pragma omp target + { + for (int i = 0; i < len; i++) { + data[i] += 5.0; + } + } + } + + printf("device values after update to:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + + return 0; +} + +// CHECK: Test 1: Update FROM - Variable count and stride +// CHECK: original values: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: from target results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 16.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: Test 2: Update TO - Variable count and stride +// CHECK: original values: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: device values after update to: +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 diff --git a/offload/test/offloading/strided_update_variable_stride.c b/offload/test/offloading/strided_update_variable_stride.c new file mode 100644 index 0000000000000..94723d91734a6 --- /dev/null +++ b/offload/test/offloading/strided_update_variable_stride.c @@ -0,0 +1,135 @@ +// This test checks "update from" and "update to" with variable stride. +// Tests data[0:5:stride] where stride is a variable, making it non-contiguous. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int stride = 2; + double data[10]; + + // ==================================================================== + // TEST 1: Update FROM device (device -> host) + // ==================================================================== + +#pragma omp target map(tofrom : stride, data[0 : 10]) + { + for (int i = 0; i < 10; i++) { + data[i] = i; + } + } + + printf("Test 1: Update FROM device\n"); + printf("original values:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(to : stride, data[0 : 10]) + { +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data[i] += i; + } + } + +#pragma omp target update from(data[0 : 5 : stride]) + } + + printf("from target results:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data[i]); + + // ==================================================================== + // TEST 2: Update TO device (host -> device) + // ==================================================================== + + for (int i = 0; i < 10; i++) { + data[i] = i; + } + + printf("\nTest 2: Update TO device\n"); + printf("original values:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(tofrom : stride, data[0 : 10]) + { +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data[i] = 50.0; + } + } + + for (int i = 0; i < 10; i += 2) { + data[i] = 10.0; + } + +#pragma omp target update to(data[0 : 5 : stride]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data[i] += 5.0; + } + } + } + + printf("device values after update to:\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data[i]); + + return 0; +} + +// CHECK: Test 1: Update FROM device +// CHECK: original values: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: from target results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 16.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: Test 2: Update TO device +// CHECK: original values: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: device values after update to: +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 55.000000 diff --git a/offload/test/offloading/strided_update_variable_stride_complex.c b/offload/test/offloading/strided_update_variable_stride_complex.c new file mode 100644 index 0000000000000..3c9857ec22178 --- /dev/null +++ b/offload/test/offloading/strided_update_variable_stride_complex.c @@ -0,0 +1,293 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests complex variable stride patterns with multiple arrays and offsets. + +#include <omp.h> +#include <stdio.h> + +struct Data { + int offset; + int stride; + double arr[20]; +}; + +int main() { + struct Data d1, d2; + int len1 = 10; + int len2 = 10; + + // Test 1: Complex stride expressions + int base_stride = 1; + int multiplier = 2; + d1.stride = 2; + d2.stride = 3; + + // Initialize on device +#pragma omp target map(tofrom : d1, d2, base_stride, multiplier) + { + for (int i = 0; i < len1; i++) { + d1.arr[i] = i * 3; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] = i * 30; + } + } + + // Test FROM: Complex stride expressions +#pragma omp target data map(to : d1, d2, base_stride, multiplier) + { +#pragma omp target + { + for (int i = 0; i < len1; i++) { + d1.arr[i] += i * 3; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] += i * 30; + } + } + + // Stride expressions: base_stride*multiplier and (d2.stride+1)/2 +#pragma omp target update from(d1.arr[0 : 5 : base_stride * multiplier], \ + d2.arr[0 : 3 : (d2.stride + 1) / 2]) + } + + printf("Test 1 - complex stride expressions (from):\n"); + printf("d1 results (stride=%d*%d=%d):\n", base_stride, multiplier, + base_stride * multiplier); + for (int i = 0; i < len1; i++) + printf("%f\n", d1.arr[i]); + + printf("d2 results (stride=(%d+1)/2=%d):\n", d2.stride, (d2.stride + 1) / 2); + for (int i = 0; i < len2; i++) + printf("%f\n", d2.arr[i]); + + // Reset for TO test +#pragma omp target map(tofrom : d1, d2) + { + for (int i = 0; i < len1; i++) { + d1.arr[i] = i * 4; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] = i * 40; + } + } + + // Modify host data with stride expressions + int stride1 = base_stride * multiplier; + int stride2 = (d2.stride + 1) / 2; + for (int i = 0; i < 5; i++) { + d1.arr[i * stride1] = i + 200; + } + for (int i = 0; i < 3; i++) { + d2.arr[i * stride2] = i + 150; + } + + // Test TO: Update with complex stride expressions +#pragma omp target data map(to : d1, d2, base_stride, multiplier) + { +#pragma omp target update to(d1.arr[0 : 5 : base_stride * multiplier], \ + d2.arr[0 : 3 : (d2.stride + 1) / 2]) + +#pragma omp target + { + for (int i = 0; i < len1; i++) { + d1.arr[i] += 200; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] += 200; + } + } + } + + printf("Test 1 - complex stride expressions (to):\n"); + printf("d1 results (stride=%d*%d=%d):\n", base_stride, multiplier, + base_stride * multiplier); + for (int i = 0; i < len1; i++) + printf("%f\n", d1.arr[i]); + + printf("d2 results (stride=(%d+1)/2=%d):\n", d2.stride, (d2.stride + 1) / 2); + for (int i = 0; i < len2; i++) + printf("%f\n", d2.arr[i]); + + // Test 2: Variable stride with non-zero offset + d1.offset = 2; + d1.stride = 2; + d2.offset = 1; + d2.stride = 2; + + // Initialize on device +#pragma omp target map(tofrom : d1, d2, len1, len2) + { + for (int i = 0; i < len1; i++) { + d1.arr[i] = i; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] = i * 10; + } + } + + // Test FROM: Variable stride with offset +#pragma omp target data map(to : d1, d2, len1, len2) + { +#pragma omp target + { + for (int i = 0; i < len1; i++) { + d1.arr[i] += i; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] += i * 10; + } + } + +#pragma omp target update from(d1.arr[d1.offset : 4 : d1.stride], \ + d2.arr[d2.offset : 4 : d2.stride]) + } + + printf("Test 2 - variable stride with offset (from):\n"); + printf("d1 results:\n"); + for (int i = 0; i < len1; i++) + printf("%f\n", d1.arr[i]); + + printf("d2 results:\n"); + for (int i = 0; i < len2; i++) + printf("%f\n", d2.arr[i]); + + // Reset for TO test +#pragma omp target map(tofrom : d1, d2) + { + for (int i = 0; i < len1; i++) { + d1.arr[i] = i * 2; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] = i * 20; + } + } + + // Modify host data + for (int i = 0; i < 4; i++) { + d1.arr[d1.offset + i * d1.stride] = i + 100; + } + for (int i = 0; i < 4; i++) { + d2.arr[d2.offset + i * d2.stride] = i + 50; + } + + // Test TO: Update with variable stride and offset +#pragma omp target data map(to : d1, d2) + { +#pragma omp target update to(d1.arr[d1.offset : 4 : d1.stride], \ + d2.arr[d2.offset : 4 : d2.stride]) + +#pragma omp target + { + for (int i = 0; i < len1; i++) { + d1.arr[i] += 100; + } + for (int i = 0; i < len2; i++) { + d2.arr[i] += 100; + } + } + } + + printf("Test 2 - variable stride with offset (to):\n"); + printf("d1 results:\n"); + for (int i = 0; i < len1; i++) + printf("%f\n", d1.arr[i]); + + printf("d2 results:\n"); + for (int i = 0; i < len2; i++) + printf("%f\n", d2.arr[i]); + + return 0; +} + +// CHECK: Test 1 - complex stride expressions (from): +// CHECK: d1 results (stride=1*2=2): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 18.000000 +// CHECK-NEXT: 24.000000 +// CHECK-NEXT: 15.000000 +// CHECK-NEXT: 18.000000 +// CHECK-NEXT: 21.000000 +// CHECK-NEXT: 24.000000 +// CHECK-NEXT: 27.000000 +// CHECK: d2 results (stride=(3+1)/2=2): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 90.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 150.000000 +// CHECK-NEXT: 180.000000 +// CHECK-NEXT: 210.000000 +// CHECK-NEXT: 240.000000 +// CHECK-NEXT: 270.000000 +// CHECK: Test 1 - complex stride expressions (to): +// CHECK: d1 results (stride=1*2=2): +// CHECK-NEXT: 200.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 201.000000 +// CHECK-NEXT: 12.000000 +// CHECK-NEXT: 202.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 203.000000 +// CHECK-NEXT: 28.000000 +// CHECK-NEXT: 204.000000 +// CHECK-NEXT: 36.000000 +// CHECK: d2 results (stride=(3+1)/2=2): +// CHECK-NEXT: 150.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 151.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 152.000000 +// CHECK-NEXT: 200.000000 +// CHECK-NEXT: 240.000000 +// CHECK-NEXT: 280.000000 +// CHECK-NEXT: 320.000000 +// CHECK-NEXT: 360.000000 +// CHECK: Test 2 - variable stride with offset (from): +// CHECK: d1 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 18.000000 +// CHECK: d2 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 20.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 60.000000 +// CHECK-NEXT: 140.000000 +// CHECK-NEXT: 80.000000 +// CHECK-NEXT: 90.000000 +// CHECK: Test 2 - variable stride with offset (to): +// CHECK: d1 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 101.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 103.000000 +// CHECK-NEXT: 18.000000 +// CHECK: d2 results: +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 50.000000 +// CHECK-NEXT: 40.000000 +// CHECK-NEXT: 51.000000 +// CHECK-NEXT: 80.000000 +// CHECK-NEXT: 52.000000 +// CHECK-NEXT: 120.000000 +// CHECK-NEXT: 53.000000 +// CHECK-NEXT: 160.000000 +// CHECK-NEXT: 180.000000 diff --git a/offload/test/offloading/strided_update_variable_stride_misc.c b/offload/test/offloading/strided_update_variable_stride_misc.c new file mode 100644 index 0000000000000..d27ae0123bfa8 --- /dev/null +++ b/offload/test/offloading/strided_update_variable_stride_misc.c @@ -0,0 +1,94 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Miscellaneous variable stride tests: stride=1, stride=array_size, stride from +// array subscript. + +#include <omp.h> +#include <stdio.h> + +int main() { + // ==================================================================== + // TEST 1: Variable stride = 1 (contiguous, but detected as variable) + // ==================================================================== + + int stride_one = 1; + double data1[10]; + +#pragma omp target map(tofrom : stride_one, data1[0 : 10]) + { + for (int i = 0; i < 10; i++) { + data1[i] = i; + } + } + +#pragma omp target data map(to : stride_one, data1[0 : 10]) + { +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data1[i] += i; + } + } + +#pragma omp target update from(data1[0 : 10 : stride_one]) + } + + printf("Test 1: Variable stride = 1\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data1[i]); + + // ==================================================================== + // TEST 2: Variable stride = array size (only 2 elements) + // ==================================================================== + + int stride_large = 5; + double data2[10]; + +#pragma omp target map(tofrom : stride_large, data2[0 : 10]) + { + for (int i = 0; i < 10; i++) { + data2[i] = i; + } + } + +#pragma omp target data map(to : stride_large, data2[0 : 10]) + { +#pragma omp target + { + for (int i = 0; i < 10; i++) { + data2[i] += i; + } + } + +#pragma omp target update from(data2[0 : 2 : stride_large]) + } + + printf("\nTest 2: Variable stride = 5\n"); + for (int i = 0; i < 10; i++) + printf("%f\n", data2[i]); + + return 0; +} + +// CHECK: Test 1: Variable stride = 1 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 5.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 + +// CHECK: Test 2: Variable stride = 5 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 1.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 3.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 7.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 diff --git a/offload/test/offloading/target_update_ptr_count_expression.c b/offload/test/offloading/target_update_ptr_count_expression.c new file mode 100644 index 0000000000000..c4b9fd566d401 --- /dev/null +++ b/offload/test/offloading/target_update_ptr_count_expression.c @@ -0,0 +1,99 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests non-contiguous array sections with expression-based count on +// heap-allocated pointer arrays with both FROM and TO directives. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int len = 10; + double *result = (double *)malloc(len * sizeof(double)); + + // Initialize host array to zero + for (int i = 0; i < len; i++) { + result[i] = 0; + } + + // Initialize on device +#pragma omp target enter data map(to : len, result[0 : len]) + +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] = i; + } + } + + // Test FROM: Modify on device, then update from device +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] += i * 10; + } + } + + // Update from device with expression-based count: len/2 elements +#pragma omp target update from(result[0 : len / 2 : 2]) + + printf("heap ptr count expression (from):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", result[i]); + + // Test TO: Reset, modify host, update to device +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] = i * 2; + } + } + + // Modify host data + for (int i = 0; i < len / 2; i++) { + result[i * 2] = i + 100; + } + + // Update to device with expression-based count +#pragma omp target update to(result[0 : len / 2 : 2]) + + // Read back full array +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] += 100; + } + } + +#pragma omp target update from(result[0 : len]) + + printf("heap ptr count expression (to):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", result[i]); + +#pragma omp target exit data map(delete : len, result[0 : len]) + free(result); + return 0; +} + +// CHECK: heap ptr count expression (from): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 44.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 66.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 88.000000 +// CHECK-NEXT: 0.000000 +// CHECK: heap ptr count expression (to): +// CHECK-NEXT: 200.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 201.000000 +// CHECK-NEXT: 106.000000 +// CHECK-NEXT: 202.000000 +// CHECK-NEXT: 110.000000 +// CHECK-NEXT: 203.000000 +// CHECK-NEXT: 114.000000 +// CHECK-NEXT: 204.000000 +// CHECK-NEXT: 118.000000 diff --git a/offload/test/offloading/target_update_ptr_variable_count_and_stride.c b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c new file mode 100644 index 0000000000000..1a28595969c69 --- /dev/null +++ b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c @@ -0,0 +1,94 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests heap-allocated pointers with both variable count expression and +// variable stride. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int len = 10; + int stride = 2; + double *result = (double *)malloc(len * sizeof(double)); + + for (int i = 0; i < len; i++) { + result[i] = 0; + } + +#pragma omp target enter data map(to : len, stride, result[0 : len]) + +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] = i; + } + } + + // Test FROM: Variable count and stride +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] += i * 10; + } + } + +#pragma omp target update from(result[0 : len / 2 : stride]) + + printf("heap ptr variable count and stride (from):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", result[i]); + + // Test TO: Reset, modify host, update to device +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] = i * 2; + } + } + + for (int i = 0; i < len / 2; i++) { + result[i * stride] = i + 100; + } + +#pragma omp target update to(result[0 : len / 2 : stride]) + +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] += 100; + } + } + +#pragma omp target update from(result[0 : len]) + + printf("heap ptr variable count and stride (to):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", result[i]); + +#pragma omp target exit data map(delete : len, stride, result[0 : len]) + free(result); + return 0; +} + +// CHECK: heap ptr variable count and stride (from): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 44.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 66.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 88.000000 +// CHECK-NEXT: 0.000000 +// CHECK: heap ptr variable count and stride (to): +// CHECK-NEXT: 200.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 201.000000 +// CHECK-NEXT: 106.000000 +// CHECK-NEXT: 202.000000 +// CHECK-NEXT: 110.000000 +// CHECK-NEXT: 203.000000 +// CHECK-NEXT: 114.000000 +// CHECK-NEXT: 204.000000 +// CHECK-NEXT: 118.000000 diff --git a/offload/test/offloading/target_update_ptr_variable_stride.c b/offload/test/offloading/target_update_ptr_variable_stride.c new file mode 100644 index 0000000000000..bea396065b760 --- /dev/null +++ b/offload/test/offloading/target_update_ptr_variable_stride.c @@ -0,0 +1,95 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests non-contiguous array sections with variable stride on heap-allocated +// pointers. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int stride = 2; + int len = 10; + double *result = (double *)malloc(len * sizeof(double)); + + // Initialize + for (int i = 0; i < len; i++) { + result[i] = 0; + } + +#pragma omp target enter data map(to : stride, len, result[0 : len]) + +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] = i; + } + } + + // Test FROM +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] += i * 10; + } + } + +#pragma omp target update from(result[0 : 5 : stride]) + + printf("heap ptr variable stride (from):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", result[i]); + + // Test TO: Reset, modify host, update to device +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] = i * 2; + } + } + + for (int i = 0; i < 5; i++) { + result[i * stride] = i + 100; + } + +#pragma omp target update to(result[0 : 5 : stride]) + +#pragma omp target map(alloc : result[0 : len]) + { + for (int i = 0; i < len; i++) { + result[i] += 100; + } + } + +#pragma omp target update from(result[0 : len]) + + printf("heap ptr variable stride (to):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", result[i]); + +#pragma omp target exit data map(delete : stride, len, result[0 : len]) + free(result); + return 0; +} + +// CHECK: heap ptr variable stride (from): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 22.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 44.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 66.000000 +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 88.000000 +// CHECK-NEXT: 0.000000 +// CHECK: heap ptr variable stride (to): +// CHECK-NEXT: 200.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 201.000000 +// CHECK-NEXT: 106.000000 +// CHECK-NEXT: 202.000000 +// CHECK-NEXT: 110.000000 +// CHECK-NEXT: 203.000000 +// CHECK-NEXT: 114.000000 +// CHECK-NEXT: 204.000000 +// CHECK-NEXT: 118.000000 diff --git a/offload/test/offloading/target_update_strided_struct_count_expression.c b/offload/test/offloading/target_update_strided_struct_count_expression.c new file mode 100644 index 0000000000000..1c1fd005c405f --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_count_expression.c @@ -0,0 +1,97 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests non-contiguous array sections with expression-based count on struct +// member arrays with both FROM and TO directives. + +#include <omp.h> +#include <stdio.h> + +struct S { + int len; + double data[20]; +}; + +int main() { + struct S s; + s.len = 10; + + // Initialize on device +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = i; + } + } + + // Test FROM: Modify on device, then update from device +#pragma omp target data map(to : s) + { +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] += i * 10; + } + } + + // Update from device with expression-based count: len/2 elements +#pragma omp target update from(s.data[0 : s.len / 2 : 2]) + } + + printf("struct count expression (from):\n"); + for (int i = 0; i < s.len; i++) + printf("%f\n", s.data[i]); + + // Test TO: Reset, modify host, update to device +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = i * 2; + } + } + + // Modify host data + for (int i = 0; i < s.len / 2; i++) { + s.data[i * 2] = i + 100; + } + + // Update to device with expression-based count +#pragma omp target data map(to : s) + { +#pragma omp target update to(s.data[0 : s.len / 2 : 2]) + +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] += 100; + } + } + } + + printf("struct count expression (to):\n"); + for (int i = 0; i < s.len; i++) + printf("%f\n", s.data[i]); + + return 0; +} + +// CHECK: struct count expression (from): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 11.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 33.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 77.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 +// CHECK: struct count expression (to): +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 101.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 103.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 104.000000 +// CHECK-NEXT: 18.000000 diff --git a/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c new file mode 100644 index 0000000000000..6daf10383e921 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c @@ -0,0 +1,96 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests struct member arrays with both variable count expression and variable +// stride. + +#include <omp.h> +#include <stdio.h> + +struct S { + int len; + int stride; + double data[20]; +}; + +int main() { + struct S s; + s.len = 10; + s.stride = 2; + + // Initialize +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = i; + } + } + + // Test FROM: Variable count and stride +#pragma omp target data map(to : s) + { +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] += i * 10; + } + } + +#pragma omp target update from(s.data[0 : s.len / 2 : s.stride]) + } + + printf("struct variable count and stride (from):\n"); + for (int i = 0; i < s.len; i++) + printf("%f\n", s.data[i]); + + // Test TO: Reset, modify host, update to device +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = i * 2; + } + } + + for (int i = 0; i < s.len / 2; i++) { + s.data[i * s.stride] = i + 100; + } + +#pragma omp target data map(to : s) + { +#pragma omp target update to(s.data[0 : s.len / 2 : s.stride]) + +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] += 100; + } + } + } + + printf("struct variable count and stride (to):\n"); + for (int i = 0; i < s.len; i++) + printf("%f\n", s.data[i]); + + return 0; +} + +// CHECK: struct variable count and stride (from): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 11.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 33.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 77.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 +// CHECK: struct variable count and stride (to): +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 101.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 103.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 104.000000 +// CHECK-NEXT: 18.000000 diff --git a/offload/test/offloading/target_update_strided_struct_variable_stride.c b/offload/test/offloading/target_update_strided_struct_variable_stride.c new file mode 100644 index 0000000000000..4cd9da629ca93 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_variable_stride.c @@ -0,0 +1,95 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// Tests non-contiguous array sections with variable stride on struct member +// arrays. + +#include <omp.h> +#include <stdio.h> + +struct S { + int stride; + double data[20]; +}; + +int main() { + struct S s; + s.stride = 2; + int len = 10; + + // Initialize +#pragma omp target map(tofrom : s, len) + { + for (int i = 0; i < len; i++) { + s.data[i] = i; + } + } + + // Test FROM +#pragma omp target data map(to : s, len) + { +#pragma omp target + { + for (int i = 0; i < len; i++) { + s.data[i] += i * 10; + } + } + +#pragma omp target update from(s.data[0 : 5 : s.stride]) + } + + printf("struct variable stride (from):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", s.data[i]); + + // Test TO: Reset, modify host, update to device +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < len; i++) { + s.data[i] = i * 2; + } + } + + for (int i = 0; i < 5; i++) { + s.data[i * s.stride] = i + 100; + } + +#pragma omp target data map(to : s) + { +#pragma omp target update to(s.data[0 : 5 : s.stride]) + +#pragma omp target + { + for (int i = 0; i < len; i++) { + s.data[i] += 100; + } + } + } + + printf("struct variable stride (to):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", s.data[i]); + + return 0; +} + +// CHECK: struct variable stride (from): +// CHECK-NEXT: 0.000000 +// CHECK-NEXT: 11.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 33.000000 +// CHECK-NEXT: 4.000000 +// CHECK-NEXT: 55.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 77.000000 +// CHECK-NEXT: 8.000000 +// CHECK-NEXT: 9.000000 +// CHECK: struct variable stride (to): +// CHECK-NEXT: 100.000000 +// CHECK-NEXT: 2.000000 +// CHECK-NEXT: 101.000000 +// CHECK-NEXT: 6.000000 +// CHECK-NEXT: 102.000000 +// CHECK-NEXT: 10.000000 +// CHECK-NEXT: 103.000000 +// CHECK-NEXT: 14.000000 +// CHECK-NEXT: 104.000000 +// CHECK-NEXT: 18.000000 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
