[clang] [llvm] [OpenMP] Fix non-contiguous array omp target update (PR #156889)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 4 07:08:48 PDT 2025
github-actions[bot] wrote:
<!--LLVM CODE FORMAT COMMENT: {clang-format}-->
:warning: C/C++ code formatter, clang-format found issues in your code. :warning:
<details>
<summary>
You can test this locally with the following command:
</summary>
``````````bash
git-clang-format --diff origin/main HEAD --extensions cpp -- clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/target_update_codegen.cpp offload/libomptarget/omptarget.cpp offload/test/offloading/non_contiguous_update.cpp
``````````
:warning:
The reproduction instructions above might return results for more than one PR
in a stack if you are using a stacked PR workflow. You can limit the results by
changing `origin/main` to the base branch/commit you want to compare against.
:warning:
</details>
<details>
<summary>
View the diff from clang-format here.
</summary>
``````````diff
diff --git a/offload/test/offloading/non_contiguous_update.cpp b/offload/test/offloading/non_contiguous_update.cpp
index 3973174bf..fa4d4f544 100644
--- a/offload/test/offloading/non_contiguous_update.cpp
+++ b/offload/test/offloading/non_contiguous_update.cpp
@@ -18,110 +18,111 @@ enum tgt_map_type { OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000 };
#ifdef __cplusplus
extern "C" {
#endif
- void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
- void **args_base, void **args, int64_t *arg_sizes,
- int64_t *arg_types);
+void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args, int64_t *arg_sizes,
+ int64_t *arg_types);
#ifdef __cplusplus
}
#endif
int main() {
{
- // case 1
- // int arr[3][4][5][6];
- // #pragma omp target update to(arr[0:2][1:3][1:2][:])
- // set up descriptor
- __tgt_target_non_contig non_contig[5] = {
- {0, 2, 480}, {120, 3, 120}, {24, 2, 24}, {0, 6, 4}, {0, 4, 1}};
- int64_t size = sizeof(non_contig) / sizeof(non_contig[0]), type = OMP_TGT_MAPTYPE_NON_CONTIG;
+ // case 1
+ // int arr[3][4][5][6];
+ // #pragma omp target update to(arr[0:2][1:3][1:2][:])
+ // set up descriptor
+ __tgt_target_non_contig non_contig[5] = {
+ {0, 2, 480}, {120, 3, 120}, {24, 2, 24}, {0, 6, 4}, {0, 4, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
- void *base;
- void *begin = &non_contig;
- int64_t *sizes = &size;
- int64_t *types = &type;
+ void *base;
+ void *begin = &non_contig;
+ int64_t *sizes = &size;
+ int64_t *types = &type;
- // The below diagram is the visualization of the non-contiguous transfer after
- // optimization. Note that each element represent the merged innermost
- // dimension (unit size = 24) since the stride * count of last dimension is
- // equal to the stride of second last dimension.
- //
- // OOOOO OOOOO OOOOO
- // OXXOO OXXOO OOOOO
- // OXXOO OXXOO OOOOO
- // OXXOO OXXOO OOOOO
- __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
- sizes, types);
- // DEBUG: offset 144 len 48
- // DEBUG: offset 264 len 48
- // DEBUG: offset 384 len 48
- // DEBUG: offset 624 len 48
- // DEBUG: offset 744 len 48
- // DEBUG: offset 864 len 48
+ // The below diagram is the visualization of the non-contiguous transfer
+ // after optimization. Note that each element represent the merged innermost
+ // dimension (unit size = 24) since the stride * count of last dimension is
+ // equal to the stride of second last dimension.
+ //
+ // OOOOO OOOOO OOOOO
+ // OXXOO OXXOO OOOOO
+ // OXXOO OXXOO OOOOO
+ // OXXOO OXXOO OOOOO
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ sizes, types);
+ // DEBUG: offset 144 len 48
+ // DEBUG: offset 264 len 48
+ // DEBUG: offset 384 len 48
+ // DEBUG: offset 624 len 48
+ // DEBUG: offset 744 len 48
+ // DEBUG: offset 864 len 48
}
{
- // case 2
- // double darr[3][4][5];
- // #pragma omp target update to(darr[0:2:2][2:2][:2:2])
- // set up descriptor
- __tgt_target_non_contig non_contig[4] = {
- {0, 2, 320}, {80, 2, 40}, {0, 2, 16}, {0, 8, 1}};
- int64_t size = sizeof(non_contig) / sizeof(non_contig[0]), type = OMP_TGT_MAPTYPE_NON_CONTIG;
+ // case 2
+ // double darr[3][4][5];
+ // #pragma omp target update to(darr[0:2:2][2:2][:2:2])
+ // set up descriptor
+ __tgt_target_non_contig non_contig[4] = {
+ {0, 2, 320}, {80, 2, 40}, {0, 2, 16}, {0, 8, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
- void *base;
- void *begin = &non_contig;
- int64_t *sizes = &size;
- int64_t *types = &type;
+ void *base;
+ void *begin = &non_contig;
+ int64_t *sizes = &size;
+ int64_t *types = &type;
- // The below diagram is the visualization of the non-contiguous transfer after
- // optimization. Note that each element represent the innermost dimension
- // (unit size = 8).
- //
- // OOOOO OOOOO OOOOO
- // OOOOO OOOOO OOOOO
- // XOXOO OOOOO XOXOO
- // XOXOO OOOOO XOXOO
- __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
- sizes, types);
- // DEBUG: offset 80 len 8
- // DEBUG: offset 96 len 8
- // DEBUG: offset 120 len 8
- // DEBUG: offset 136 len 8
- // DEBUG: offset 400 len 8
- // DEBUG: offset 416 len 8
- // DEBUG: offset 440 len 8
- // DEBUG: offset 456 len 8
+ // The below diagram is the visualization of the non-contiguous transfer
+ // after optimization. Note that each element represent the innermost
+ // dimension (unit size = 8).
+ //
+ // OOOOO OOOOO OOOOO
+ // OOOOO OOOOO OOOOO
+ // XOXOO OOOOO XOXOO
+ // XOXOO OOOOO XOXOO
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ sizes, types);
+ // DEBUG: offset 80 len 8
+ // DEBUG: offset 96 len 8
+ // DEBUG: offset 120 len 8
+ // DEBUG: offset 136 len 8
+ // DEBUG: offset 400 len 8
+ // DEBUG: offset 416 len 8
+ // DEBUG: offset 440 len 8
+ // DEBUG: offset 456 len 8
}
{
- // case 3
- // int darr[6][6];
- // #pragma omp target update to(darr[1:2:2][2:3])
- // set up descriptor
- __tgt_target_non_contig non_contig[3] = {
- {24, 2, 48}, {8, 3, 4}, {0, 4, 1}};
- int64_t size = sizeof(non_contig) / sizeof(non_contig[0]), type = OMP_TGT_MAPTYPE_NON_CONTIG;
+ // case 3
+ // int darr[6][6];
+ // #pragma omp target update to(darr[1:2:2][2:3])
+ // set up descriptor
+ __tgt_target_non_contig non_contig[3] = {{24, 2, 48}, {8, 3, 4}, {0, 4, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
- void *base;
- void *begin = &non_contig;
- int64_t *sizes = &size;
- int64_t *types = &type;
-
- // The below diagram is the visualization of the non-contiguous transfer after
- // optimization. Note that each element represent the merged innermost
- // dimension (unit size = 12).
- //
- // OOOOOO
- // OOXXXO
- // OOOOOO
- // OOXXXO
- // OOOOOO
- // OOOOOO
- __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
- sizes, types);
- // DEBUG: offset 24 len 12
- // DEBUG: offset 72 len 12
+ void *base;
+ void *begin = &non_contig;
+ int64_t *sizes = &size;
+ int64_t *types = &type;
+ // The below diagram is the visualization of the non-contiguous transfer
+ // after optimization. Note that each element represent the merged innermost
+ // dimension (unit size = 12).
+ //
+ // OOOOOO
+ // OOXXXO
+ // OOOOOO
+ // OOXXXO
+ // OOOOOO
+ // OOOOOO
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ sizes, types);
+ // DEBUG: offset 24 len 12
+ // DEBUG: offset 72 len 12
}
return 0;
``````````
</details>
https://github.com/llvm/llvm-project/pull/156889
More information about the llvm-commits
mailing list