https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/176914
>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/7] 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/7] 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/7] 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/7] 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/7] 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/7] 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; +} >From bc0f93420faf245ea2c9849c8fe4d57fead73911 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Tue, 20 Jan 2026 07:44:30 -0500 Subject: [PATCH 7/7] formatted --- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 578e63351ca5f..6f44a9649abeb 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9776,8 +9776,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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
