llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-offload Author: Amit Tiwari (amitamd7) <details> <summary>Changes</summary> **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` --- Patch is 32.17 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/176914.diff 15 Files Affected: - (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+17-1) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_codegen.cpp (+73) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_messages_from.c (+40) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_messages_to.c (+40) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_from.c (+47) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_multiple_messages_to.c (+47) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_from.c (+32) - (added) clang/test/OpenMP/target_update_strided_struct_ptr_partial_messages_to.c (+32) - (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+21-8) - (added) offload/test/offloading/target_update_strided_struct_ptr_from.c (+87) - (added) offload/test/offloading/target_update_strided_struct_ptr_multiple_from.c (+81) - (added) offload/test/offloading/target_update_strided_struct_ptr_multiple_to.c (+101) - (added) offload/test/offloading/target_update_strided_struct_ptr_partial_from.c (+67) - (added) offload/test/offloading/target_update_strided_struct_ptr_partial_to.c (+76) - (added) offload/test/offloading/target_update_strided_struct_ptr_to.c (+106) ``````````diff diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8981a0de6d0e4..acb429c3ce7b3 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(); }); @@ -8391,7 +8407,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; 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_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_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_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_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_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/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/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; } } 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. + +#inc... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/176914 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
