https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/176708
### Issue 1: Dimension override missing with variable count expressions When variable count expressions were used with stride, the constant subsection path computed size first. This marked `ArgSizes` with byte size semantics. Variable expression logic later triggered, but reused `ArgSizes` assuming "bytes" semantics **Result**: `ArgSizes` wasn't overwritten with dimension count, breaking non-contiguous mapping. ### Issue 2: Variable stride not recognized as non-contiguous `CGOpenMPRuntime.cpp` failed to detect `DeclRefExpr`, `MemberExpr`, `ArraySubscriptExpr` as non-contiguous. Issue 3: Missing expression semantics in `OMPIRBuilder` `OMPIRBuilder.cpp` didn't handle dimension count for `OMP_MAP_NON_CONTIG` flag **Fixes**: `clang/lib/CodeGen/CGOpenMPRuntime.cpp` - Variable stride detection + dimension count logic `llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp` - Expression semantics for non-contiguous Detect variable stride expressions (`DeclRefExpr/MemberExpr/ArraySubscriptExpr`) as non-contiguous Set `OMP_MAP_NON_CONTIG` flag (`0x100000000000`) for variable stride/count. Generate 3D descriptor structures with runtime dimensions. Fix dimension override to use dimension count instead of byte size. Added testcases to cover stack arrays, heap pointers, struct members, etc., for expression semantics in non-contiguous update. >From 87d7113b3e9979bc31b0ba7d959652aa3b57e700 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 05:05:38 -0500 Subject: [PATCH 1/3] expression_semantics_patch --- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 29 ++++++++++++++++------- 1 file changed, 21 insertions(+), 8 deletions(-) diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8d7a207a91f5a..578e63351ca5f 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9773,16 +9773,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; } } >From a9dcaede9ba9b50e1e281ecdeaba16e400d17ecb Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 06:34:09 -0500 Subject: [PATCH 2/3] variable_stride_fix --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8981a0de6d0e4..027d1fb26bc97 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8033,12 +8033,28 @@ class MappableExprsHandler { const Expr *StrideExpr = OASE->getStride(); 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(); }); >From 74dcd62d5836cbb6820e3e35a79bf92203afe7c4 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 19 Jan 2026 03:25:48 -0500 Subject: [PATCH 3/3] cleanup --- ...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 +++++ .../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 ++++++ 20 files changed, 2401 insertions(+) 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/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/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
