https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/175505
>From f7d8aa461270d858190344ce0ce17274b3375f62 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 19 Jan 2026 02:32:05 -0500 Subject: [PATCH] test: minimal formatting - only logical changes in CGOpenMPRuntime.cpp - Added variable stride detection for non-contiguous arrays - 16 lines of logical changes only - Formatted only the changed region - No formatting cascade to other parts of the file - Includes 20 test files (5 clang + 15 offload tests) --- 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 +++++ .../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 ++++++ 21 files changed, 2417 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/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/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
