Author: Amit Tiwari Date: 2025-12-23T12:57:12+05:30 New Revision: 230b437d05c080b54f74708d6c09d4ccf568951d
URL: https://github.com/llvm/llvm-project/commit/230b437d05c080b54f74708d6c09d4ccf568951d DIFF: https://github.com/llvm/llvm-project/commit/230b437d05c080b54f74708d6c09d4ccf568951d.diff LOG: [Clang][OpenMP] Handle check for non-contiguous mapping in pointer-based array sections (#157443) ### 1. ElementType deduction for pointer-based array sections Problem: Pointer-based array sections were previously ignored during `ElementType` deduction, leading to incorrect assumptions about array item types. This often resulted in out-of-bounds access, as seen in the assertion failure: ``` Assertion `idx < size()' failed. llvm-project/llvm/include/llvm/ADT/SmallVector.h:292: reference llvm::SmallVectorTemplateCommon<llvm::Value *>::operatorsize_type [T = llvm::Value *] ``` Fix: Added a check in clang/lib/CodeGen/CGOpenMPRuntime.cpp to ensure `ElementType` is correctly detected for cases involving non-contiguous updates with a base pointer. Impact: Resolves failures in OpenMP_VV (formerly sollve_vv) and other offload/clang-OpenMP tests: All tests under: https://github.com/OpenMP-Validation-and-Verification/OpenMP_VV/tree/master/tests/5.0/target_update test_target_update_mapper_from_discontiguous.c test_target_update_mapper_to_discontiguous.c test_target_update_to_discontiguous.c test_target_update_from_discontiguous.c ### 2. Zero-dimension propagation in struct member mappings Problem: A zero-dimension entry for struct members introduced inconsistencies in complex mapping logic within OMPIRBuilder.cpp. Placeholder zeros propagated to emitNonContiguousDescriptor(), breaking reverse indexing logic and corrupting IR: Loops assume `Dims[I] >= 1`. When `Dims[I] == 0`: Reverse indexing still stores pointers to uninitialized allocas or mismatched slots. Runtime interprets `ArgSizes[I]` (derived from `Dims[I])` as dimensionality, causing size/offset calculations to collapse to zero → results in `size=0` async copy and plugin interface errors. Fix: Prepend a synthetic dimension of size 1 instead of appending a zero, preserving correctness in `targetDataUpdate()` for non-contiguous updates. Impact: Added dedicated test cases that previously failed on main. Added: clang/test/OpenMP/target_update_strided_ptr_messages_from.c clang/test/OpenMP/target_update_strided_ptr_messages_to.c clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c clang/test/OpenMP/target_update_strided_struct_messages_from.c clang/test/OpenMP/target_update_strided_struct_messages_to.c clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.c clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c offload/test/offloading/strided_ptr_multiple_update_from.c offload/test/offloading/strided_ptr_multiple_update_to.c offload/test/offloading/strided_ptr_partial_update_from.c offload/test/offloading/strided_ptr_partial_update_to.c offload/test/offloading/target_update_from.c offload/test/offloading/target_update_strided_struct_from.c offload/test/offloading/target_update_strided_struct_multiple_from.c offload/test/offloading/target_update_strided_struct_multiple_to.c offload/test/offloading/target_update_strided_struct_partial_from.c offload/test/offloading/target_update_strided_struct_partial_to.c offload/test/offloading/target_update_strided_struct_to.c offload/test/offloading/target_update_to.c Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/target_update_codegen.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index bf3af0571d7ac..b8ee701c482bb 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8334,13 +8334,25 @@ class MappableExprsHandler { ElementType = CAT->getElementType().getTypePtr(); else if (VAT) ElementType = VAT->getElementType().getTypePtr(); - else - assert(&Component == &*Components.begin() && - "Only expect pointer (non CAT or VAT) when this is the " - "first Component"); - // If ElementType is null, then it means the base is a pointer - // (neither CAT nor VAT) and we'll attempt to get ElementType again - // for next iteration. + else if (&Component == &*Components.begin()) { + // If the base is a raw pointer (e.g. T *data with data[a:b:c]), + // there was no earlier CAT/VAT/array handling to establish + // ElementType. Capture the pointee type now so that subsequent + // components (offset/length/stride) have a concrete element type to + // work with. This makes pointer-backed sections behave consistently + // with CAT/VAT/array bases. + if (const auto *PtrType = Ty->getAs<PointerType>()) + ElementType = PtrType->getPointeeType().getTypePtr(); + } else { + // Any component after the first should never have a raw pointer type; + // by this point. ElementType must already be known (set above or in + // prior array / CAT / VAT handling). + assert(!Ty->isPointerType() && + "Non-first components should not be raw pointers"); + } + + // At this stage, if ElementType was a base pointer and we are in the + // first iteration, it has been computed. if (ElementType) { // For the case that having pointer as base, we need to remove one // level of indirection. @@ -9090,12 +9102,20 @@ class MappableExprsHandler { // If there is an entry in PartialStruct it means we have a struct with // individual members mapped. Emit an extra combined entry. if (PartialStruct.Base.isValid()) { - GroupUnionCurInfo.NonContigInfo.Dims.push_back(0); + // Prepend a synthetic dimension of length 1 to represent the + // aggregated struct object. Using 1 (not 0, as 0 produced an + // incorrect non-contiguous descriptor (DimSize==1), causing the + // non-contiguous motion clause path to be skipped.) is important: + // * It preserves the correct rank so targetDataUpdate() computes + // DimSize == 2 for cases like strided array sections originating + // from user-defined mappers (e.g. test with s.data[0:8:2]). + GroupUnionCurInfo.NonContigInfo.Dims.insert( + GroupUnionCurInfo.NonContigInfo.Dims.begin(), 1); emitCombinedEntry( CurInfo, GroupUnionCurInfo.Types, PartialStruct, AttachInfo, - /*IsMapThis*/ !VD, OMPBuilder, VD, + /*IsMapThis=*/!VD, OMPBuilder, VD, /*OffsetForMemberOfFlag=*/CombinedInfo.BasePointers.size(), - /*NotTargetParam=*/true); + /*NotTargetParams=*/true); } // Append this group's results to the overall CurInfo in the correct diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp index 4d9c49ebcf03a..df238f97f608c 100644 --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -1187,7 +1187,7 @@ void foo(int arg) { // CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x ptr]]] } // CK21: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } -// CK21: [[SIZE:@.+]] = private unnamed_addr constant [2 x i64] zeroinitializer +// CK21: [[SIZE:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4] // CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 299067162755073] struct ST { @@ -1234,8 +1234,9 @@ struct ST { // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPSZ:%.+]], ptr [[MTYPE]]{{.+}}) // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0 - // CK21-DAG: store ptr [[DIMS]], ptr [[PTRS]], + // CK21-DAG: store ptr [[DIMS]], ptr [[PTRS1:%.+]], + // CK21-DAG: [[PTRS1]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1 + // CK21: ret void #pragma omp target update to(dptr[0:2][1:3][0:4]) } }; diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c new file mode 100644 index 0000000000000..c7e8a9efafc9b --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_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 + +int main(int argc, char **argv) { + int len = 16; + double *data; + + // Valid strided array sections with FROM + #pragma omp target update from(data[0:8:2]) // OK - even indices + {} + + #pragma omp target update from(data[1:4:3]) // OK - odd start with stride + {} + + #pragma omp target update from(data[2:3:5]) // OK - large stride + {} + + // Missing stride (default = 1) + #pragma omp target update from(data[0:8]) // OK - default stride + {} + + #pragma omp target update from(data[4:len-4]) // OK - computed length + {} + + // Invalid stride expressions + #pragma omp target update from(data[0:8:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:5:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Syntax errors + #pragma omp target update from(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'}} + {} + + #pragma omp target update from(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; +} \ No newline at end of file diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c new file mode 100644 index 0000000000000..47a549c9bc0dc --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_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 + +int main(int argc, char **argv) { + int len = 16; + double *data; + + // Valid strided array sections with TO + #pragma omp target update to(data[0:8:2]) // OK - even indices + {} + + #pragma omp target update to(data[1:4:3]) // OK - odd start with stride + {} + + #pragma omp target update to(data[2:3:5]) // OK - large stride + {} + + // Missing stride (default = 1) + #pragma omp target update to(data[0:8]) // OK - default stride + {} + + #pragma omp target update to(data[4:len-4]) // OK - computed length + {} + + // Invalid stride expressions + #pragma omp target update to(data[0:8: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(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'}} + + #pragma omp target update to(data[1:5:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Syntax errors + #pragma omp target update to(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'}} + {} + + #pragma omp target update to(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; +} \ No newline at end of file diff --git a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c new file mode 100644 index 0000000000000..af5dfec1b7941 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 12; + double *data; + double *data1; + double *data2; + + // Valid multiple strided array sections + #pragma omp target update from(data1[0:6:2], data2[0:4:3]) // OK - diff erent strides + {} + + #pragma omp target update from(data1[1:2:3], data2[2:3:2]) // OK - with offsets + {} + + // Mixed strided and regular sections + #pragma omp target update from(data1[0:len], data2[0:4:2]) // OK - mixed + {} + + #pragma omp target update from(data1[1:3:2], data2[0:len]) // OK - reversed mix + {} + + // Using the single data pointer with strides + #pragma omp target update from(data[0:4:2]) // OK - single pointer + {} + + // Invalid stride in one of multiple sections + #pragma omp target update from(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + #pragma omp target update from(data1[0:3:-1], data2[0:2:2]) // expected-error {{section stride is evaluated to a non-positive value -1}} + + #pragma omp target update from(data[0:4:0], data1[0:2:1]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + // Complex expressions in multiple arrays + int stride1 = 2, stride2 = 3; + #pragma omp target update from(data1[1:4:stride1+1], data2[0:3:stride2-1]) // OK - expressions + {} + + // Mix all three pointers + #pragma omp target update from(data[0:2:3], data1[1:3:2], data2[2:2:4]) // OK - three arrays + {} + + // Syntax errors in multiple arrays + #pragma omp target update from(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + #pragma omp target update from(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + #pragma omp target update from(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + return 0; +} \ No newline at end of file diff --git a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c new file mode 100644 index 0000000000000..08c7581598612 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 12; + double *data; + double *data1; + double *data2; + + // Valid multiple strided array sections + #pragma omp target update to(data1[0:6:2], data2[0:4:3]) // OK - diff erent strides + {} + + #pragma omp target update to(data1[1:2:3], data2[2:3:2]) // OK - with offsets + {} + + // Mixed strided and regular sections + #pragma omp target update to(data1[0:len], data2[0:4:2]) // OK - mixed + {} + + #pragma omp target update to(data1[1:3:2], data2[0:len]) // OK - reversed mix + {} + + // Using the single data pointer with strides + #pragma omp target update to(data[0:4:2]) // OK - single pointer + {} + + // Invalid stride in one of multiple sections + #pragma omp target update to(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + #pragma omp target update to(data1[0:3:-1], data2[0:2:2]) // expected-error {{section stride is evaluated to a non-positive value -1}} + + #pragma omp target update to(data[0:4:0], data1[0:2:1]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + // Complex expressions in multiple arrays + int stride1 = 2, stride2 = 3; + #pragma omp target update to(data1[1:4:stride1+1], data2[0:3:stride2-1]) // OK - expressions + {} + + // Mix all three pointers + #pragma omp target update to(data[0:2:3], data1[1:3:2], data2[2:2:4]) // OK - three arrays + {} + + // Syntax errors in multiple arrays + #pragma omp target update to(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + #pragma omp target update to(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + #pragma omp target update to(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + return 0; +} \ No newline at end of file diff --git a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c new file mode 100644 index 0000000000000..5937184d1f358 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c @@ -0,0 +1,37 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 11; + double *data; + + // Valid partial strided sections with FROM + #pragma omp target update from(data[0:2:3]) // OK - partial coverage + {} + + #pragma omp target update from(data[1:3:4]) // OK - offset with partial stride + {} + + #pragma omp target update from(data[2:2:5]) // OK - large partial stride + {} + + // Stride larger than remaining elements + #pragma omp target update from(data[0:2:10]) // OK - stride > array size + {} + + #pragma omp target update from(data[0:3:len]) // OK - stride = len + {} + + // Complex expressions + int offset = 1; + int stride = 2; + + // Runtime-dependent invalid strides + #pragma omp target update from(data[0:4:offset-1]) // OK if offset > 1 + {} + + // Compile-time invalid strides + #pragma omp target update from(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} \ No newline at end of file diff --git a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c new file mode 100644 index 0000000000000..848ef95509f1c --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 11; + double *data; + + // Valid partial strided sections with TO + #pragma omp target update to(data[0:2:3]) // OK - partial coverage + {} + + #pragma omp target update to(data[1:3:4]) // OK - offset with partial stride + {} + + #pragma omp target update to(data[2:2:5]) // OK - large partial stride + {} + + // Stride larger than remaining elements + #pragma omp target update to(data[0:2:10]) // OK - stride > array size + {} + + #pragma omp target update to(data[0:3:len]) // OK - stride = len + {} + + int offset = 1; + int stride = 2; + + // Runtime-dependent invalid strides + #pragma omp target update to(data[0:4:offset-1]) // OK if offset > 1 + {} + + // Compile-time invalid strides + #pragma omp target update to(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} \ No newline at end of file diff --git a/clang/test/OpenMP/target_update_strided_struct_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_messages_from.c new file mode 100644 index 0000000000000..b70a2453dc122 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_messages_from.c @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 8 +typedef struct { + double data[N]; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + + // Valid strided array sections + #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_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_messages_to.c new file mode 100644 index 0000000000000..39e774014fd44 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_messages_to.c @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 8 +typedef struct { + double data[N]; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + + // Valid strided array sections + #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_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c new file mode 100644 index 0000000000000..3005d5bb28702 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 12 +typedef struct { + double data[N]; + int len; +} T; + +int main(int argc, char **argv) { + T s1, s2; + s1.len = N; + s2.len = N; + + // Valid multiple strided array sections + #pragma omp target update from(s1.data[0:4:2], s2.data[0:2:5]) // OK + {} + + // Mixed strided and regular array sections + #pragma omp target update from(s1.data[0:s1.len], s2.data[0:4:2]) // OK + {} + + // Invalid stride in one of multiple sections + #pragma omp target update from(s1.data[0:3:4], s2.data[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} + + // Missing colon + #pragma omp target update from(s1.data[0:4:2], s2.data[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.c new file mode 100644 index 0000000000000..c22c903472ce2 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_multiple_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 + +#define N 12 +typedef struct { + double data[N]; + int len; + } T; + +int main(int argc, char **argv) { + T s1, s2; + s1.len = N; + s2.len = N; + + // Valid multiple strided array sections + #pragma omp target update to(s1.data[0:4:2], s2.data[0:2:5]) // OK + {} + + // Mixed strided and regular array sections + #pragma omp target update to(s1.data[0:s1.len], s2.data[0:4:2]) // OK + {} + + // Invalid stride in one of multiple sections + #pragma omp target update to(s1.data[0:3:4], s2.data[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} + {} + + // Missing colon + #pragma omp target update to(s1.data[0:4:2], s2.data[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}} + {} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c new file mode 100644 index 0000000000000..0ae5790680437 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c @@ -0,0 +1,31 @@ +// 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 11 +typedef struct { + double data[N]; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + + // Valid partial strided updates + #pragma omp target update from(s.data[0:4:3]) // OK + {} + + // Stride larger than length + #pragma omp target update from(s.data[0:2:10]) // OK + {} + + // Valid: complex expressions + int offset = 1; + + // Invalid stride expressions (runtime-dependent) + #pragma omp target update from(s.data[0:4:offset-1]) // OK if offset > 1 + {} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c new file mode 100644 index 0000000000000..c003bbadaca5c --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c @@ -0,0 +1,31 @@ +// 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 11 +typedef struct { + double data[N]; + int len; +} T; + +int main(int argc, char **argv) { + T s; + s.len = N; + + // Valid partial strided updates + #pragma omp target update to(s.data[0:4:3]) // OK + {} + + // Stride larger than length + #pragma omp target update to(s.data[0:2:10]) // OK + {} + + // Valid: complex expressions + int offset = 1; + + // Potentially invalid stride expressions depending on runtime values + #pragma omp target update to(s.data[0:4:offset-1]) // OK if offset > 1 + {} + + return 0; +} diff --git a/offload/test/offloading/strided_ptr_multiple_update_from.c b/offload/test/offloading/strided_ptr_multiple_update_from.c new file mode 100644 index 0000000000000..b671a1e2b688d --- /dev/null +++ b/offload/test/offloading/strided_ptr_multiple_update_from.c @@ -0,0 +1,114 @@ +// This test checks that #pragma omp target update from(data1[0:6:2], +// data2[0:4:3]) correctly updates strided sections covering full arrays +// from the device to the host using dynamically allocated memory. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int len = 12; + double *data1 = (double *)calloc(len, sizeof(double)); + double *data2 = (double *)calloc(len, sizeof(double)); + +// Initial values +#pragma omp target map(tofrom : data1[0 : len], data2[0 : len]) + { + for (int i = 0; i < len; i++) { + data1[i] = i; + data2[i] = i * 10; + } + } + + printf("original host array values:\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data2[i]); + +#pragma omp target data map(to : data1[0 : len], data2[0 : len]) + { +#pragma omp target + { + for (int i = 0; i < len; i++) + data1[i] += i; + for (int i = 0; i < len; i++) + data2[i] += 100; + } + +// data1[0:6:2] covers all even indices: 0,2,4,6,8,10 (6 elements) +// data2[0:4:3] covers every 3rd: 0,3,6,9 (4 elements) +#pragma omp target update from(data1[0 : 6 : 2], data2[0 : 4 : 3]) + } + + printf("device array values after update from:\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data2[i]); + + // CHECK: original host array values: + // CHECK-NEXT: data1: + // 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: data2: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 30.0 + // CHECK-NEXT: 40.0 + // CHECK-NEXT: 50.0 + // CHECK-NEXT: 60.0 + // CHECK-NEXT: 70.0 + // CHECK-NEXT: 80.0 + // CHECK-NEXT: 90.0 + // CHECK-NEXT: 100.0 + // CHECK-NEXT: 110.0 + + // CHECK: device array values after update from: + // CHECK-NEXT: data1: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 4.0 + // CHECK-NEXT: 3.0 + // CHECK-NEXT: 8.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 12.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 16.0 + // CHECK-NEXT: 9.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 11.0 + // CHECK-NEXT: data2: + // CHECK-NEXT: 100.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 130.0 + // CHECK-NEXT: 40.0 + // CHECK-NEXT: 50.0 + // CHECK-NEXT: 160.0 + // CHECK-NEXT: 70.0 + // CHECK-NEXT: 80.0 + // CHECK-NEXT: 190.0 + // CHECK-NEXT: 100.0 + // CHECK-NEXT: 110.0 + + free(data1); + free(data2); + return 0; +} diff --git a/offload/test/offloading/strided_ptr_multiple_update_to.c b/offload/test/offloading/strided_ptr_multiple_update_to.c new file mode 100644 index 0000000000000..7ae9189ead8e3 --- /dev/null +++ b/offload/test/offloading/strided_ptr_multiple_update_to.c @@ -0,0 +1,130 @@ +// This test checks that #pragma omp target update to(data1[0:6:2], +// data2[0:4:3]) correctly updates strided sections covering full arrays +// from the host to the device using dynamically allocated memory. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int len = 12; + double *data1 = (double *)calloc(len, sizeof(double)); + double *data2 = (double *)calloc(len, sizeof(double)); + + // Initialize host arrays + for (int i = 0; i < len; i++) { + data1[i] = i; + data2[i] = i * 10; + } + + printf("original host array values:\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data2[i]); + +#pragma omp target data map(tofrom : data1[0 : len], data2[0 : len]) + { +// Initialize device arrays to 20 +#pragma omp target + { + for (int i = 0; i < len; i++) { + data1[i] = 20.0; + data2[i] = 20.0; + } + } + + // data1: even indices + for (int i = 0; i < 6; i++) { + data1[i * 2] = 10.0; + } + // data2: every 3rd index + for (int i = 0; i < 4; i++) { + data2[i * 3] = 10.0; + } + +// data1[0:6:2] updates all even indices: 0,2,4,6,8,10 (6 elements) +// data2[0:4:3] updates every 3rd: 0,3,6,9 (4 elements) +#pragma omp target update to(data1[0 : 6 : 2], data2[0 : 4 : 3]) + +// Verify on device by adding 5 +#pragma omp target + { + for (int i = 0; i < len; i++) { + data1[i] += 5.0; + data2[i] += 5.0; + } + } + } + + printf("device array values after update to:\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data2[i]); + + // CHECK: original host array values: + // CHECK-NEXT: data1: + // 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: data2: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 30.0 + // CHECK-NEXT: 40.0 + // CHECK-NEXT: 50.0 + // CHECK-NEXT: 60.0 + // CHECK-NEXT: 70.0 + // CHECK-NEXT: 80.0 + // CHECK-NEXT: 90.0 + // CHECK-NEXT: 100.0 + // CHECK-NEXT: 110.0 + + // CHECK: device array values after update to: + // CHECK-NEXT: data1: + // 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: data2: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + + free(data1); + free(data2); + return 0; +} diff --git a/offload/test/offloading/strided_ptr_partial_update_from.c b/offload/test/offloading/strided_ptr_partial_update_from.c new file mode 100644 index 0000000000000..d30a671393002 --- /dev/null +++ b/offload/test/offloading/strided_ptr_partial_update_from.c @@ -0,0 +1,68 @@ +// This test checks that #pragma omp target update from(data[0:2:3]) correctly +// updates every third element (stride 3) from the device to the host, partially +// across the array using dynamically allocated memory. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int len = 11; + double *data = (double *)calloc(len, sizeof(double)); + +#pragma omp target map(tofrom : data[0 : len]) + { + for (int i = 0; i < len; i++) + data[i] = i; + } + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + +#pragma omp target data map(to : data[0 : len]) + { +#pragma omp target + for (int i = 0; i < len; i++) + data[i] += i; + +#pragma omp target update from(data[0 : 2 : 3]) // indices 0,3 only + } + + printf("device array values after update from:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: original host array values: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 3.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 10.000000 + + // CHECK: device array values after update from: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 10.000000 + + free(data); + return 0; +} diff --git a/offload/test/offloading/strided_ptr_partial_update_to.c b/offload/test/offloading/strided_ptr_partial_update_to.c new file mode 100644 index 0000000000000..37c165351a760 --- /dev/null +++ b/offload/test/offloading/strided_ptr_partial_update_to.c @@ -0,0 +1,81 @@ +// This test checks that #pragma omp target update to(data[0:2:3]) correctly +// updates every third element (stride 3) from the host to the device, partially +// across the array using dynamically allocated memory. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +int main() { + int len = 11; + double *data = (double *)calloc(len, sizeof(double)); + + // Initialize on host + for (int i = 0; i < len; i++) + data[i] = i; + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : data[0 : len]) + { +// Initialize device array to 20 +#pragma omp target + { + for (int i = 0; i < len; i++) + data[i] = 20.0; + } + + // Modify host data for elements that will be updated + data[0] = 10.0; + data[3] = 10.0; + +// This creates Host→Device transfer for indices 0,3 only +#pragma omp target update to(data[0 : 2 : 3]) + +// Verify on device by adding 5 to all elements +#pragma omp target + { + for (int i = 0; i < len; i++) + data[i] += 5.0; + } + } + + printf("device array values after update to:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: original host array values: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 3.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 10.000000 + + // CHECK: device array values after update to: + // CHECK-NEXT: 15.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 15.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + // CHECK-NEXT: 25.000000 + + free(data); + return 0; +} diff --git a/offload/test/offloading/target_update_from.c b/offload/test/offloading/target_update_from.c new file mode 100644 index 0000000000000..4f84968f6d9a9 --- /dev/null +++ b/offload/test/offloading/target_update_from.c @@ -0,0 +1,74 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update from" clause in OpenMP supports strided +// sections. #pragma omp target update from(result[0:N/2:2]) updates every other +// element from device +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 32 + +int main() { + double *result = (double *)calloc(N, sizeof(double)); + + printf("initial host array values:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + printf("\n"); + +#pragma omp target data map(to : result[0 : N]) + { +#pragma omp target map(alloc : result[0 : N]) + for (int i = 0; i < N; i++) + result[i] += i; + + // Update strided elements from device: even indices 0,2,4,...,30 +#pragma omp target update from(result[0 : 16 : 2]) + } + + printf("after target update from (even indices up to 30 updated):\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + printf("\n"); + + // Expected: even indices i, odd indices 0 + // CHECK: after target update from + // CHECK: 0.000000 + // CHECK: 0.000000 + // CHECK: 2.000000 + // CHECK: 0.000000 + // CHECK: 4.000000 + // CHECK: 0.000000 + // CHECK: 6.000000 + // CHECK: 0.000000 + // CHECK: 8.000000 + // CHECK: 0.000000 + // CHECK: 10.000000 + // CHECK: 0.000000 + // CHECK: 12.000000 + // CHECK: 0.000000 + // CHECK: 14.000000 + // CHECK: 0.000000 + // CHECK: 16.000000 + // CHECK: 0.000000 + // CHECK: 18.000000 + // CHECK: 0.000000 + // CHECK: 20.000000 + // CHECK: 0.000000 + // CHECK: 22.000000 + // CHECK: 0.000000 + // CHECK: 24.000000 + // CHECK: 0.000000 + // CHECK: 26.000000 + // CHECK: 0.000000 + // CHECK: 28.000000 + // CHECK: 0.000000 + // CHECK: 30.000000 + // CHECK: 0.000000 + // CHECK-NOT: 1.000000 + // CHECK-NOT: 3.000000 + // CHECK-NOT: 31.000000 + + free(result); + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_from.c b/offload/test/offloading/target_update_strided_struct_from.c new file mode 100644 index 0000000000000..46448d10189f2 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_from.c @@ -0,0 +1,86 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update from" with user-defined mapper supports strided +// sections using fixed-size arrays in structs. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +typedef struct { + double data[N]; + size_t len; +} T; + +#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len]) + +int main() { + T s; + s.len = N; + + for (int i = 0; i < N; i++) { + s.data[i] = i; + } + + printf("original host array values:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(mapper(custom), to : s) + { +// Execute on device with explicit mapper +#pragma omp target map(mapper(custom), tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] += i; + } + } + +// Update strided elements from device: indices 0,2,4,6,8,10,12,14 +#pragma omp target update from(s.data[0 : 8 : 2]) + } + + printf("from target array results:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", s.data[i]); + + // CHECK: original host array values: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 3.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 10.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 14.000000 + // CHECK-NEXT: 15.000000 + + // CHECK: from target array results: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 3.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 16.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 20.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 24.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 28.000000 + // CHECK-NEXT: 15.000000 + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_multiple_from.c b/offload/test/offloading/target_update_strided_struct_multiple_from.c new file mode 100644 index 0000000000000..fa57d35b86a4a --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_multiple_from.c @@ -0,0 +1,130 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update from(s1.data[0:6:2], +// s2.data[0:4:3]) correctly updates strided sections covering the full arrays +// from device to host. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define LEN 12 + +typedef struct { + double data[LEN]; + int len; +} T; + +int main() { + T s1, s2; + s1.len = LEN; + s2.len = LEN; + + // Initialize struct arrays on host with simple sequential values + for (int i = 0; i < LEN; i++) { + s1.data[i] = i; + s2.data[i] = i; + } + + printf("original host struct array values:\n"); + printf("s1.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s1.data[i]); + printf("s2.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s2.data[i]); + +#pragma omp target data map(to : s1, s2) + { +// Initialize all device values to 20 +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.data[i] = 20.0; + s2.data[i] = 20.0; + } + } + +// Modify specific strided elements on device to 10 +#pragma omp target map(tofrom : s1, s2) + { + // s1: modify even indices (0,2,4,6,8,10) + for (int i = 0; i < 6; i++) { + s1.data[i * 2] = 10.0; + } + // s2: modify every 3rd index (0,3,6,9) + for (int i = 0; i < 4; i++) { + s2.data[i * 3] = 10.0; + } + } + +// s1.data[0:6:2] updates only even indices: 0,2,4,6,8,10 +// s2.data[0:4:3] updates only every 3rd: 0,3,6,9 +#pragma omp target update from(s1.data[0 : 6 : 2], s2.data[0 : 4 : 3]) + } + + printf("host struct array values after update from:\n"); + printf("s1.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s1.data[i]); + printf("s2.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s2.data[i]); + + // CHECK: original host struct array values: + // CHECK-NEXT: s1.data: + // 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: s2.data: + // 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: host struct array values after update from: + // CHECK-NEXT: s1.data: + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 3.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 9.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 11.0 + // CHECK-NEXT: s2.data: + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 2.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 4.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 8.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 11.0 + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_multiple_to.c b/offload/test/offloading/target_update_strided_struct_multiple_to.c new file mode 100644 index 0000000000000..69c4015153580 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_multiple_to.c @@ -0,0 +1,136 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update to(s1.data[0:6:2], +// s2.data[0:4:3]) correctly updates strided sections covering the full arrays +// from host to device. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define LEN 12 + +typedef struct { + double data[LEN]; + int len; +} T; + +int main() { + T s1, s2; + s1.len = LEN; + s2.len = LEN; + + // Initialize struct arrays on host with simple sequential values + for (int i = 0; i < LEN; i++) { + s1.data[i] = i; + s2.data[i] = i; + } + + printf("original host struct array values:\n"); + printf("s1.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s1.data[i]); + printf("s2.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s2.data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : s1, s2) + { +// Initialize device struct arrays to 20 +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.data[i] = 20.0; + s2.data[i] = 20.0; + } + } + + // s1: even indices (0,2,4,6,8,10) + for (int i = 0; i < 6; i++) { + s1.data[i * 2] = 10.0; + } + // s2: every 3rd index (0,3,6,9) + for (int i = 0; i < 4; i++) { + s2.data[i * 3] = 10.0; + } + +// s1.data[0:6:2] updates all even indices: 0,2,4,6,8,10 (6 elements, stride 2) +// s2.data[0:4:3] updates every 3rd: 0,3,6,9 (4 elements, stride 3) +#pragma omp target update to(s1.data[0 : 6 : 2], s2.data[0 : 4 : 3]) + +// Verify update on device by adding 5 +#pragma omp target map(tofrom : s1, s2) + { + for (int i = 0; i < s1.len; i++) { + s1.data[i] += 5.0; + s2.data[i] += 5.0; + } + } + } + + printf("device struct array values after update to:\n"); + printf("s1.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s1.data[i]); + printf("s2.data:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s2.data[i]); + + // CHECK: original host struct array values: + // CHECK-NEXT: s1.data: + // 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: s2.data: + // 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: device struct array values after update to: + // CHECK-NEXT: s1.data: + // 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: s2.data: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_partial_from.c b/offload/test/offloading/target_update_strided_struct_partial_from.c new file mode 100644 index 0000000000000..c0ffc72f900a6 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_partial_from.c @@ -0,0 +1,86 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update from(s.data[0:2:3]) correctly +// updates every third element (stride 3) from the device to the host +// using a struct with fixed-size array member. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define LEN 11 + +typedef struct { + double data[LEN]; + size_t len; +} T; + +#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len]) + +int main() { + T s; + s.len = LEN; + + // Initialize struct data on host + for (int i = 0; i < LEN; i++) { + s.data[i] = i; + } + + printf("original host array values:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(mapper(custom), to : s) + { +// Execute on device with mapper +#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 2 stride positions) +#pragma omp target map(mapper(custom), tofrom : s) + { + s.data[0] = 10.0; + s.data[3] = 10.0; + } + +// indices 0,3 only +#pragma omp target update from(s.data[0 : 2 : 3]) + } + + printf("device array values after update from:\n"); + for (int i = 0; i < LEN; 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: device array values after update from: + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 2.0 + // CHECK-NEXT: 10.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 + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_partial_to.c b/offload/test/offloading/target_update_strided_struct_partial_to.c new file mode 100644 index 0000000000000..018dcac7cdcf8 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_partial_to.c @@ -0,0 +1,85 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that #pragma omp target update to(s.data[0:2:3]) correctly +// updates every third element (stride 3) from the host to the device +// for struct member arrays. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define LEN 11 + +typedef struct { + double data[LEN]; + int len; +} T; + +int main() { + T s; + s.len = LEN; + + // Initialize struct array on host with simple sequential values + for (int i = 0; i < LEN; i++) + s.data[i] = i; + + printf("original host struct array values:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : s) + { +// Initialize all elements on device to 20 +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) + s.data[i] = 20.0; + } + + // Modify host struct data for elements that will be updated (set to 10) + s.data[0] = 10.0; + s.data[3] = 10.0; + +// indices 0,3 only +#pragma omp target update to(s.data[0 : 2 : 3]) + +// Verify on device by adding 5 to all elements +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) + s.data[i] += 5.0; + } + } + + printf("device struct array values after update to:\n"); + for (int i = 0; i < LEN; i++) + printf("%.1f\n", s.data[i]); + + // CHECK: original host struct array values: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 2.0 + // CHECK-NEXT: 3.0 + // CHECK-NEXT: 4.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 6.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 8.0 + // CHECK-NEXT: 9.0 + // CHECK-NEXT: 10.0 + + // CHECK: device struct array values after update to: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + + return 0; +} diff --git a/offload/test/offloading/target_update_strided_struct_to.c b/offload/test/offloading/target_update_strided_struct_to.c new file mode 100644 index 0000000000000..90b3a5fe26989 --- /dev/null +++ b/offload/test/offloading/target_update_strided_struct_to.c @@ -0,0 +1,98 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update to" with struct member arrays supports strided +// sections using fixed-size arrays in structs. + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +typedef struct { + double data[N]; + int len; +} T; + +int main() { + T s; + s.len = N; + + // Initialize struct array on host with simple sequential values + for (int i = 0; i < N; i++) { + s.data[i] = i; + } + + printf("original host struct array values:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : s) + { +// Initialize device struct array to 20 +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] = 20.0; + } + } + + // Modify host struct data for strided elements (set to 10) + for (int i = 0; i < 8; i++) { + s.data[i * 2] = 10.0; // Set even indices to 10 + } + +// indices 0,2,4,6,8,10,12,14 +#pragma omp target update to(s.data[0 : 8 : 2]) + +// Execute on device - add 5 to verify update worked +#pragma omp target map(tofrom : s) + { + for (int i = 0; i < s.len; i++) { + s.data[i] += 5.0; + } + } + } + + printf("after target update to struct:\n"); + for (int i = 0; i < N; i++) + printf("%.1f\n", s.data[i]); + + // CHECK: original host struct array values: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 2.0 + // CHECK-NEXT: 3.0 + // CHECK-NEXT: 4.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 6.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 8.0 + // CHECK-NEXT: 9.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 11.0 + // CHECK-NEXT: 12.0 + // CHECK-NEXT: 13.0 + // CHECK-NEXT: 14.0 + // CHECK-NEXT: 15.0 + + // CHECK: after target update to struct: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + + return 0; +} diff --git a/offload/test/offloading/target_update_to.c b/offload/test/offloading/target_update_to.c new file mode 100644 index 0000000000000..d45d3182a166d --- /dev/null +++ b/offload/test/offloading/target_update_to.c @@ -0,0 +1,94 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// This test checks that "update to" clause in OpenMP supports strided sections. +// #pragma omp target update to(result[0:8:2]) updates every other element +// (stride 2) + +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#define N 16 + +int main() { + double *result = (double *)calloc(N, sizeof(double)); + + // Initialize on host + for (int i = 0; i < N; i++) { + result[i] = i; + } + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + printf("\n"); + +#pragma omp target data map(from : result[0 : N]) + { +// Initialize device array to 0 +#pragma omp target + { + for (int i = 0; i < N; i++) { + result[i] = 0.0; + } + } + + // Modify host strided elements (even indices) + for (int i = 0; i < 8; i++) { + result[i * 2] = 100.0; + } + +// Update strided elements to device: indices 0,2,4,6,8,10,12,14 +#pragma omp target update to(result[0 : 8 : 2]) + +#pragma omp target + { + for (int i = 0; i < N; i++) { + result[i] += i; + } + } + } + + printf("from target array results:\n"); + for (int i = 0; i < N; i++) + printf("%f\n", result[i]); + + // CHECK: original host array values: + // CHECK-NEXT: 0.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 3.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 10.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 14.000000 + // CHECK-NEXT: 15.000000 + + // CHECK: from target array results: + // CHECK-NEXT: 100.000000 + // CHECK-NEXT: 1.000000 + // CHECK-NEXT: 102.000000 + // CHECK-NEXT: 3.000000 + // CHECK-NEXT: 104.000000 + // CHECK-NEXT: 5.000000 + // CHECK-NEXT: 106.000000 + // CHECK-NEXT: 7.000000 + // CHECK-NEXT: 108.000000 + // CHECK-NEXT: 9.000000 + // CHECK-NEXT: 110.000000 + // CHECK-NEXT: 11.000000 + // CHECK-NEXT: 112.000000 + // CHECK-NEXT: 13.000000 + // CHECK-NEXT: 114.000000 + // CHECK-NEXT: 15.000000 + + free(result); + return 0; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
