https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/170287
>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 b651261614a57e22c4cc240e90fbdd34c9d54adf Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Mon, 24 Nov 2025 05:55:21 -0500 Subject: [PATCH 3/4] testing_extern --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 2 +- clang/test/OpenMP/target_update_codegen.cpp | 9 +- clang/test/OpenMP/target_update_codegen.ll | 26 ++++ .../target_update_strided_ptr_messages_from.c | 44 ++++++ .../target_update_strided_ptr_messages_to.c | 41 ++++++ ...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 +++++++++++ 21 files changed, 1492 insertions(+), 7 deletions(-) create mode 100644 clang/test/OpenMP/target_update_codegen.ll 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_codegen.ll b/clang/test/OpenMP/target_update_codegen.ll new file mode 100644 index 0000000000000..ae34f51835c9d --- /dev/null +++ b/clang/test/OpenMP/target_update_codegen.ll @@ -0,0 +1,26 @@ +; ModuleID = '/tmp/target_update_codegen-91f92b.bc' +source_filename = "target_update_codegen.cpp" +target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128" +target triple = "x86_64-unknown-linux-gnu" + [email protected] = private constant [19160 x i8] c"\10\FF\10\AD\01\00\00\00\D8J\00\00\00\00\00\00 \00\00\00\00\00\00\00(\00\00\00\00\00\00\00\02\00\01\00\00\00\00\00H\00\00\00\00\00\00\00\02\00\00\00\00\00\00\00\90\00\00\00\00\00\00\00BJ\00\00\00\00\00\00i\00\00\00\00\00\00\00\87\00\00\00\00\00\00\00n\00\00\00\00\00\00\00u\00\00\00\00\00\00\00\00arch\00triple\00amdgcn-amd-amdhsa\00gfx90a\00\00\00; ModuleID = 'target_update_codegen.cpp'\0Asource_filename = \22target_update_codegen.cpp\22\0Atarget datalayout = \22e-m:e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9\22\0Atarget triple = \22amdgcn-amd-amdhsa\22\0A\0A@__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0\0A@__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0\0A@__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0\0A@__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0\0A@__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0\0A\0A!llvm.dbg.cu = !{!0}\0A!llvm.module.flags = !{!213, !214, !215, !216, !217, !218, !219, !220}\0A!llvm.ident = !{!221, !222}\0A!opencl.ocl.version = !{!223}\0A\0A!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus_14, file: !1, producer: \22clang version 22.0.0git (https://github.com/llvm/llvm-project.git e38e0feaf7bb13e6147eca002cd616499253c595)\22, isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, imports: !2, splitDebugInlining: false, nameTableKind: None)\0A!1 = !DIFile(filename: \22target_update_codegen.cpp\22, directory: \22/work/amtiwari/llvm-project/clang/test/OpenMP\22, checksumkind: CSK_MD5, checksum: \227308366b6b2010802bc22bb2e9c0e7c4\22)\0A!2 = !{!3, !11, !15, !22, !26, !34, !39, !41, !49, !53, !57, !67, !69, !73, !77, !81, !86, !90, !94, !98, !102, !110, !114, !118, !120, !124, !128, !133, !139, !143, !147, !149, !157, !161, !169, !171, !175, !179, !183, !187, !192, !197, !202, !203, !204, !205, !207, !208, !209, !210, !211, !212}\0A!3 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !5, file: !10, line: 52)\0A!4 = !DINamespace(name: \22std\22, scope: null)\0A!5 = !DISubprogram(name: \22abs\22, scope: !6, file: !6, line: 848, type: !7, flags: DIFlagPrototyped, spFlags: 0)\0A!6 = !DIFile(filename: \22/usr/include/stdlib.h\22, directory: \22\22, checksumkind: CSK_MD5, checksum: \2202258fad21adf111bb9df9825e61954a\22)\0A!7 = !DISubroutineType(types: !8)\0A!8 = !{!9, !9}\0A!9 = !DIBasicType(name: \22int\22, size: 32, encoding: DW_ATE_signed)\0A!10 = !DIFile(filename: \22/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/bits/std_abs.h\22, directory: \22\22)\0A!11 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !12, file: !14, line: 127)\0A!12 = !DIDerivedType(tag: DW_TAG_typedef, name: \22div_t\22, file: !6, line: 63, baseType: !13)\0A!13 = !DICompositeType(tag: DW_TAG_structure_type, file: !6, line: 59, size: 64, flags: DIFlagFwdDecl, identifier: \22_ZTS5div_t\22)\0A!14 = !DIFile(filename: \22/usr/lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/cstdlib\22, directory: \22\22)\0A!15 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !16, file: !14, line: 128)\0A!16 = !DIDerivedType(tag: DW_TAG_typedef, name: \22ldiv_t\22, file: !6, line: 71, baseType: !17)\0A!17 = distinct !DICompositeType(tag: DW_TAG_structure_type, file: !6, line: 67, size: 128, flags: DIFlagTypePassByValue, elements: !18, identifier: \22_ZTS6ldiv_t\22)\0A!18 = !{!19, !21}\0A!19 = !DIDerivedType(tag: DW_TAG_member, name: \22quot\22, scope: !17, file: !6, line: 69, baseType: !20, size: 64)\0A!20 = !DIBasicType(name: \22long\22, size: 64, encoding: DW_ATE_signed)\0A!21 = !DIDerivedType(tag: DW_TAG_member, name: \22rem\22, scope: !17, file: !6, line: 70, baseType: !20, size: 64, offset: 64)\0A!22 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !23, file: !14, line: 130)\0A!23 = !DISubprogram(name: \22abort\22, scope: !6, file: !6, line: 598, type: !24, flags: DIFlagPrototyped | DIFlagNoReturn, spFlags: 0)\0A!24 = !DISubroutineType(types: !25)\0A!25 = !{null}\0A!26 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !27, file: !14, line: 132)\0A!27 = !DISubprogram(name: \22aligned_alloc\22, scope: !6, file: !6, line: 592, type: !28, flags: DIFlagPrototyped, spFlags: 0)\0A!28 = !DISubroutineType(types: !29)\0A!29 = !{!30, !31, !31}\0A!30 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: null, size: 64, dwarfAddressSpace: 1)\0A!31 = !DIDerivedType(tag: DW_TAG_typedef, name: \22size_t\22, file: !32, line: 18, baseType: !33)\0A!32 = !DIFile(filename: \22lib/clang/22/include/__stddef_size_t.h\22, directory: \22/work/amtiwari/llvm-project\22, checksumkind: CSK_MD5, checksum: \222c44e821a2b1951cde2eb0fb2e656867\22)\0A!33 = !DIBasicType(name: \22unsigned long\22, size: 64, encoding: DW_ATE_unsigned)\0A!34 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !35, file: !14, line: 134)\0A!35 = !DISubprogram(name: \22atexit\22, scope: !6, file: !6, line: 602, type: !36, flags: DIFlagPrototyped, spFlags: 0)\0A!36 = !DISubroutineType(types: !37)\0A!37 = !{!9, !38}\0A!38 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !24, size: 64, dwarfAddressSpace: 1)\0A!39 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !40, file: !14, line: 137)\0A!40 = !DISubprogram(name: \22at_quick_exit\22, scope: !6, file: !6, line: 607, type: !36, flags: DIFlagPrototyped, spFlags: 0)\0A!41 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !42, file: !14, line: 140)\0A!42 = !DISubprogram(name: \22atof\22, scope: !6, file: !6, line: 102, type: !43, flags: DIFlagPrototyped, spFlags: 0)\0A!43 = !DISubroutineType(types: !44)\0A!44 = !{!45, !46}\0A!45 = !DIBasicType(name: \22double\22, size: 64, encoding: DW_ATE_float)\0A!46 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !47, size: 64, dwarfAddressSpace: 1)\0A!47 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !48)\0A!48 = !DIBasicType(name: \22char\22, size: 8, encoding: DW_ATE_signed_char)\0A!49 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !50, file: !14, line: 141)\0A!50 = !DISubprogram(name: \22atoi\22, scope: !6, file: !6, line: 105, type: !51, flags: DIFlagPrototyped, spFlags: 0)\0A!51 = !DISubroutineType(types: !52)\0A!52 = !{!9, !46}\0A!53 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !54, file: !14, line: 142)\0A!54 = !DISubprogram(name: \22atol\22, scope: !6, file: !6, line: 108, type: !55, flags: DIFlagPrototyped, spFlags: 0)\0A!55 = !DISubroutineType(types: !56)\0A!56 = !{!20, !46}\0A!57 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !58, file: !14, line: 143)\0A!58 = !DISubprogram(name: \22bsearch\22, scope: !6, file: !6, line: 828, type: !59, flags: DIFlagPrototyped, spFlags: 0)\0A!59 = !DISubroutineType(types: !60)\0A!60 = !{!30, !61, !61, !31, !31, !63}\0A!61 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !62, size: 64, dwarfAddressSpace: 1)\0A!62 = !DIDerivedType(tag: DW_TAG_const_type, baseType: null)\0A!63 = !DIDerivedType(tag: DW_TAG_typedef, name: \22__compar_fn_t\22, file: !6, line: 816, baseType: !64)\0A!64 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !65, size: 64, dwarfAddressSpace: 1)\0A!65 = !DISubroutineType(types: !66)\0A!66 = !{!9, !61, !61}\0A!67 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !68, file: !14, line: 144)\0A!68 = !DISubprogram(name: \22calloc\22, scope: !6, file: !6, line: 543, type: !28, flags: DIFlagPrototyped, spFlags: 0)\0A!69 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !70, file: !14, line: 145)\0A!70 = !DISubprogram(name: \22div\22, scope: !6, file: !6, line: 860, type: !71, flags: DIFlagPrototyped, spFlags: 0)\0A!71 = !DISubroutineType(types: !72)\0A!72 = !{!12, !9, !9}\0A!73 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !74, file: !14, line: 146)\0A!74 = !DISubprogram(name: \22exit\22, scope: !6, file: !6, line: 624, type: !75, flags: DIFlagPrototyped | DIFlagNoReturn, spFlags: 0)\0A!75 = !DISubroutineType(types: !76)\0A!76 = !{null, !9}\0A!77 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !78, file: !14, line: 147)\0A!78 = !DISubprogram(name: \22free\22, scope: !6, file: !6, line: 555, type: !79, flags: DIFlagPrototyped, spFlags: 0)\0A!79 = !DISubroutineType(types: !80)\0A!80 = !{null, !30}\0A!81 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !82, file: !14, line: 148)\0A!82 = !DISubprogram(name: \22getenv\22, scope: !6, file: !6, line: 641, type: !83, flags: DIFlagPrototyped, spFlags: 0)\0A!83 = !DISubroutineType(types: !84)\0A!84 = !{!85, !46}\0A!85 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !48, size: 64, dwarfAddressSpace: 1)\0A!86 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !87, file: !14, line: 149)\0A!87 = !DISubprogram(name: \22labs\22, scope: !6, file: !6, line: 849, type: !88, flags: DIFlagPrototyped, spFlags: 0)\0A!88 = !DISubroutineType(types: !89)\0A!89 = !{!20, !20}\0A!90 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !91, file: !14, line: 150)\0A!91 = !DISubprogram(name: \22ldiv\22, scope: !6, file: !6, line: 862, type: !92, flags: DIFlagPrototyped, spFlags: 0)\0A!92 = !DISubroutineType(types: !93)\0A!93 = !{!16, !20, !20}\0A!94 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !95, file: !14, line: 151)\0A!95 = !DISubprogram(name: \22malloc\22, scope: !6, file: !6, line: 540, type: !96, flags: DIFlagPrototyped, spFlags: 0)\0A!96 = !DISubroutineType(types: !97)\0A!97 = !{!30, !31}\0A!98 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !99, file: !14, line: 153)\0A!99 = !DISubprogram(name: \22mblen\22, scope: !6, file: !6, line: 930, type: !100, flags: DIFlagPrototyped, spFlags: 0)\0A!100 = !DISubroutineType(types: !101)\0A!101 = !{!9, !46, !31}\0A!102 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !103, file: !14, line: 154)\0A!103 = !DISubprogram(name: \22mbstowcs\22, scope: !6, file: !6, line: 941, type: !104, flags: DIFlagPrototyped, spFlags: 0)\0A!104 = !DISubroutineType(types: !105)\0A!105 = !{!31, !106, !109, !31}\0A!106 = !DIDerivedType(tag: DW_TAG_restrict_type, baseType: !107)\0A!107 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !108, size: 64, dwarfAddressSpace: 1)\0A!108 = !DIBasicType(name: \22wchar_t\22, size: 32, encoding: DW_ATE_signed)\0A!109 = !DIDerivedType(tag: DW_TAG_restrict_type, baseType: !46)\0A!110 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !111, file: !14, line: 155)\0A!111 = !DISubprogram(name: \22mbtowc\22, scope: !6, file: !6, line: 933, type: !112, flags: DIFlagPrototyped, spFlags: 0)\0A!112 = !DISubroutineType(types: !113)\0A!113 = !{!9, !106, !109, !31}\0A!114 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !115, file: !14, line: 157)\0A!115 = !DISubprogram(name: \22qsort\22, scope: !6, file: !6, line: 838, type: !116, flags: DIFlagPrototyped, spFlags: 0)\0A!116 = !DISubroutineType(types: !117)\0A!117 = !{null, !30, !31, !31, !63}\0A!118 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !119, file: !14, line: 160)\0A!119 = !DISubprogram(name: \22quick_exit\22, scope: !6, file: !6, line: 630, type: !75, flags: DIFlagPrototyped | DIFlagNoReturn, spFlags: 0)\0A!120 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !121, file: !14, line: 163)\0A!121 = !DISubprogram(name: \22rand\22, scope: !6, file: !6, line: 454, type: !122, flags: DIFlagPrototyped, spFlags: 0)\0A!122 = !DISubroutineType(types: !123)\0A!123 = !{!9}\0A!124 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !125, file: !14, line: 164)\0A!125 = !DISubprogram(name: \22realloc\22, scope: !6, file: !6, line: 551, type: !126, flags: DIFlagPrototyped, spFlags: 0)\0A!126 = !DISubroutineType(types: !127)\0A!127 = !{!30, !30, !31}\0A!128 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !129, file: !14, line: 165)\0A!129 = !DISubprogram(name: \22srand\22, scope: !6, file: !6, line: 456, type: !130, flags: DIFlagPrototyped, spFlags: 0)\0A!130 = !DISubroutineType(types: !131)\0A!131 = !{null, !132}\0A!132 = !DIBasicType(name: \22unsigned int\22, size: 32, encoding: DW_ATE_unsigned)\0A!133 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !134, file: !14, line: 166)\0A!134 = !DISubprogram(name: \22strtod\22, scope: !6, file: !6, line: 118, type: !135, flags: DIFlagPrototyped, spFlags: 0)\0A!135 = !DISubroutineType(types: !136)\0A!136 = !{!45, !109, !137}\0A!137 = !DIDerivedType(tag: DW_TAG_restrict_type, baseType: !138)\0A!138 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !85, size: 64, dwarfAddressSpace: 1)\0A!139 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !140, file: !14, line: 167)\0A!140 = !DISubprogram(name: \22strtol\22, scope: !6, file: !6, line: 177, type: !141, flags: DIFlagPrototyped, spFlags: 0)\0A!141 = !DISubroutineType(types: !142)\0A!142 = !{!20, !109, !137, !9}\0A!143 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !144, file: !14, line: 168)\0A!144 = !DISubprogram(name: \22strtoul\22, scope: !6, file: !6, line: 181, type: !145, flags: DIFlagPrototyped, spFlags: 0)\0A!145 = !DISubroutineType(types: !146)\0A!146 = !{!33, !109, !137, !9}\0A!147 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !148, file: !14, line: 169)\0A!148 = !DISubprogram(name: \22system\22, scope: !6, file: !6, line: 791, type: !51, flags: DIFlagPrototyped, spFlags: 0)\0A!149 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !150, file: !14, line: 171)\0A!150 = !DISubprogram(name: \22wcstombs\22, scope: !6, file: !6, line: 945, type: !151, flags: DIFlagPrototyped, spFlags: 0)\0A!151 = !DISubroutineType(types: !152)\0A!152 = !{!31, !153, !154, !31}\0A!153 = !DIDerivedType(tag: DW_TAG_restrict_type, baseType: !85)\0A!154 = !DIDerivedType(tag: DW_TAG_restrict_type, baseType: !155)\0A!155 = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !156, size: 64, dwarfAddressSpace: 1)\0A!156 = !DIDerivedType(tag: DW_TAG_const_type, baseType: !108)\0A!157 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !158, file: !14, line: 172)\0A!158 = !DISubprogram(name: \22wctomb\22, scope: !6, file: !6, line: 937, type: !159, flags: DIFlagPrototyped, spFlags: 0)\0A!159 = !DISubroutineType(types: !160)\0A!160 = !{!9, !85, !108}\0A!161 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !163, file: !14, line: 200)\0A!162 = !DINamespace(name: \22__gnu_cxx\22, scope: null)\0A!163 = !DIDerivedType(tag: DW_TAG_typedef, name: \22lldiv_t\22, file: !6, line: 81, baseType: !164)\0A!164 = distinct !DICompositeType(tag: DW_TAG_structure_type, file: !6, line: 77, size: 128, flags: DIFlagTypePassByValue, elements: !165, identifier: \22_ZTS7lldiv_t\22)\0A!165 = !{!166, !168}\0A!166 = !DIDerivedType(tag: DW_TAG_member, name: \22quot\22, scope: !164, file: !6, line: 79, baseType: !167, size: 64)\0A!167 = !DIBasicType(name: \22long long\22, size: 64, encoding: DW_ATE_signed)\0A!168 = !DIDerivedType(tag: DW_TAG_member, name: \22rem\22, scope: !164, file: !6, line: 80, baseType: !167, size: 64, offset: 64)\0A!169 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !170, file: !14, line: 206)\0A!170 = !DISubprogram(name: \22_Exit\22, scope: !6, file: !6, line: 636, type: !75, flags: DIFlagPrototyped | DIFlagNoReturn, spFlags: 0)\0A!171 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !172, file: !14, line: 210)\0A!172 = !DISubprogram(name: \22llabs\22, scope: !6, file: !6, line: 852, type: !173, flags: DIFlagPrototyped, spFlags: 0)\0A!173 = !DISubroutineType(types: !174)\0A!174 = !{!167, !167}\0A!175 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !176, file: !14, line: 216)\0A!176 = !DISubprogram(name: \22lldiv\22, scope: !6, file: !6, line: 866, type: !177, flags: DIFlagPrototyped, spFlags: 0)\0A!177 = !DISubroutineType(types: !178)\0A!178 = !{!163, !167, !167}\0A!179 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !180, file: !14, line: 227)\0A!180 = !DISubprogram(name: \22atoll\22, scope: !6, file: !6, line: 113, type: !181, flags: DIFlagPrototyped, spFlags: 0)\0A!181 = !DISubroutineType(types: !182)\0A!182 = !{!167, !46}\0A!183 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !184, file: !14, line: 228)\0A!184 = !DISubprogram(name: \22strtoll\22, scope: !6, file: !6, line: 201, type: !185, flags: DIFlagPrototyped, spFlags: 0)\0A!185 = !DISubroutineType(types: !186)\0A!186 = !{!167, !109, !137, !9}\0A!187 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !188, file: !14, line: 229)\0A!188 = !DISubprogram(name: \22strtoull\22, scope: !6, file: !6, line: 206, type: !189, flags: DIFlagPrototyped, spFlags: 0)\0A!189 = !DISubroutineType(types: !190)\0A!190 = !{!191, !109, !137, !9}\0A!191 = !DIBasicType(name: \22unsigned long long\22, size: 64, encoding: DW_ATE_unsigned)\0A!192 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !193, file: !14, line: 231)\0A!193 = !DISubprogram(name: \22strtof\22, scope: !6, file: !6, line: 124, type: !194, flags: DIFlagPrototyped, spFlags: 0)\0A!194 = !DISubroutineType(types: !195)\0A!195 = !{!196, !109, !137}\0A!196 = !DIBasicType(name: \22float\22, size: 32, encoding: DW_ATE_float)\0A!197 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !162, entity: !198, file: !14, line: 232)\0A!198 = !DISubprogram(name: \22strtold\22, scope: !6, file: !6, line: 127, type: !199, flags: DIFlagPrototyped, spFlags: 0)\0A!199 = !DISubroutineType(types: !200)\0A!200 = !{!201, !109, !137}\0A!201 = !DIBasicType(name: \22long double\22, size: 128, encoding: DW_ATE_float)\0A!202 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !163, file: !14, line: 240)\0A!203 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !170, file: !14, line: 242)\0A!204 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !172, file: !14, line: 244)\0A!205 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !206, file: !14, line: 245)\0A!206 = !DISubprogram(name: \22div\22, linkageName: \22_ZN9__gnu_cxx3divExx\22, scope: !162, file: !14, line: 213, type: !177, flags: DIFlagPrototyped, spFlags: 0)\0A!207 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !176, file: !14, line: 246)\0A!208 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !180, file: !14, line: 248)\0A!209 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !193, file: !14, line: 249)\0A!210 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !184, file: !14, line: 250)\0A!211 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !188, file: !14, line: 251)\0A!212 = !DIImportedEntity(tag: DW_TAG_imported_declaration, scope: !4, entity: !198, file: !14, line: 252)\0A!213 = !{i32 1, !\22amdhsa_code_object_version\22, i32 600}\0A!214 = !{i32 7, !\22Dwarf Version\22, i32 5}\0A!215 = !{i32 2, !\22Debug Info Version\22, i32 3}\0A!216 = !{i32 1, !\22wchar_size\22, i32 4}\0A!217 = !{i32 7, !\22openmp\22, i32 51}\0A!218 = !{i32 7, !\22openmp-device\22, i32 51}\0A!219 = !{i32 8, !\22PIC Level\22, i32 2}\0A!220 = !{i32 7, !\22frame-pointer\22, i32 2}\0A!221 = !{!\22clang version 22.0.0git (https://github.com/llvm/llvm-project.git e38e0feaf7bb13e6147eca002cd616499253c595)\22}\0A!222 = !{!\22AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.1 24491 1e0fda770a2079fbd71e4b70974d74f62fd3af10)\22}\0A!223 = !{i32 2, i32 0}\0A\00\00\00\00\00\00", section ".llvm.offloading", align 8, !exclude !0 [email protected] = appending global [1 x ptr] [ptr @llvm.embedded.object], section "llvm.metadata" + +!llvm.dbg.cu = !{!1} +!llvm.module.flags = !{!3, !4, !5, !6, !7, !8, !9, !10} +!llvm.ident = !{!11} +!llvm.embedded.objects = !{!12} + +!0 = !{} +!1 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus_14, file: !2, producer: "clang version 22.0.0git (https://github.com/llvm/llvm-project.git e38e0feaf7bb13e6147eca002cd616499253c595)", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, splitDebugInlining: false, nameTableKind: None) +!2 = !DIFile(filename: "target_update_codegen.cpp", directory: "/work/amtiwari/llvm-project/clang/test/OpenMP", checksumkind: CSK_MD5, checksum: "7308366b6b2010802bc22bb2e9c0e7c4") +!3 = !{i32 7, !"Dwarf Version", i32 5} +!4 = !{i32 2, !"Debug Info Version", i32 3} +!5 = !{i32 1, !"wchar_size", i32 4} +!6 = !{i32 7, !"openmp", i32 51} +!7 = !{i32 8, !"PIC Level", i32 2} +!8 = !{i32 7, !"PIE Level", i32 2} +!9 = !{i32 7, !"uwtable", i32 2} +!10 = !{i32 7, !"frame-pointer", i32 2} +!11 = !{!"clang version 22.0.0git (https://github.com/llvm/llvm-project.git e38e0feaf7bb13e6147eca002cd616499253c595)"} +!12 = !{ptr @llvm.embedded.object, !".llvm.offloading"} 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..6459c79149326 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_messages_to.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 + +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 dd6dd25b079ddf5189540c255fc7ec25f5bb047a Mon Sep 17 00:00:00 2001 From: amtiwari <[email protected]> Date: Fri, 5 Dec 2025 04:25:05 -0500 Subject: [PATCH 4/4] test_w/o_extern --- .../test/OpenMP/target_update_strided_ptr_messages_from.c | 4 ---- clang/test/OpenMP/target_update_strided_ptr_messages_to.c | 1 - .../target_update_strided_ptr_multiple_messages_from.c | 8 -------- .../target_update_strided_ptr_multiple_messages_to.c | 8 -------- .../target_update_strided_ptr_partial_messages_from.c | 4 ---- .../target_update_strided_ptr_partial_messages_to.c | 4 ---- 6 files changed, 29 deletions(-) diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c index cd158ff10e4de..65757083dd50d 100644 --- a/clang/test/OpenMP/target_update_strided_ptr_messages_from.c +++ b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c @@ -1,12 +1,8 @@ // 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 diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c index 6459c79149326..e074fa6cbbfc7 100644 --- a/clang/test/OpenMP/target_update_strided_ptr_messages_to.c +++ b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c @@ -3,7 +3,6 @@ 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 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 index 2462ae1c840c3..3aa8777676a66 100644 --- a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c +++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c @@ -1,9 +1,6 @@ // 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; @@ -11,11 +8,6 @@ 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 {} 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 index d8cebefb02606..ce26ad2d827e0 100644 --- a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c +++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c @@ -1,9 +1,6 @@ // 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; @@ -11,11 +8,6 @@ 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 {} 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 index 2f36dfa84cd1b..3d8f26cdecc0b 100644 --- a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c +++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c @@ -1,12 +1,8 @@ // 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 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 index ce4b315bec1ae..8df0b46dfb272 100644 --- a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c +++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c @@ -1,12 +1,8 @@ // 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
