https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/176914
**Problem:** The non‑contiguous update is being applied relative to `&s` (the struct on the host stack) instead of relative to `p = s.data` (the heap array). As a result, almost every slice in the strided copy becomes a no‑op (no mapping found), and nothing on the host changes. The base address the runtime uses for the non‑contiguous slices is therefore incorrect for this case. Data transfer issue from device to host. **IR log:** Upon modifying the IR appropriately, the DEBUG showed the Base address matches the point where the op should begin: Old: Base0 = Base1 != Begin0 ==> wrong `omptarget --> Entry 0: Base=0x00007fffddf31f40, Begin=0x0000558055e40fd0, Size=8, Type=0x0, Name=s omptarget --> Entry 1: Base=0x00007fffddf31f40, Begin=0x00007fffddf31ff8, Size=2, Type=0x1100000000002,` New: Corrected `omptarget --> Entry 0: Base=0x000055e90f1e4fd0, Begin=0x000055e90f1e4fd0, Size=8, Type=0x0, Name=unknown omptarget --> Entry 1: Base=0x000055e90f1e4fd0, Begin=0x00007ffc27f3ab18, Size=2, Type=0x1100000000002, ` **Fix:** Least modification done to BP pointing the array data, not the mapped pointer. Added testcases to validate the working. OpenMP_VV sollve target-update tests pass now: `test_target_update_mapper_from_discontiguous.c test_target_update_mapper_to_discontiguous.c` >From 6f610bd8d4f6b149b07a1a4e7bfa66100cfcbbbd Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 05:05:38 -0500 Subject: [PATCH 1/6] expression_semantics_patch --- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 29 ++++++++++++++++------- 1 file changed, 21 insertions(+), 8 deletions(-) diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8d7a207a91f5a..578e63351ca5f 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9773,16 +9773,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays( ConstantInt::get(Int64Ty, 0)); SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size()); for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) { + bool IsNonContigEntry = + IsNonContiguous && + (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + CombinedInfo.Types[I] & + OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0); + // For NON_CONTIG entries ArgSizes must carry the dimension count + // (number of descriptor_dim records) – NOT the byte size expression. + // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a + // non-constant size so we marked them runtime and stored the byte size, + // leading the runtime to treat it as DimSize and overrun descriptors. + if (IsNonContigEntry) { + // Dims must be long enough and positive. + assert(I < CombinedInfo.NonContigInfo.Dims.size() && + "Induction variable is in-bounds with the NON_CONTIG Dims array"); + const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I]; + assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0"); + ConstSizes[I] = + ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); + continue; + } if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) { if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) { - if (IsNonContiguous && - static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - CombinedInfo.Types[I] & - OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG)) - ConstSizes[I] = - ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); - else - ConstSizes[I] = CI; + ConstSizes[I] = CI; continue; } } >From 6f16202c8a14ac3125a38a47be83f41ea54bb2f3 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Thu, 4 Dec 2025 06:34:09 -0500 Subject: [PATCH 2/6] variable_stride_fix --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8981a0de6d0e4..027d1fb26bc97 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8033,12 +8033,28 @@ class MappableExprsHandler { const Expr *StrideExpr = OASE->getStride(); if (!StrideExpr) return false; + + assert(StrideExpr->getType()->isIntegerType() && + "Stride expression must be of integer type"); + + // If the stride is a variable (not a constant), it's non-contiguous. + const Expr *S = StrideExpr->IgnoreParenImpCasts(); + if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) { + if (isa<VarDecl>(DRE->getDecl()) || + isa<ParmVarDecl>(DRE->getDecl())) + return true; + } + if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S)) + return true; + // If stride is not evaluatable as a constant, treat as + // non-contiguous. const auto Constant = StrideExpr->getIntegerConstantExpr(CGF.getContext()); if (!Constant) return false; + // Treat non-unitary strides as non-contiguous. return !Constant->isOne(); }); >From daebaaead7f5b432d63e677d49f0c80ba8abe491 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Tue, 20 Jan 2026 06:15:41 -0500 Subject: [PATCH 3/6] fix base pointer alignment --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 027d1fb26bc97..f4ae2d7d261bc 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8299,7 +8299,9 @@ class MappableExprsHandler { (Next == CE && MapType != OMPC_MAP_unknown)) { if (!IsMappingWholeStruct) { CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr); - CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF)); + CombinedInfo.BasePointers.push_back(IsNonContiguous + ? LB.emitRawPointer(CGF) + : BP.emitRawPointer(CGF)); CombinedInfo.DevicePtrDecls.push_back(nullptr); CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None); CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF)); @@ -8407,7 +8409,7 @@ class MappableExprsHandler { break; // The pointer becomes the base for the next element. - if (Next != CE) + if (Next != CE && !IsNonContiguous) BP = IsMemberReference ? LowestElem : LB; if (!IsPartialMapped) IsExpressionFirstInfo = false; >From 761cbba6df6575dc4f53e264f40ffb2299de5ad0 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Tue, 20 Jan 2026 06:16:50 -0500 Subject: [PATCH 4/6] tests added --- ..._update_strided_struct_ptr_messages_from.c | 40 +++++++++ ...trided_struct_ptr_multiple_messages_from.c | 47 ++++++++++ ...strided_struct_ptr_partial_messages_from.c | 32 +++++++ .../target_update_strided_struct_ptr_from.c | 87 +++++++++++++++++++ ..._update_strided_struct_ptr_multiple_from.c | 81 +++++++++++++++++ ...t_update_strided_struct_ptr_partial_from.c | 67 ++++++++++++++ 6 files changed, 354 insertions(+) create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_partial_from.c diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c new file mode 100644 index 0000000000000..d86ce9e89766b --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c @@ -0,0 +1,40 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 16 +typedef struct { + double *data; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + s.data = (double *)__builtin_alloca(N * sizeof(double)); + + // Valid strided array sections with pointer member + #pragma omp target update from(s.data[0:4:2]) // OK + {} + + #pragma omp target update from(s.data[1:3:2]) // OK + {} + + // Missing stride (default = 1) + #pragma omp target update from(s.data[0:4]) // OK + {} + + // Invalid stride expressions + #pragma omp target update from(s.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(s.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'}} + + // Missing colon + #pragma omp target update from(s.data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} + + // Too many colons + #pragma omp target update from(s.data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} 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_ptr_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c new file mode 100644 index 0000000000000..7020ccb77d231 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c @@ -0,0 +1,47 @@ +// 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; + int len; +} T; + +int main(int argc, char **argv) { + T s1, s2; + s1.len = N; + s1.data = (double *)__builtin_alloca(N * sizeof(double)); + s2.len = N; + s2.data = (double *)__builtin_alloca(N * sizeof(double)); + + // Multiple valid strided updates + #pragma omp target update from(s1.data[0:10:2], s2.data[0:7:3]) // OK + {} + + // Mixed: one with stride, one without + #pragma omp target update from(s1.data[0:N], s2.data[0:5:2]) // OK + {} + + int stride1 = 2; + int stride2 = 3; + + // Multiple with expression strides + #pragma omp target update from(s1.data[1:5:stride1], s2.data[0:4:stride2]) // OK + {} + + // One valid, one invalid + #pragma omp target update from(s1.data[0:5:2], s2.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + #pragma omp target update from(s1.data[0:5:-1], s2.data[0:4:2]) // expected-error {{section stride is evaluated to a non-positive value -1}} + + #pragma omp target update from(s1.data[0:5:0], s2.data[0:4:1]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + // Syntax errors + #pragma omp target update from(s1.data[0:5:2], s2.data[0:4 3]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + {} + + #pragma omp target update from(s1.data[0:5:2:3], s2.data[0:4:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + {} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c new file mode 100644 index 0000000000000..4c835d3bef6f0 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized +// expected-no-diagnostics + +#define N 24 +typedef struct { + double *data; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + s.data = (double *)__builtin_alloca(N * sizeof(double)); + + // Valid partial strided updates with pointer member + #pragma omp target update from(s.data[0:2:10]) // OK - partial coverage + {} + + // Stride larger than length + #pragma omp target update from(s.data[0:2:20]) // OK + {} + + // Valid: complex expressions + int offset = 1; + + // Runtime-dependent stride expressions + #pragma omp target update from(s.data[0:4:offset+1]) // OK + {} + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_ptr_from.c b/offload/test/offloading/target_update_strided_struct_ptr_from.c new file mode 100644 index 0000000000000..9785e494ed2ec --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_ptr_from.c @@ -0,0 +1,87 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update from(s.data[0:s.len/2:2]) +// correctly updates every second element (stride 2) from the device to the host +// using a struct with pointer-to-array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +typedef struct { + double *data; + int 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; + s.data = (double *)calloc(N, sizeof(double)); + + printf("original host array values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(mapper(custom), to : s) + { +// Execute on device - modify even-indexed elements +#pragma omp target + { + for (int i = 0; i < s.len; i += 2) { + s.data[i] = 10.0; + } + } + +// Update only even indices (0,2,4,6,8,10,12,14) - s.len/2 elements with stride +// 2 +#pragma omp target update from(s.data[0 : s.len / 2 : 2]) + } + + printf("device array values after update from:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + + // CHECK: original host array values: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 0.0 + + // CHECK: device array values after update from: + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 0.0 + + free(s.data); + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c b/offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c new file mode 100644 index 0000000000000..6ad84c39b717b --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c @@ -0,0 +1,81 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that multiple strided target updates work correctly +// with struct containing pointer-to-array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 12 + +typedef struct { + double *data; + int len; +} T; + +#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len]) + +int main() { + T s1, s2; + s1.len = N; + s1.data = (double *)calloc(N, sizeof(double)); + s2.len = N; + s2.data = (double *)calloc(N, sizeof(double)); + + printf("original s1 values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s1.data[i]); + printf("\n"); + + printf("original s2 values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s2.data[i]); + printf("\n\n"); + +#pragma omp target data map(mapper(custom), to : s1, s2) + { +// Modify on device +#pragma omp target + { + // s1: set even indices to 10 + for (int i = 0; i < s1.len; i += 2) { + s1.data[i] = 10.0; + } + // s2: set multiples of 3 to 10 + for (int i = 0; i < s2.len; i += 3) { + s2.data[i] = 10.0; + } + } + +// Multiple strided updates: s1 even (s1.len/2 elements, stride 2), s2 multiples +// of 3 (s2.len/3 elements, stride 3) +#pragma omp target update from(s1.data[0 : s1.len / 2 : 2], \ + s2.data[0 : s2.len / 3 : 3]) + } + + printf("s1 after update (even indices):\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s1.data[i]); + printf("\n"); + + printf("s2 after update (multiples of 3):\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s2.data[i]); + printf("\n"); + + // CHECK: original s1 values: + // CHECK-NEXT: 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 + + // CHECK: original s2 values: + // CHECK-NEXT: 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 + + // CHECK: s1 after update (even indices): + // CHECK-NEXT: 10.0 0.0 10.0 0.0 10.0 0.0 10.0 0.0 10.0 0.0 10.0 0.0 + + // CHECK: s2 after update (multiples of 3): + // CHECK-NEXT: 10.0 0.0 0.0 10.0 0.0 0.0 10.0 0.0 0.0 10.0 0.0 0.0 + + free(s1.data); + free(s2.data); + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_ptr_partial_from.c b/offload/test/offloading/target_update_strided_struct_ptr_partial_from.c new file mode 100644 index 0000000000000..192a8caf7c125 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_ptr_partial_from.c @@ -0,0 +1,67 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update from(s.data[0:N/5:5]) +// correctly updates partial strided elements (stride larger than update count) +// from device to host using a struct with pointer-to-array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 20 + +typedef struct { + double *data; + int 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; + s.data = (double *)calloc(N, sizeof(double)); + + printf("original host array values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s.data[i]); + printf("\n\n"); + +#pragma omp target data map(mapper(custom), tofrom : s) + { +// Set all elements to 20 on device +#pragma omp target map(mapper(custom), tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = 20.0; // Set all to 20 on device + } + } + +// Modify specific elements on device (only first 4 stride positions) +#pragma omp target map(mapper(custom), tofrom : s) + { + s.data[0] = 10.0; + s.data[5] = 10.0; + s.data[10] = 10.0; + s.data[15] = 10.0; + } + +// Update indices 0, 5, 10, 15 only (N/5 = 4 elements with stride 5) +#pragma omp target update from(s.data[0 : N / 5 : 5]) + } + + printf("device array values after partial stride update:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s.data[i]); + printf("\n"); + + // CHECK: original host array values: + // CHECK-NEXT: 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 + // 0.0 0.0 0.0 0.0 + + // CHECK: device array values after partial stride update: + // CHECK-NEXT: 10.0 0.0 0.0 0.0 0.0 10.0 0.0 0.0 0.0 0.0 10.0 0.0 0.0 0.0 + // 0.0 10.0 0.0 0.0 0.0 0.0 + + free(s.data); + return 0; +} >From 487cdedc995d72029954fe34407d06b769bd49e0 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Tue, 20 Jan 2026 06:18:41 -0500 Subject: [PATCH 5/6] refined pointer fix --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index f4ae2d7d261bc..acb429c3ce7b3 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8299,9 +8299,7 @@ class MappableExprsHandler { (Next == CE && MapType != OMPC_MAP_unknown)) { if (!IsMappingWholeStruct) { CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr); - CombinedInfo.BasePointers.push_back(IsNonContiguous - ? LB.emitRawPointer(CGF) - : BP.emitRawPointer(CGF)); + CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF)); CombinedInfo.DevicePtrDecls.push_back(nullptr); CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None); CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF)); >From 75f49b84d69f8be7c3fa4ea803782f853d667cbc Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Tue, 20 Jan 2026 06:26:13 -0500 Subject: [PATCH 6/6] all tests covered --- ...rget_update_strided_struct_ptr_codegen.cpp | 73 ++++++++++++ ...et_update_strided_struct_ptr_messages_to.c | 40 +++++++ ..._strided_struct_ptr_multiple_messages_to.c | 47 ++++++++ ...e_strided_struct_ptr_partial_messages_to.c | 32 ++++++ ...et_update_strided_struct_ptr_multiple_to.c | 101 +++++++++++++++++ ...get_update_strided_struct_ptr_partial_to.c | 76 +++++++++++++ .../target_update_strided_struct_ptr_to.c | 106 ++++++++++++++++++ 7 files changed, 475 insertions(+) create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c create mode 100644 clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_partial_to.c create mode 100644 offload/test/offloading/target_update_strided_struct_ptr_to.c diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp b/clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp new file mode 100644 index 0000000000000..a2e5978a81eb7 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp @@ -0,0 +1,73 @@ +// Test codegen for strided target update with struct containing pointer-to-array member +// RUN: %clang_cc1 -DCK27 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK27 +// RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK27 + +// RUN: %clang_cc1 -DCK27 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK27 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} + +// expected-no-diagnostics + +#ifdef CK27 +#ifndef CK27_INCLUDED +#define CK27_INCLUDED + +// Verify that non-contiguous map type flag is set (bit 48) +// 17592186044418 = 0x1000000000002 (OMP_MAP_NON_CONTIG | OMP_MAP_FROM) +// 17592186044417 = 0x1000000000001 (OMP_MAP_NON_CONTIG | OMP_MAP_TO) +// CK27-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 17592186044418] +// CK27-DAG: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 17592186044417] + +struct T { + double *data; + int len; +}; + +// CK27-LABEL: define {{.*}}void @{{.*}}test_strided_update_from{{.*}}( +void test_strided_update_from(int arg) { + T s; + s.len = 16; + s.data = new double[16]; + + for (int i = 0; i < 16; i++) { + s.data[i] = i; + } + + // Verify the stride descriptor is created with correct values: + // - offset = 0 + // - count = 4 (number of elements to update) + // - stride = 16 (2 * sizeof(double) = 2 * 8 = 16 bytes) + // CK27-DAG: store i64 0, ptr %{{.+}}, align 8 + // CK27-DAG: store i64 4, ptr %{{.+}}, align 8 + // CK27-DAG: store i64 16, ptr %{{.+}}, align 8 + + // Verify __tgt_target_data_update_mapper is called + // CK27: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 {{1|signext 1}}, ptr %{{.+}}, ptr %{{.+}}, ptr @{{.+}}, ptr @.offload_maptypes{{.*}}, ptr null, ptr null) + + #pragma omp target update from(s.data[0:4:2]) + + delete[] s.data; +} + +// CK27-LABEL: define {{.*}}void @{{.*}}test_strided_update_to{{.*}}( +void test_strided_update_to(int arg) { + T s; + s.len = 16; + s.data = new double[16]; + + for (int i = 0; i < 16; i++) { + s.data[i] = i; + } + + // Verify __tgt_target_data_update_mapper is called with TO map type + // CK27: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 {{1|signext 1}}, ptr %{{.+}}, ptr %{{.+}}, ptr @{{.+}}, ptr @.offload_maptypes{{.*}}, ptr null, ptr null) + + #pragma omp target update to(s.data[0:4:2]) + + delete[] s.data; +} + +#endif // CK27_INCLUDED +#endif // CK27 diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c new file mode 100644 index 0000000000000..012f484a6ae36 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c @@ -0,0 +1,40 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 16 +typedef struct { + double *data; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + s.data = (double *)__builtin_alloca(N * sizeof(double)); + + // Valid strided array sections with pointer member + #pragma omp target update to(s.data[0:4:2]) // OK + {} + + #pragma omp target update to(s.data[1:3:2]) // OK + {} + + // Missing stride (default = 1) + #pragma omp target update to(s.data[0:4]) // OK + {} + + // Invalid stride expressions + #pragma omp target update to(s.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 to(s.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'}} + + // Missing colon + #pragma omp target update to(s.data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} + + // Too many colons + #pragma omp target update to(s.data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} 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_ptr_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c new file mode 100644 index 0000000000000..05278308ad173 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c @@ -0,0 +1,47 @@ +// 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; + int len; +} T; + +int main(int argc, char **argv) { + T s1, s2; + s1.len = N; + s1.data = (double *)__builtin_alloca(N * sizeof(double)); + s2.len = N; + s2.data = (double *)__builtin_alloca(N * sizeof(double)); + + // Multiple valid strided updates (to clause) + #pragma omp target update to(s1.data[0:10:2], s2.data[0:7:3]) // OK + {} + + // Mixed: one with stride, one without + #pragma omp target update to(s1.data[0:N], s2.data[0:5:2]) // OK + {} + + int stride1 = 2; + int stride2 = 3; + + // Multiple with expression strides + #pragma omp target update to(s1.data[1:5:stride1], s2.data[0:4:stride2]) // OK + {} + + // One valid, one invalid + #pragma omp target update to(s1.data[0:5:2], s2.data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + #pragma omp target update to(s1.data[0:5:-1], s2.data[0:4:2]) // expected-error {{section stride is evaluated to a non-positive value -1}} + + #pragma omp target update to(s1.data[0:5:0], s2.data[0:4:1]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + // Syntax errors + #pragma omp target update to(s1.data[0:5:2], s2.data[0:4 3]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + {} + + #pragma omp target update to(s1.data[0:5:2:3], s2.data[0:4:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + {} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c new file mode 100644 index 0000000000000..d62a6c640d0b3 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized +// expected-no-diagnostics + +#define N 24 +typedef struct { + double *data; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + s.data = (double *)__builtin_alloca(N * sizeof(double)); + + // Valid partial strided updates with pointer member (to clause) + #pragma omp target update to(s.data[0:2:10]) // OK - partial coverage + {} + + // Stride larger than length + #pragma omp target update to(s.data[0:2:20]) // OK + {} + + // Valid: complex expressions + int offset = 1; + + // Runtime-dependent stride expressions + #pragma omp target update to(s.data[0:4:offset+1]) // OK + {} + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c b/offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c new file mode 100644 index 0000000000000..267c9f1db1b9f --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c @@ -0,0 +1,101 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that multiple strided target updates to device work +// correctly with struct containing pointer-to-array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 12 + +typedef struct { + double *data; + int len; +} T; + +#pragma omp declare mapper(custom : T v) \ + map(tofrom : v, v.len, v.data[0 : v.len]) + +int main() { + T s1, s2; + s1.len = N; + s1.data = (double *)calloc(N, sizeof(double)); + s2.len = N; + s2.data = (double *)calloc(N, sizeof(double)); + + // Initialize structs on host + for (int i = 0; i < N; i++) { + s1.data[i] = i; + s2.data[i] = i; + } + + printf("original s1 values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s1.data[i]); + printf("\n"); + + printf("original s2 values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s2.data[i]); + printf("\n\n"); + +#pragma omp target data map(tofrom : s1, s2) + { +// Initialize device struct arrays to 20 +#pragma omp target + { + for (int i = 0; i < s1.len; i++) { + s1.data[i] = 20.0; + s2.data[i] = 20.0; + } + } + + // Modify host: s1 even indices, s2 multiples of 3 + for (int i = 0; i < s1.len; i += 2) { + s1.data[i] = 10.0; + } + for (int i = 0; i < s2.len; i += 3) { + s2.data[i] = 10.0; + } + +// Multiple strided updates to device: s1 even (s1.len/2 elements, stride 2), s2 +// multiples of 3 (s2.len/3 elements, stride 3) +#pragma omp target update to(s1.data[0 : s1.len / 2 : 2], \ + s2.data[0 : s2.len / 3 : 3]) + +// Verify update on device by adding 5 +#pragma omp target + { + for (int i = 0; i < s1.len; i++) { + s1.data[i] += 5.0; + s2.data[i] += 5.0; + } + } + } // Exit target data - tofrom mapper copies data back + + printf("s1 after update to device (even indices):\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s1.data[i]); + printf("\n"); + + printf("s2 after update to device (multiples of 3):\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s2.data[i]); + printf("\n"); + + // CHECK: original s1 values: + // CHECK-NEXT: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0 + + // CHECK: original s2 values: + // CHECK-NEXT: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0 + + // CHECK: s1 after update to device (even indices): + // CHECK-NEXT: 15.0 25.0 15.0 25.0 15.0 25.0 15.0 25.0 15.0 25.0 15.0 25.0 + + // CHECK: s2 after update to device (multiples of 3): + // CHECK-NEXT: 15.0 25.0 25.0 15.0 25.0 25.0 15.0 25.0 25.0 15.0 25.0 25.0 + + free(s1.data); + free(s2.data); + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_ptr_partial_to.c b/offload/test/offloading/target_update_strided_struct_ptr_partial_to.c new file mode 100644 index 0000000000000..35772f7c642e4 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_ptr_partial_to.c @@ -0,0 +1,76 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update to(s.data[0:N/5:5]) correctly +// updates partial strided elements (stride larger than update count) from host +// to device using a struct with pointer-to-array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 20 + +typedef struct { + double *data; + int len; +} T; + +#pragma omp declare mapper(custom : T v) \ + map(tofrom : v, v.len, v.data[0 : v.len]) + +int main() { + T s; + s.len = N; + s.data = (double *)calloc(N, sizeof(double)); + + // Initialize struct data on host + 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("%.1f ", s.data[i]); + printf("\n\n"); + +#pragma omp target data map(tofrom : s) + { +// Initialize device struct arrays to 20 +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] = 20.0; + } + } + + // Modify host elements: indices 0, 5, 10, 15 only + s.data[0] = 10.0; + s.data[5] = 10.0; + s.data[10] = 10.0; + s.data[15] = 10.0; + +// Update indices 0, 5, 10, 15 only (N/5 = 4 elements with stride 5) to device +#pragma omp target update to(s.data[0 : N / 5 : 5]) + +// Execute on device - add 5 to verify update worked +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] += 5.0; + } + } + } // Exit target data - tofrom mapper copies data back + + printf("device array values after partial stride update to:\n"); + for (int i = 0; i < N; i++) + printf("%.1f ", s.data[i]); + printf("\n"); + + // CHECK: original host array values: + // CHECK-NEXT: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0 12.0 13.0 14.0 15.0 16.0 17.0 18.0 19.0 + + // CHECK: device array values after partial stride update to: + // CHECK-NEXT: 15.0 25.0 25.0 25.0 25.0 15.0 25.0 25.0 25.0 25.0 15.0 25.0 25.0 25.0 25.0 15.0 25.0 25.0 25.0 25.0 + + free(s.data); + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_ptr_to.c b/offload/test/offloading/target_update_strided_struct_ptr_to.c new file mode 100644 index 0000000000000..690802e202065 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_ptr_to.c @@ -0,0 +1,106 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update to(s.data[0:s.len/2:2]) +// correctly updates every second element (stride 2) from the host to the device +// using a struct with pointer-to-array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +typedef struct { + double *data; + int len; +} T; + +#pragma omp declare mapper(custom : T v) \ + map(tofrom : v, v.len, v.data[0 : v.len]) + +int main() { + T s; + s.len = N; + s.data = (double *)calloc(N, sizeof(double)); + + // Initialize struct data on host + 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("%.1f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : s) + { +// Set device data to 20 +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] = 20.0; + } + } + + // Modify host even-indexed elements + for (int i = 0; i < N; i += 2) { + s.data[i] = 10.0; + } + +// Update only even indices (0,2,4,6,8,10,12,14) to device - s.len/2 elements +// with stride 2 +#pragma omp target update to(s.data[0 : s.len / 2 : 2]) + +// Execute on device - add 5 to verify update worked +#pragma omp target + { + for (int i = 0; i < s.len; i++) { + s.data[i] += 5.0; + } + } + } // Exit target data - tofrom mapper copies data back + + printf("device array values after update to:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + + // CHECK: original host 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: device array values after update to: + // 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 + + free(s.data); + return 0; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
