Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[libomptarget] Add support for target update non-contiguous
This patch is the runtime support for https://reviews.llvm.org/D84192. In order not to modify the tgt_target_data_update information but still be able to pass the extra information for non-contiguous map item (offset, count, and stride for each dimension), this patch overload arg when the maptype is set as OMP_TGT_MAPTYPE_DESCRIPTOR. The origin arg is for passing the pointer information, however, the overloaded arg is an array of descriptor_dim: ``` struct descriptor_dim { int64_t offset; int64_t count; int64_t stride }; ``` and the array size is the dimension size. In addition, since we have count and stride information in descriptor_dim, we can replace/overload the arg_size parameter by using dimension size. Reviewed By: grokos Differential Revision: https://reviews.llvm.org/D82245
- Loading branch information
1 parent
e6f3ec6
commit 6847bce
Showing
3 changed files
with
253 additions
and
81 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
100 changes: 100 additions & 0 deletions
100
openmp/libomptarget/test/offloading/non_contiguous_update.cpp
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,100 @@ | ||
| // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG | ||
| // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG | ||
| // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG | ||
| // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG | ||
| // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG | ||
| // REQUIRES: libomptarget-debug | ||
|
|
||
| #include <cstdio> | ||
| #include <cstdlib> | ||
| #include <cassert> | ||
|
|
||
| // Data structure definitions copied from OpenMP RTL. | ||
| struct __tgt_target_non_contig { | ||
| int64_t offset; | ||
| int64_t width; | ||
| int64_t stride; | ||
| }; | ||
|
|
||
| enum tgt_map_type { | ||
| OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000 | ||
| }; | ||
|
|
||
| // OpenMP RTL interfaces | ||
| #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); | ||
| #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}, {1, 3, 120}, {1, 2, 24}, {0, 6, 4}, {0, 1, 4}}; | ||
| int64_t size = 4, 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 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 | ||
| // DEBUG: offset 264 | ||
| // DEBUG: offset 384 | ||
| // DEBUG: offset 624 | ||
| // DEBUG: offset 744 | ||
| // DEBUG: offset 864 | ||
|
|
||
|
|
||
| // 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_2[4] = { | ||
| {0, 2, 320}, {2, 2, 40}, {0, 2, 16}, {0, 1, 8}}; | ||
| int64_t size_2 = 4, type_2 = OMP_TGT_MAPTYPE_NON_CONTIG; | ||
|
|
||
| void *base_2; | ||
| void *begin_2 = &non_contig_2; | ||
| int64_t *sizes_2 = &size_2; | ||
| int64_t *types_2 = &type_2; | ||
|
|
||
| // The below diagram is the visualization of the non-contiguous transfer after | ||
| // optimization. Note that each element represent the innermost dimension | ||
| // (unit size = 24) since the stride * count of last dimension is equal to the | ||
| // stride of second last dimension. | ||
| // | ||
| // OOOOO OOOOO OOOOO | ||
| // OOOOO OOOOO OOOOO | ||
| // XOXOO OOOOO XOXOO | ||
| // XOXOO OOOOO XOXOO | ||
| __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base_2, &begin_2, | ||
| sizes_2, types_2); | ||
| // DEBUG: offset 80 | ||
| // DEBUG: offset 96 | ||
| // DEBUG: offset 120 | ||
| // DEBUG: offset 136 | ||
| // DEBUG: offset 400 | ||
| // DEBUG: offset 416 | ||
| // DEBUG: offset 440 | ||
| // DEBUG: offset 456 | ||
| return 0; | ||
| } |