https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/169889
None >From f394b935494d2f9614ff821f6e5fc0fb7fdd2361 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 8 Sep 2025 08:23:21 -0400 Subject: [PATCH 1/4] handle_array_pointer_var_check --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index a8255ac74cfcf..d9e3f5ed61e83 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8239,10 +8239,12 @@ 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"); + else if (&Component == &*Components.begin()) { + // Handle pointer-based array sections like data[a:b:c] + if (const auto *PtrType = Ty->getAs<PointerType>()) { + ElementType = PtrType->getPointeeType().getTypePtr(); + } + } // 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. >From 4762b1159d370b8722d7733fb9cdfcab950e5658 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Fri, 21 Nov 2025 06:47:59 -0500 Subject: [PATCH 2/4] plugin_interface_issue_fix --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 30 +++++++++++++++++++++------ 1 file changed, 24 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index d9e3f5ed61e83..a0b373478b19b 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8240,14 +8240,25 @@ class MappableExprsHandler { else if (VAT) ElementType = VAT->getElementType().getTypePtr(); else if (&Component == &*Components.begin()) { - // Handle pointer-based array sections like data[a:b:c] + // 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"); } - // 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. + + // 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. @@ -8957,8 +8968,15 @@ 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()) { - UnionCurInfo.NonContigInfo.Dims.push_back(0); - // Emit a combined entry: + // 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]). + UnionCurInfo.NonContigInfo.Dims.insert( + UnionCurInfo.NonContigInfo.Dims.begin(), 1); emitCombinedEntry(CombinedInfo, UnionCurInfo.Types, PartialStruct, /*IsMapThis*/ !VD, OMPBuilder, VD); } >From 3c552fd3adad3faeba591182f5805247fb645244 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 24 Nov 2025 05:55:21 -0500 Subject: [PATCH 3/4] testcases_added --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 2 +- clang/test/OpenMP/target_update_codegen.cpp | 9 +- .../target_update_strided_ptr_messages_from.c | 44 ++++++ .../target_update_strided_ptr_messages_to.c | 44 ++++++ ...pdate_strided_ptr_multiple_messages_from.c | 64 +++++++++ ..._update_strided_ptr_multiple_messages_to.c | 64 +++++++++ ...update_strided_ptr_partial_messages_from.c | 41 ++++++ ...t_update_strided_ptr_partial_messages_to.c | 40 ++++++ .../strided_ptr_multiple_update_from.c | 114 +++++++++++++++ .../strided_ptr_multiple_update_to.c | 130 +++++++++++++++++ .../strided_ptr_partial_update_from.c | 68 +++++++++ .../strided_ptr_partial_update_to.c | 81 +++++++++++ offload/test/offloading/target_update_from.c | 73 ++++++++++ .../target_update_strided_struct_from.c | 86 +++++++++++ ...rget_update_strided_struct_multiple_from.c | 130 +++++++++++++++++ ...target_update_strided_struct_multiple_to.c | 136 ++++++++++++++++++ ...arget_update_strided_struct_partial_from.c | 86 +++++++++++ .../target_update_strided_struct_partial_to.c | 85 +++++++++++ .../target_update_strided_struct_to.c | 98 +++++++++++++ offload/test/offloading/target_update_to.c | 81 +++++++++++ 20 files changed, 1469 insertions(+), 7 deletions(-) create mode 100644 clang/test/OpenMP/target_update_strided_ptr_messages_from.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_messages_to.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c create mode 100644 clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c create mode 100644 offload/test/offloading/strided_ptr_multiple_update_from.c create mode 100644 offload/test/offloading/strided_ptr_multiple_update_to.c create mode 100644 offload/test/offloading/strided_ptr_partial_update_from.c create mode 100644 offload/test/offloading/strided_ptr_partial_update_to.c create mode 100644 offload/test/offloading/target_update_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_multiple_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_multiple_to.c create mode 100644 offload/test/offloading/target_update_strided_struct_partial_from.c create mode 100644 offload/test/offloading/target_update_strided_struct_partial_to.c create mode 100644 offload/test/offloading/target_update_strided_struct_to.c create mode 100644 offload/test/offloading/target_update_to.c diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index a0b373478b19b..ae17d749dc979 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8258,7 +8258,7 @@ class MappableExprsHandler { } // At this stage, if ElementType was a base pointer and we are in the - // first iteration, it has been computed. + // first iteration, it has been computed. if (ElementType) { // For the case that having pointer as base, we need to remove one // level of indirection. diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp index c8211f475c7fc..c7ff467ad056a 100644 --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -1177,7 +1177,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 { @@ -1221,11 +1221,8 @@ struct ST { // CK21: store i64 1, ptr [[COUNT_4]], // CK21: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2 // CK21: store i64 {{4|8}}, ptr [[STRIDE_4]], - // 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: call void @__tgt_target_data_update_mapper(ptr @1, i64 -1, i32 2, ptr %{{[0-9]+}}, ptr %{{[0-9]+}}, ptr %{{[0-9]+}}, ptr @.offload_maptypes, ptr null, ptr null) + // 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..cd158ff10e4de --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c @@ -0,0 +1,44 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +extern void *malloc(__SIZE_TYPE__); +extern void free(void *); + +int main(int argc, char **argv) { + int len = 16; + double *data = (double *)malloc(len * sizeof(double)); + + // 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'}} + {} + + free(data); + 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..64839a8384425 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c @@ -0,0 +1,44 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +extern void *malloc(__SIZE_TYPE__); +extern void free(void *); + +int main(int argc, char **argv) { + int len = 16; + double *data = (double *)malloc(len * sizeof(double)); + + // 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'}} + {} + + free(data); + 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..2462ae1c840c3 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c @@ -0,0 +1,64 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +extern void *malloc(__SIZE_TYPE__); +extern void free(void *); + +double *data; +double *data1; +double *data2; + +int main(int argc, char **argv) { + int len = 12; + + // Allocate memory for explicit pointers + data = (double *)malloc(len * sizeof(double)); + data1 = (double *)malloc(len * sizeof(double)); + data2 = (double *)malloc(len * sizeof(double)); + + // Valid multiple strided array sections + #pragma omp target update from(data1[0:6:2], data2[0:4:3]) // OK - different 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 '['}} + + free(data); + free(data1); + free(data2); + 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..d8cebefb02606 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c @@ -0,0 +1,64 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +extern void *malloc(__SIZE_TYPE__); +extern void free(void *); + +double *data; +double *data1; +double *data2; + +int main(int argc, char **argv) { + int len = 12; + + // Allocate memory for explicit pointers + data = (double *)malloc(len * sizeof(double)); + data1 = (double *)malloc(len * sizeof(double)); + data2 = (double *)malloc(len * sizeof(double)); + + // Valid multiple strided array sections + #pragma omp target update to(data1[0:6:2], data2[0:4:3]) // OK - different 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 '['}} + + free(data); + free(data1); + free(data2); + 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..2f36dfa84cd1b --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c @@ -0,0 +1,41 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +extern void *malloc(__SIZE_TYPE__); +extern void free(void *); + +int main(int argc, char **argv) { + int len = 11; + double *data = (double *)malloc(len * sizeof(double)); + + // 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'}} + + free(data); + 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..ce4b315bec1ae --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_partial_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 + +extern void *malloc(__SIZE_TYPE__); +extern void free(void *); + +int main(int argc, char **argv) { + int len = 11; + double *data = (double *)malloc(len * sizeof(double)); + + // 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'}} + + free(data); + return 0; +} \ No newline at end of file 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..afb2f1fc8dd53 --- /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; +} \ No newline at end of file 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..2005117c83754 --- /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; +} \ No newline at end of file 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..e2f4398b2fbd7 --- /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; +} \ No newline at end of file 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..28bd27fd04366 --- /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; +} \ No newline at end of file diff --git a/offload/test/offloading/target_update_from.c b/offload/test/offloading/target_update_from.c new file mode 100644 index 0000000000000..089b498ada0db --- /dev/null +++ b/offload/test/offloading/target_update_from.c @@ -0,0 +1,73 @@ +// 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: 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; +} \ No newline at end of file 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..1c5759b1864c5 --- /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), tofrom : 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; +} \ No newline at end of file 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..95faf104826d5 --- /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(tofrom : 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: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: s2.data: + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 20.0 + + return 0; +} \ No newline at end of file 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..1c3646aefabec --- /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; +} \ No newline at end of file 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..d40fc85adc4f7 --- /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), tofrom : 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; +} \ No newline at end of file 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..d0c8adc6b5ea0 --- /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; +} \ No newline at end of file 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..b3fe32bbaa345 --- /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; +} \ No newline at end of file diff --git a/offload/test/offloading/target_update_to.c b/offload/test/offloading/target_update_to.c new file mode 100644 index 0000000000000..ecfc3bb69dbf4 --- /dev/null +++ b/offload/test/offloading/target_update_to.c @@ -0,0 +1,81 @@ +// 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(tofrom : result[0 : N]) + { +// Update strided elements to device: indices 0,2,4,6 +#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: 0.000000 + // CHECK-NEXT: 2.000000 + // CHECK-NEXT: 4.000000 + // CHECK-NEXT: 6.000000 + // CHECK-NEXT: 8.000000 + // CHECK-NEXT: 10.000000 + // CHECK-NEXT: 12.000000 + // CHECK-NEXT: 14.000000 + // CHECK-NEXT: 16.000000 + // CHECK-NEXT: 18.000000 + // CHECK-NEXT: 20.000000 + // CHECK-NEXT: 22.000000 + // CHECK-NEXT: 24.000000 + // CHECK-NEXT: 26.000000 + // CHECK-NEXT: 28.000000 + // CHECK-NEXT: 30.000000 + + free(result); + return 0; +} \ No newline at end of file >From 24a9ce34985c72962aa20c9ae68a12d80c01fdf5 Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Fri, 28 Nov 2025 04:47:34 -0500 Subject: [PATCH 4/4] test1 --- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 24 +++++++++++++++-------- 1 file changed, 16 insertions(+), 8 deletions(-) diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 5101717526263..a4250f2f68ad7 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9048,16 +9048,24 @@ 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) { + 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; } } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
