https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/175505
**Issue 1: Dimension override bug 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 3ce6d1973f06635bcbf810e0b1909d60918bd390 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 05:05:38 -0500 Subject: [PATCH 1/4] 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 716f8582dd7b2..fe6c755b0e504 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9365,16 +9365,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays( ConstantInt::get(Int64Ty, 0)); SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size()); for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) { + bool IsNonContigEntry = + IsNonContiguous && + (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + CombinedInfo.Types[I] & + OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0); + // For NON_CONTIG entries ArgSizes must carry the dimension count + // (number of descriptor_dim records) – NOT the byte size expression. + // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a + // non-constant size so we marked them runtime and stored the byte size, + // leading the runtime to treat it as DimSize and overrun descriptors. + if (IsNonContigEntry) { + // Dims must be long enough and positive. + assert(I < CombinedInfo.NonContigInfo.Dims.size() && + "Induction variable is in-bounds with the NON_CONTIG Dims array"); + const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I]; + assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0"); + ConstSizes[I] = + ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); + continue; + } if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) { if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) { - if (IsNonContiguous && - static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - CombinedInfo.Types[I] & - OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG)) - ConstSizes[I] = - ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); - else - ConstSizes[I] = CI; + ConstSizes[I] = CI; continue; } } >From 0ffe45560dc0e9206486eafa7aeb8a9a11755844 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 05:10:55 -0500 Subject: [PATCH 2/4] testcases_count_expression --- .../strided_update_count_expression_from.c | 54 ++++++++ .../strided_update_count_expression_to.c | 72 ++++++++++ ...xpression_stride_greater_than_count_from.c | 42 ++++++ ...ontiguous_count_expression_variable_from.c | 123 +++++++++++++++++ ...-contiguous_count_expression_variable_to.c | 125 ++++++++++++++++++ ...n-contiguous_count_expression_zero_count.c | 43 ++++++ .../target_update_ptr_count_expression_from.c | 74 +++++++++++ .../target_update_ptr_count_expression_to.c | 82 ++++++++++++ ...ate_strided_struct_count_expression_from.c | 86 ++++++++++++ ...pdate_strided_struct_count_expression_to.c | 98 ++++++++++++++ 10 files changed, 799 insertions(+) create mode 100644 offload/test/offloading/strided_update_count_expression_from.c create mode 100644 offload/test/offloading/strided_update_count_expression_to.c create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_variable_from.c create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_variable_to.c create mode 100644 offload/test/offloading/target_non-contiguous_count_expression_zero_count.c create mode 100644 offload/test/offloading/target_update_ptr_count_expression_from.c create mode 100644 offload/test/offloading/target_update_ptr_count_expression_to.c create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression_to.c diff --git a/offload/test/offloading/strided_update_count_expression_from.c b/offload/test/offloading/strided_update_count_expression_from.c new file mode 100644 index 0000000000000..d33ba9e428af5 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression_from.c @@ -0,0 +1,54 @@ +// This test checks that "update from" clause in OpenMP is supported when the +// elements are updated in a non-contiguous manner. This test checks that +// #pragma omp target update from(data[0:len/2:2]) correctly updates only every +// other element (stride 2) from the device to the host + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 8; + double data[len]; +#pragma omp target map(tofrom : len, data[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + } + } + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + +#pragma omp target data map(to : len, data[0 : len]) + { +// Modify arrays on device +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += i; + } + +#pragma omp target update from(data[0 : len/2 : 2]) + } + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 4.000000 + // CHECK: 3.000000 + // CHECK: 8.000000 + // CHECK: 5.000000 + // CHECK: 12.000000 + // CHECK: 7.000000 + // CHECK-NOT: 2.000000 + // CHECK-NOT: 6.000000 + // CHECK-NOT: 10.000000 + // CHECK-NOT: 14.000000 + + printf("from target array results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + return 0; +} diff --git a/offload/test/offloading/strided_update_count_expression_to.c b/offload/test/offloading/strided_update_count_expression_to.c new file mode 100644 index 0000000000000..0b1f179c467e3 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression_to.c @@ -0,0 +1,72 @@ +// This test checks that "update to" clause in OpenMP is supported when the +// elements are updated in a non-contiguous manner. This test checks that +// #pragma omp target update to(data[0:len/2:2]) correctly updates only every +// other element (stride 2) from the host to the device + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 8; + double data[len]; + + // Initialize on host + for (int i = 0; i < len; i++) { + data[i] = i; + } + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : len, data[0 : len]) + { + // Initialize device to 20 +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] = 20.0; + } + + // Modify host for strided elements + data[0] = 10.0; + data[2] = 10.0; + data[4] = 10.0; + data[6] = 10.0; + +#pragma omp target update to(data[0 : len/2 : 2]) + + // Verify on device by adding 5 +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += 5.0; + } + } + + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 2.000000 + // CHECK: 3.000000 + // CHECK: 4.000000 + // CHECK: 5.000000 + // CHECK: 6.000000 + // CHECK: 7.000000 + + printf("device array values after update to:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + + return 0; +} diff --git a/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c b/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c new file mode 100644 index 0000000000000..2677fe1310760 --- /dev/null +++ b/offload/test/offloading/target_non-contiguous_count_expression_stride_greater_than_count_from.c @@ -0,0 +1,42 @@ +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 8; + double data[len]; +#pragma omp target map(tofrom : len, data[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + } + } + +#pragma omp target data map(to : len, data[0 : len]) + { +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += i; + } + + int small_count = 2; +#pragma omp target update from(data[0 : small_count : 10]) + } + + printf("from target array results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + return 0; +} + +// CHECK: from target array results: +// CHECK: 0.000000 +// CHECK: 1.000000 +// CHECK: 2.000000 +// CHECK: 3.000000 +// CHECK: 4.000000 +// CHECK: 5.000000 +// CHECK: 6.000000 +// CHECK: 7.000000 diff --git a/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c b/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c new file mode 100644 index 0000000000000..74fed75049cd2 --- /dev/null +++ b/offload/test/offloading/target_non-contiguous_count_expression_variable_from.c @@ -0,0 +1,123 @@ +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 16; + double data[len]; + double data1[len], data2[len]; + + // Initialize data, data1, data2 on device +#pragma omp target map(tofrom : len, data[0 : len], data1[0 : len], data2[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + data1[i] = i; + data2[i] = i * 10; + } + } + +#pragma omp target data map(to : len, data[0 : len], data1[0 : len], data2[0 : len]) + { + // Device modifies arrays: +#pragma omp target + { + for (int i = 0; i < len; i++) { + data[i] += i; + data1[i] += i; + data2[i] += 100; + } + } + + int count = 4; + // indices: {0, 2, 4, 6} +#pragma omp target update from(data[0 : count : 2]) + + // indices: {0, 2, 4, 6, 8, 10, 12, 14} +#pragma omp target update from(data[0 : len/2 : 2]) + + // indices: {2, 4, 6, 8, 10, 12, 14} +#pragma omp target update from(data[2 : len-4 : 2]) + +int partial_count = 4; + // indices: {0, 3, 6, 9} +#pragma omp target update from(data[0 : partial_count : 3]) + +int count1 = 3; +int count2 = 2; + // data1 indices: {0, 4, 8} + // data2 indices: {0, 5} +#pragma omp target update from(data1[0 : count1 : 4], data2[0 : count2 : 5]) + } + + // Print results + printf("from target array results (data):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + printf("from target array results (data1, data2):\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data2[i]); + printf("\n"); + + return 0; +} + +// CHECK: from target array results (data): +// CHECK: 0.000000 +// CHECK: 1.000000 +// CHECK: 4.000000 +// CHECK: 6.000000 +// CHECK: 8.000000 +// CHECK: 5.000000 +// CHECK: 12.000000 +// CHECK: 7.000000 +// CHECK: 16.000000 +// CHECK: 18.000000 +// CHECK: 20.000000 +// CHECK: 11.000000 +// CHECK: 24.000000 +// CHECK: 13.000000 +// CHECK: 28.000000 +// CHECK: 15.000000 + +// CHECK: from target array results (data1, data2): +// CHECK: data1: +// CHECK: 0.000000 +// CHECK: 1.000000 +// CHECK: 2.000000 +// CHECK: 3.000000 +// CHECK: 8.000000 +// CHECK: 5.000000 +// CHECK: 6.000000 +// CHECK: 7.000000 +// CHECK: 16.000000 +// CHECK: 9.000000 +// CHECK: 10.000000 +// CHECK: 11.000000 +// CHECK: 12.000000 +// CHECK: 13.000000 +// CHECK: 14.000000 +// CHECK: 15.000000 +// CHECK: data2: +// CHECK: 100.000000 +// CHECK: 10.000000 +// CHECK: 20.000000 +// CHECK: 30.000000 +// CHECK: 40.000000 +// CHECK: 150.000000 +// CHECK: 60.000000 +// CHECK: 70.000000 +// CHECK: 80.000000 +// CHECK: 90.000000 +// CHECK: 100.000000 +// CHECK: 110.000000 +// CHECK: 120.000000 +// CHECK: 130.000000 +// CHECK: 140.000000 +// CHECK: 150.000000 diff --git a/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c b/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c new file mode 100644 index 0000000000000..3cf53f0c206da --- /dev/null +++ b/offload/test/offloading/target_non-contiguous_count_expression_variable_to.c @@ -0,0 +1,125 @@ +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 16; + double data[len]; + double data1[len], data2[len]; + + // Initialize data, data1, data2 on device (and copy back to host) +#pragma omp target map(tofrom : len, data[0 : len], data1[0 : len], data2[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + data1[i] = i; + data2[i] = i * 10; + } + } + +#pragma omp target data map(to : len, data[0 : len], data1[0 : len], data2[0 : len]) + { + int count = 4; + // indices: {0, 2, 4, 6} +#pragma omp target update to(data[0 : count : 2]) + + int half_len = len / 2; + // indices: {0, 2, 4, 6, 8, 10, 12, 14} +#pragma omp target update to(data[0 : half_len : 2]) + + int dyn_count = len - 4; + // indices: {2, 4, 6, 8, 10, 12, 14} +#pragma omp target update to(data[2 : dyn_count : 2]) + + int partial_count = 4; + // indices: {0, 3, 6, 9} +#pragma omp target update to(data[0 : partial_count : 3]) + + int count1 = 3; + int count2 = 2; + // data1 indices: {0, 4, 8} + // data2 indices: {0, 5} +#pragma omp target update to(data1[0 : count1 : 4], data2[0 : count2 : 5]) + + // Device modifies arrays +#pragma omp target + { + for (int i = 0; i < len; i++) { + data[i] += i; // becomes 2*i on device + data1[i] += i; // becomes 2*i on device + data2[i] += 100; // becomes i*10 + 100 on device + } + } + } + + // Print results + printf("from target array results (data):\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + printf("from target array results (data1, data2):\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data2[i]); + printf("\n"); + + return 0; +} + +// CHECK: from target array results (data): +// CHECK: 0.000000 +// CHECK: 1.000000 +// CHECK: 2.000000 +// CHECK: 3.000000 +// CHECK: 4.000000 +// CHECK: 5.000000 +// CHECK: 6.000000 +// CHECK: 7.000000 +// CHECK: 8.000000 +// CHECK: 9.000000 +// CHECK: 10.000000 +// CHECK: 11.000000 +// CHECK: 12.000000 +// CHECK: 13.000000 +// CHECK: 14.000000 +// CHECK: 15.000000 + +// CHECK: from target array results (data1, data2): +// CHECK: data1: +// CHECK: 0.000000 +// CHECK: 1.000000 +// CHECK: 2.000000 +// CHECK: 3.000000 +// CHECK: 4.000000 +// CHECK: 5.000000 +// CHECK: 6.000000 +// CHECK: 7.000000 +// CHECK: 8.000000 +// CHECK: 9.000000 +// CHECK: 10.000000 +// CHECK: 11.000000 +// CHECK: 12.000000 +// CHECK: 13.000000 +// CHECK: 14.000000 +// CHECK: 15.000000 +// CHECK: data2: +// CHECK: 0.000000 +// CHECK: 10.000000 +// CHECK: 20.000000 +// CHECK: 30.000000 +// CHECK: 40.000000 +// CHECK: 50.000000 +// CHECK: 60.000000 +// CHECK: 70.000000 +// CHECK: 80.000000 +// CHECK: 90.000000 +// CHECK: 100.000000 +// CHECK: 110.000000 +// CHECK: 120.000000 +// CHECK: 130.000000 +// CHECK: 140.000000 +// CHECK: 150.000000 diff --git a/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c b/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c new file mode 100644 index 0000000000000..1567e79814b84 --- /dev/null +++ b/offload/test/offloading/target_non-contiguous_count_expression_zero_count.c @@ -0,0 +1,43 @@ + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 8; + double data[len]; +#pragma omp target map(tofrom : len, data[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + } + } + +#pragma omp target data map(to : len, data[0 : len]) + { +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += i; + } + + int zero_count = 0; +#pragma omp target update from(data[0 : zero_count : 2]) + } + + printf("from target array results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + return 0; +} + +// CHECK: from target array results: +// CHECK: 0.000000 +// CHECK: 1.000000 +// CHECK: 2.000000 +// CHECK: 3.000000 +// CHECK: 4.000000 +// CHECK: 5.000000 +// CHECK: 6.000000 +// CHECK: 7.000000 diff --git a/offload/test/offloading/target_update_ptr_count_expression_from.c b/offload/test/offloading/target_update_ptr_count_expression_from.c new file mode 100644 index 0000000000000..2c5740b906f6c --- /dev/null +++ b/offload/test/offloading/target_update_ptr_count_expression_from.c @@ -0,0 +1,74 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update from" clause in OpenMP supports strided +// sections. #pragma omp target update from(result[0:len/2:2]) updates every other +// element from device +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 32 + +int main() { + double *result = (double *)calloc(N, sizeof(double)); + int len = N; + + printf("initial host array values:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + printf("\n"); + +#pragma omp target data map(to : result[0 : N]) + { +#pragma omp target map(alloc : result[0 : N]) + for (int i = 0; i < N; i++) + result[i] += i; + + // Update strided elements from device: even indices 0,2,4,...,30 +#pragma omp target update from(result[0 : len/2 : 2]) + } + + printf("after target update from (even indices up to 30 updated):\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + printf("\n"); + + // Expected: even indices i, odd indices 0 + // CHECK: 0.000000 + // CHECK: 0.000000 + // CHECK: 2.000000 + // CHECK: 0.000000 + // CHECK: 4.000000 + // CHECK: 0.000000 + // CHECK: 6.000000 + // CHECK: 0.000000 + // CHECK: 8.000000 + // CHECK: 0.000000 + // CHECK: 10.000000 + // CHECK: 0.000000 + // CHECK: 12.000000 + // CHECK: 0.000000 + // CHECK: 14.000000 + // CHECK: 0.000000 + // CHECK: 16.000000 + // CHECK: 0.000000 + // CHECK: 18.000000 + // CHECK: 0.000000 + // CHECK: 20.000000 + // CHECK: 0.000000 + // CHECK: 22.000000 + // CHECK: 0.000000 + // CHECK: 24.000000 + // CHECK: 0.000000 + // CHECK: 26.000000 + // CHECK: 0.000000 + // CHECK: 28.000000 + // CHECK: 0.000000 + // CHECK: 30.000000 + // CHECK: 0.000000 + // CHECK-NOT: 1.000000 + // CHECK-NOT: 3.000000 + // CHECK-NOT: 31.000000 + + free(result); + return 0; +} diff --git a/offload/test/offloading/target_update_ptr_count_expression_to.c b/offload/test/offloading/target_update_ptr_count_expression_to.c new file mode 100644 index 0000000000000..3900592f10f93 --- /dev/null +++ b/offload/test/offloading/target_update_ptr_count_expression_to.c @@ -0,0 +1,82 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update to" clause in OpenMP supports strided sections. +// #pragma omp target update to(result[0:len/2:2]) updates every other element +// (stride 2) + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +int main() { + double *result = (double *)calloc(N, sizeof(double)); + int len = N; + + // Initialize on host + for (int i = 0; i < N; i++) { + result[i] = i; + } + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + printf("\n"); + +#pragma omp target data map(tofrom : result[0 : N]) + { +// Update strided elements to device: indices 0,2,4,6 +#pragma omp target update to(result[0 : len/2 : 2]) + +#pragma omp target + { + for (int i = 0; i < N; i++) { + result[i] += i; + } + } + } + + printf("from target array results:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + + // 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-NEXT: 10.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 14.000000 + // CHECK-NEXT: 15.000000 + + // CHECK: from target array results: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 10.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 14.000000 + // CHECK-NEXT: 16.000000 + // CHECK-NEXT: 18.000000 + // CHECK-NEXT: 20.000000 + // CHECK-NEXT: 22.000000 + // CHECK-NEXT: 24.000000 + // CHECK-NEXT: 26.000000 + // CHECK-NEXT: 28.000000 + // CHECK-NEXT: 30.000000 + + free(result); + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_count_expression_from.c b/offload/test/offloading/target_update_strided_struct_count_expression_from.c new file mode 100644 index 0000000000000..bb3e2ba577f7f --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_count_expression_from.c @@ -0,0 +1,86 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update from" with user-defined mapper supports strided +// sections using fixed-size arrays in structs. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +typedef struct { + double data[N]; + size_t len; +} T; + +#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len]) + +int main() { + T s; + s.len = N; + + for (int i = 0; i < N; i++) { + s.data[i] = i; + } + + printf("original host array values:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(mapper(custom), tofrom : s) + { +// Execute on device with explicit mapper +#pragma omp target map(mapper(custom), tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] += i; + } + } + +// Update strided elements from device: indices 0,2,4,6,8,10,12,14 +#pragma omp target update from(s.data[0 : s.len/2 : 2]) + } + + printf("from target array results:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", s.data[i]); + + // 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-NEXT: 10.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 14.000000 + // CHECK-NEXT: 15.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-NEXT: 20.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 24.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 28.000000 + // CHECK-NEXT: 15.000000 + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_count_expression_to.c b/offload/test/offloading/target_update_strided_struct_count_expression_to.c new file mode 100644 index 0000000000000..c4bc7faf367a0 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_count_expression_to.c @@ -0,0 +1,98 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update to" with struct member arrays supports strided +// sections using fixed-size arrays in structs. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +typedef struct { + double data[N]; + int len; +} T; + +int main() { + T s; + s.len = N; + + // Initialize struct array on host with simple sequential values + for (int i = 0; i < N; i++) { + s.data[i] = i; + } + + printf("original host struct array values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : s) + { +// Initialize device struct array to 20 +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = 20.0; + } + } + + // Modify host struct data for strided elements (set to 10) + for (int i = 0; i < 8; i++) { + s.data[i * 2] = 10.0; // Set even indices to 10 + } + +// indices 0,2,4,6,8,10,12,14 +#pragma omp target update to(s.data[0 : s.len/2 : 2]) + +// Execute on device - add 5 to verify update worked +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] += 5.0; + } + } + } + + printf("after target update to struct:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + + // CHECK: original host struct array values: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 2.0 + // CHECK-NEXT: 3.0 + // CHECK-NEXT: 4.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 6.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 8.0 + // CHECK-NEXT: 9.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 11.0 + // CHECK-NEXT: 12.0 + // CHECK-NEXT: 13.0 + // CHECK-NEXT: 14.0 + // CHECK-NEXT: 15.0 + + // CHECK: after target update to struct: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + + return 0; +} >From bbb8410ab124efb1c2fedbb5a707b8eb609f1190 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 06:34:09 -0500 Subject: [PATCH 3/4] 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 b8ee701c482bb..32b76f5d74d46 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7907,12 +7907,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 e14356943b68f3b2e883e09353fa6c230157caa8 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 12 Jan 2026 03:12:42 -0500 Subject: [PATCH 4/4] expression_semantics_all_tests --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 4 +- ...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, 2403 insertions(+), 2 deletions(-) create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c create mode 100644 clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c create mode 100644 clang/test/OpenMP/target_update_variable_count_and_stride_messages.c create mode 100644 offload/test/offloading/strided_update_count_expression.c create mode 100644 offload/test/offloading/strided_update_count_expression_complex.c create mode 100644 offload/test/offloading/strided_update_count_expression_misc.c create mode 100644 offload/test/offloading/strided_update_multiple_arrays_count_expression.c create mode 100644 offload/test/offloading/strided_update_multiple_arrays_variable_stride.c create mode 100644 offload/test/offloading/strided_update_variable_count_and_stride.c create mode 100644 offload/test/offloading/strided_update_variable_stride.c create mode 100644 offload/test/offloading/strided_update_variable_stride_complex.c create mode 100644 offload/test/offloading/strided_update_variable_stride_misc.c create mode 100644 offload/test/offloading/target_update_ptr_count_expression.c create mode 100644 offload/test/offloading/target_update_ptr_variable_count_and_stride.c create mode 100644 offload/test/offloading/target_update_ptr_variable_stride.c create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression.c create mode 100644 offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c create mode 100644 offload/test/offloading/target_update_strided_struct_variable_stride.c diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 32b76f5d74d46..766a9d24409b3 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7907,9 +7907,9 @@ class MappableExprsHandler { const Expr *StrideExpr = OASE->getStride(); if (!StrideExpr) return false; - + assert(StrideExpr->getType()->isIntegerType() && - "Stride expression must be of integer type"); + "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(); 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
