https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/181987
>From fe3b4e0cbaaebbdc8b5a59590e363d42c4ccf8b0 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 05:05:38 -0500 Subject: [PATCH 1/3] expression_semantics_patch --- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 29 ++++++++++++++++------- 1 file changed, 21 insertions(+), 8 deletions(-) diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 25f4da7c90d90..da03864cd6f15 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9945,16 +9945,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 2dcfc9a02d196c07901a99599bab26a8a3fa33fb Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Wed, 18 Feb 2026 05:29:00 -0500 Subject: [PATCH 2/3] tests_added --- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 4 +- .../strided_update_count_expression.c | 132 ++++++++ .../strided_update_count_expression_complex.c | 287 ++++++++++++++++++ ..._update_multiple_arrays_count_expression.c | 162 ++++++++++ .../target_update_ptr_count_expression.c | 99 ++++++ ...t_update_strided_struct_count_expression.c | 97 ++++++ 6 files changed, 779 insertions(+), 2 deletions(-) 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_multiple_arrays_count_expression.c create mode 100644 offload/test/offloading/target_update_ptr_count_expression.c create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression.c diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index da03864cd6f15..e58cf251c396f 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9948,8 +9948,8 @@ Error OpenMPIRBuilder::emitOffloadingArrays( bool IsNonContigEntry = IsNonContiguous && (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - CombinedInfo.Types[I] & - OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0); + 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 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..e18fc881a8021 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression.c @@ -0,0 +1,132 @@ +// 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> + +void test_1_update_from() { + int len = 10; + double data[len]; + + // Initialize data on host + 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]); +} + +void test_2_update_to() { + int len = 10; + double data[len]; + + 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]); +} + +int main() { + test_1_update_from(); + test_2_update_to(); + 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..560c4e767f882 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression_complex.c @@ -0,0 +1,287 @@ +// 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]; +}; + +void test_1_complex_count_expressions() { + struct Data s1, s2; + 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 - initialize on host + 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]); +} + +void test_2_complex_count_with_offset() { + struct Data s1, s2; + 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 - initialize on host + 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]); +} + +// 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 + +int main() { + test_1_complex_count_expressions(); + test_2_complex_count_with_offset(); + return 0; +} 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..a1472cacc4a39 --- /dev/null +++ b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c @@ -0,0 +1,162 @@ +// 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> + +void test_1_update_from_multiple() { + int n1 = 10, n2 = 10; + double arr1[n1], arr2[n2]; + +#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]); +} + +void test_2_update_to_multiple() { + int n1 = 10, n2 = 10; + double arr1[n1], arr2[n2]; + + 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]); +} + +int main() { + test_1_update_from_multiple(); + test_2_update_to_multiple(); + 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/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_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 >From 2f194562c128f7eb62d2cdf9973519399167a1eb Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 19 Feb 2026 06:02:56 -0500 Subject: [PATCH 3/3] revised --- .../ASTMatchers/ASTMatchersNarrowingTest.cpp | 88 +++++++++++++++++++ .../ASTMatchers/ASTMatchersNodeTest.cpp | 46 ++++++++++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 13 +-- .../Frontend/OpenMPIRBuilderTest.cpp | 48 ++++++++++ ...t_update_strided_struct_count_expression.c | 2 +- 5 files changed, 187 insertions(+), 10 deletions(-) diff --git a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp index b59b93bd9fcd4..d9647c201fc30 100644 --- a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp +++ b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp @@ -5168,6 +5168,94 @@ TEST_P(ASTMatchersTest, OMPFromClause_DoesNotMatchMapClause) { EXPECT_TRUE(notMatchesWithOpenMP(Source0, Matcher)); } +TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_ArraySection_CountExpression) { + StringRef Source0 = R"( + void foo() { + int count = 8; + int arr[100]; + #pragma omp target update from(arr[0:count:2]) + ; + } + )"; + + auto astUnit = + tooling::buildASTFromCodeWithArgs(Source0, {"-fopenmp=libomp"}); + ASSERT_TRUE(astUnit); + + auto Results = match(ompTargetUpdateDirective().bind("directive"), + astUnit->getASTContext()); + ASSERT_FALSE(Results.empty()); + + const auto *Directive = + Results[0].getNodeAs<OMPTargetUpdateDirective>("directive"); + ASSERT_TRUE(Directive); + + OMPFromClause *FromClause = nullptr; + for (auto *Clause : Directive->clauses()) { + if ((FromClause = dyn_cast<OMPFromClause>(Clause))) { + break; + } + } + ASSERT_TRUE(FromClause); + + for (const auto *VarExpr : FromClause->varlist()) { + const auto *ArraySection = dyn_cast<ArraySectionExpr>(VarExpr); + if (!ArraySection) + continue; + + // Verify length expression exists and is not a constant + const Expr *Length = ArraySection->getLength(); + ASSERT_TRUE(Length); + EXPECT_FALSE(isa<IntegerLiteral>(Length)) + << "Expected length to be a variable expression (count)"; + } +} + +TEST_P(ASTMatchersTest, + OMPTargetUpdateDirective_ArraySection_ComplexCountExpression) { + StringRef Source0 = R"( + void foo() { + int len = 16; + int count = 8; + int arr[100]; + #pragma omp target update from(arr[0:(len+count)/2:2]) + ; + } + )"; + + auto astUnit = + tooling::buildASTFromCodeWithArgs(Source0, {"-fopenmp=libomp"}); + ASSERT_TRUE(astUnit); + + auto Results = match(ompTargetUpdateDirective().bind("directive"), + astUnit->getASTContext()); + ASSERT_FALSE(Results.empty()); + + const auto *Directive = + Results[0].getNodeAs<OMPTargetUpdateDirective>("directive"); + ASSERT_TRUE(Directive); + + OMPFromClause *FromClause = nullptr; + for (auto *Clause : Directive->clauses()) { + if ((FromClause = dyn_cast<OMPFromClause>(Clause))) { + break; + } + } + ASSERT_TRUE(FromClause); + + for (const auto *VarExpr : FromClause->varlist()) { + const auto *ArraySection = dyn_cast<ArraySectionExpr>(VarExpr); + if (!ArraySection) + continue; + + // Verify length is a complex expression, not a constant + const Expr *Length = ArraySection->getLength(); + ASSERT_TRUE(Length); + EXPECT_FALSE(isa<IntegerLiteral>(Length)) + << "Expected length to be a complex expression ((len+count)/2)"; + } +} + TEST_P(ASTMatchersTest, HasAnyBase_DirectBase) { if (!GetParam().isCXX()) { return; diff --git a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp index 96f1b5249d859..7338ff5f302f6 100644 --- a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp +++ b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp @@ -3057,6 +3057,52 @@ void x(int x) { EXPECT_TRUE(notMatchesWithOpenMP(Source5, Matcher)); } +TEST(ASTMatchersTestOpenMP, OMPTargetUpdateDirective_CountExpression) { + auto Matcher = ompTargetUpdateDirective(hasAnyClause(ompFromClause())); + + StringRef Source0 = R"( + void foo() { + int len = 16; + int arr[100]; + #pragma omp target update from(arr[0:len/2:2]) + ; + } + )"; + EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher)); + + auto astUnit = + tooling::buildASTFromCodeWithArgs(Source0, {"-fopenmp=libomp"}); + ASSERT_TRUE(astUnit); + + auto Results = match(ompTargetUpdateDirective().bind("directive"), + astUnit->getASTContext()); + ASSERT_FALSE(Results.empty()); + + const auto *Directive = + Results[0].getNodeAs<OMPTargetUpdateDirective>("directive"); + ASSERT_TRUE(Directive); + + OMPFromClause *FromClause = nullptr; + for (auto *Clause : Directive->clauses()) { + if ((FromClause = dyn_cast<OMPFromClause>(Clause))) { + break; + } + } + ASSERT_TRUE(FromClause); + + for (const auto *VarExpr : FromClause->varlist()) { + const auto *ArraySection = dyn_cast<ArraySectionExpr>(VarExpr); + if (!ArraySection) + continue; + + const Expr *Length = ArraySection->getLength(); + ASSERT_TRUE(Length); + const auto *LengthLiteral = dyn_cast<IntegerLiteral>(Length); + EXPECT_FALSE(LengthLiteral) + << "Expected length to be a variable expression"; + } +} + TEST(ASTMatchersTest, Finder_DynamicOnlyAcceptsSomeMatchers) { MatchFinder Finder; EXPECT_TRUE(Finder.addDynamicMatcher(decl(), nullptr)); diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index e58cf251c396f..a1d5fa7cff992 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9950,19 +9950,14 @@ Error OpenMPIRBuilder::emitOffloadingArrays( (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. + // For NON_CONTIG entries, ArgSizes stores the dimension count (number of + // descriptor_dim records), not the byte size. 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"); + "Index must be in-bounds for 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]); + ConstSizes[I] = ConstantInt::get(Int64Ty, DimCount); continue; } if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) { diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp index 003a831833482..c5e9901463a8f 100644 --- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp +++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp @@ -7940,4 +7940,52 @@ TEST_F(OpenMPIRBuilderTest, spliceBBWithEmptyBB) { EXPECT_FALSE(Terminator->getDbgRecordRange().empty()); } +// Test that NON_CONTIG entries with count expressions store dimension count +TEST_F(OpenMPIRBuilderTest, EmitOffloadingArraysNonContigCountExpression) { + OpenMPIRBuilder OMPBuilder(*M); + OMPBuilder.initialize(); + IRBuilder<> Builder(BB); + + Value *BasePtr = ConstantPointerNull::get(Builder.getPtrTy()); + Value *CountExpr = + Builder.CreateUDiv(Builder.getInt64(16), Builder.getInt64(2)); + Value *ByteSize = + Builder.CreateMul(CountExpr, Builder.getInt64(4)); // Non-constant + + // Set up NON_CONTIG entry with count expression + OpenMPIRBuilder::MapInfosTy CombinedInfo; + CombinedInfo.BasePointers.push_back(BasePtr); + CombinedInfo.Pointers.push_back(BasePtr); + CombinedInfo.DevicePointers.push_back(OpenMPIRBuilder::DeviceInfoTy::Pointer); + CombinedInfo.Sizes.push_back(ByteSize); + CombinedInfo.Types.push_back(static_cast<omp::OpenMPOffloadMappingFlags>( + omp::OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG | + omp::OpenMPOffloadMappingFlags::OMP_MAP_TO)); + CombinedInfo.Names.push_back( + Builder.CreateGlobalString("data", "data_name", 0, M.get())); + + CombinedInfo.NonContigInfo.IsNonContiguous = true; + CombinedInfo.NonContigInfo.Dims.push_back(1); + CombinedInfo.NonContigInfo.Offsets.push_back({Builder.getInt64(0)}); + CombinedInfo.NonContigInfo.Counts.push_back({CountExpr}); + CombinedInfo.NonContigInfo.Strides.push_back({Builder.getInt64(2)}); + + OpenMPIRBuilder::TargetDataInfo Info(true, false); + Info.NumberOfPtrs = 1; + OpenMPIRBuilder::TargetDataRTArgs RTArgs; + using InsertPointTy = OpenMPIRBuilder::InsertPointTy; + + EXPECT_FALSE(OMPBuilder.emitOffloadingArraysAndArgs( + InsertPointTy(Builder.saveIP()), InsertPointTy(Builder.saveIP()), Info, + RTArgs, CombinedInfo, + [](unsigned) -> Expected<Function *> { + return static_cast<Function *>(nullptr); + }, + /*IsNonContiguous=*/true)); + + // Verify SizesArray is constant, not runtime + EXPECT_NE(dyn_cast<GlobalVariable>(RTArgs.SizesArray), nullptr); + EXPECT_EQ(dyn_cast<AllocaInst>(RTArgs.SizesArray), nullptr); +} + } // namespace diff --git a/offload/test/offloading/target_update_strided_struct_count_expression.c b/offload/test/offloading/target_update_strided_struct_count_expression.c index 1c1fd005c405f..89b0e1b4c7ea7 100644 --- a/offload/test/offloading/target_update_strided_struct_count_expression.c +++ b/offload/test/offloading/target_update_strided_struct_count_expression.c @@ -54,7 +54,7 @@ int main() { } // Update to device with expression-based count -#pragma omp target data map(to : s) +#pragma omp target data map(alloc : s) { #pragma omp target update to(s.data[0 : s.len / 2 : 2]) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
