diff options
Diffstat (limited to 'offload/test')
46 files changed, 1835 insertions, 57 deletions
diff --git a/offload/test/CMakeLists.txt b/offload/test/CMakeLists.txt index 711621d..c317394 100644 --- a/offload/test/CMakeLists.txt +++ b/offload/test/CMakeLists.txt @@ -61,7 +61,7 @@ add_offload_testsuite(check-offload "Running libomptarget tests" ${LIBOMPTARGET_LIT_TESTSUITES} EXCLUDE_FROM_CHECK_ALL - DEPENDS llvm-offload-device-info omptarget ${OMP_DEPEND} ${LIBOMPTARGET_TESTED_PLUGINS} + DEPENDS llvm-offload-device-info omptarget ${OMP_DEPEND} ${LIBOMPTARGET_TESTED_PLUGINS} check-offload-unit ARGS ${LIBOMPTARGET_LIT_ARG_LIST}) # Add liboffload unit tests - the test binary will run on all available devices diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg index 800a63b..f3e8e9a 100644 --- a/offload/test/lit.cfg +++ b/offload/test/lit.cfg @@ -121,6 +121,7 @@ if config.libomptarget_test_pgo: # For all other targets, we currently assume it is. supports_unified_shared_memory = True supports_apu = False +supports_large_allocation_memory_pool = False if config.libomptarget_current_target.startswith('nvptx'): try: cuda_arch = int(config.cuda_test_arch[:3]) @@ -132,9 +133,11 @@ if config.libomptarget_current_target.startswith('nvptx'): elif config.libomptarget_current_target.startswith('amdgcn'): # amdgpu_test_arch contains a list of AMD GPUs in the system # only check the first one assuming that we will run the test on it. - if not (config.amdgpu_test_arch.startswith("gfx90a") or - config.amdgpu_test_arch.startswith("gfx942") or - config.amdgpu_test_arch.startswith("gfx950")): + if (config.amdgpu_test_arch.startswith("gfx90a") or + config.amdgpu_test_arch.startswith("gfx942") or + config.amdgpu_test_arch.startswith("gfx950")): + supports_large_allocation_memory_pool = True + else: supports_unified_shared_memory = False # check if AMD architecture is an APU: if ((config.amdgpu_test_arch.startswith("gfx942") and @@ -144,6 +147,8 @@ if supports_unified_shared_memory: config.available_features.add('unified_shared_memory') if supports_apu: config.available_features.add('apu') +if supports_large_allocation_memory_pool: + config.available_features.add('large_allocation_memory_pool') # Setup environment to find dynamic library at runtime if config.operating_system == 'Windows': diff --git a/offload/test/mapping/data_member_ref.cpp b/offload/test/mapping/data_member_ref.cpp index fdb8abc..7947a62 100644 --- a/offload/test/mapping/data_member_ref.cpp +++ b/offload/test/mapping/data_member_ref.cpp @@ -60,7 +60,8 @@ int main() { printf("Host %d %d.\n", Bar.VRef.Data, V.Data); // CHECK: Host 123456. printf("Host %d.\n", *Baz.VRef.Data); -#pragma omp target map(*Baz.VRef.Data) map(from : D1, D2) +#pragma omp target map(Baz.VRef.Data) map(*Baz.VRef.Data) map(V1.Data[0 : 0]) \ + map(from : D1, D2) { // CHECK: Device 123456. D1 = *Baz.VRef.Data; diff --git a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp index c6c5657..45fd042 100644 --- a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp +++ b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp @@ -44,8 +44,8 @@ int main() { int spp00fa = -1, spp00fca = -1, spp00fb_r = -1; __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]); -#pragma omp target map(tofrom: spp[0][0]) firstprivate(p) \ - map(from: spp00fa, spp00fca, spp00fb_r) +#pragma omp target map(tofrom : spp[0][0]) map(alloc : spp[0]) firstprivate(p) \ + map(from : spp00fa, spp00fca, spp00fb_r) { spp00fa = spp[0][0].f.a; spp00fca = spp[0][0].f.c.a; diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp index a9e3f05..a59ed69 100644 --- a/offload/test/mapping/declare_mapper_nested_mappers.cpp +++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp @@ -42,8 +42,8 @@ int main() { int spp00fa = -1, spp00fb_r = -1, spp00fg1 = -1, spp00fg_r = -1; __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]), p1 = reinterpret_cast<__intptr_t>(&y[0]); -#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1) \ - map(from: spp00fa, spp00fb_r, spp00fg1, spp00fg_r) +#pragma omp target map(tofrom : spp[0][0]) map(alloc : spp[0]) \ + firstprivate(p, p1) map(from : spp00fa, spp00fb_r, spp00fg1, spp00fg_r) { spp00fa = spp[0][0].f.a; spp00fb_r = spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0; diff --git a/offload/test/mapping/map_ptr_and_star_global.c b/offload/test/mapping/map_ptr_and_star_global.c index c3b0dd2..869fb8c 100644 --- a/offload/test/mapping/map_ptr_and_star_global.c +++ b/offload/test/mapping/map_ptr_and_star_global.c @@ -1,5 +1,7 @@ // RUN: %libomptarget-compilexx-run-and-check-generic +// REQUIRES: libc + #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_ptr_and_star_local.c b/offload/test/mapping/map_ptr_and_star_local.c index f0ca84d..cc826b3 100644 --- a/offload/test/mapping/map_ptr_and_star_local.c +++ b/offload/test/mapping/map_ptr_and_star_local.c @@ -1,5 +1,7 @@ // RUN: %libomptarget-compilexx-run-and-check-generic +// REQUIRES: libc + #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_ptr_and_subscript_global.c b/offload/test/mapping/map_ptr_and_subscript_global.c index a3a10b6..839db06 100644 --- a/offload/test/mapping/map_ptr_and_subscript_global.c +++ b/offload/test/mapping/map_ptr_and_subscript_global.c @@ -1,5 +1,7 @@ // RUN: %libomptarget-compilexx-run-and-check-generic +// REQUIRES: libc + #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_ptr_and_subscript_local.c b/offload/test/mapping/map_ptr_and_subscript_local.c index bb44999..68ac9dc 100644 --- a/offload/test/mapping/map_ptr_and_subscript_local.c +++ b/offload/test/mapping/map_ptr_and_subscript_local.c @@ -1,5 +1,7 @@ // RUN: %libomptarget-compilexx-run-and-check-generic +// REQUIRES: libc + #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_structptr_and_member_global.c b/offload/test/mapping/map_structptr_and_member_global.c index 10e72e0..960eea4 100644 --- a/offload/test/mapping/map_structptr_and_member_global.c +++ b/offload/test/mapping/map_structptr_and_member_global.c @@ -1,5 +1,7 @@ // RUN: %libomptarget-compilexx-run-and-check-generic +// REQUIRES: libc + #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_structptr_and_member_local.c b/offload/test/mapping/map_structptr_and_member_local.c index 9e59551..bd75940 100644 --- a/offload/test/mapping/map_structptr_and_member_local.c +++ b/offload/test/mapping/map_structptr_and_member_local.c @@ -1,5 +1,7 @@ // RUN: %libomptarget-compilexx-run-and-check-generic +// REQUIRES: libc + #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/ptr_and_obj_motion.c b/offload/test/mapping/ptr_and_obj_motion.c index 8fa2c98..a94c07aa 100644 --- a/offload/test/mapping/ptr_and_obj_motion.c +++ b/offload/test/mapping/ptr_and_obj_motion.c @@ -17,7 +17,7 @@ void init(double vertexx[]) { } void change(DV *dvptr) { -#pragma omp target map(dvptr->dataptr[0 : 100]) +#pragma omp target map(dvptr->dataptr[0 : 100]) map(alloc : dvptr -> dataptr) { printf("In change: %lf, expected 77.0\n", dvptr->dataptr[77]); dvptr->dataptr[77] += 1.0; diff --git a/offload/test/mapping/target_derefence_array_pointrs.cpp b/offload/test/mapping/target_derefence_array_pointrs.cpp index a6dd4069..d213c87 100644 --- a/offload/test/mapping/target_derefence_array_pointrs.cpp +++ b/offload/test/mapping/target_derefence_array_pointrs.cpp @@ -18,23 +18,24 @@ void foo(int **t1d) { for (j = 0; j < 3; j++) (*t1d)[j] = 0; -#pragma omp target map(tofrom : (*t1d)[0 : 3]) +#pragma omp target map(tofrom : (*t1d)[0 : 3]) map(alloc : *t1d) { (*t1d)[1] = 1; } // CHECK: 1 printf("%d\n", (*t1d)[1]); -#pragma omp target map(tofrom : (**t2d)[0 : 3]) +#pragma omp target map(tofrom : (**t2d)[0 : 3]) map(alloc : **t2d, *t2d) { (**t2d)[1] = 2; } // CHECK: 2 printf("%d\n", (**t2d)[1]); -#pragma omp target map(tofrom : (***t3d)[0 : 3]) +#pragma omp target map(tofrom : (***t3d)[0 : 3]) \ + map(alloc : ***t3d, **t3d, *t3d) { (***t3d)[1] = 3; } // CHECK: 3 printf("%d\n", (***t3d)[1]); -#pragma omp target map(tofrom : (**t1d)) +#pragma omp target map(tofrom : (**t1d)) map(alloc : *t1d) { (*t1d)[0] = 4; } // CHECK: 4 printf("%d\n", (*t1d)[0]); -#pragma omp target map(tofrom : (*(*(t1d + a) + b))) +#pragma omp target map(tofrom : (*(*(t1d + a) + b))) map(to : *(t1d + a)) { *(*(t1d + a) + b) = 5; } // CHECK: 5 printf("%d\n", *(*(t1d + a) + b)); @@ -49,7 +50,7 @@ void bar() { for (int i = 0; i < 3; i++) { (**a)[1] = i; } -#pragma omp target map((**a)[ : 3]) +#pragma omp target map((**a)[ : 3]) map(alloc : **a, *a) { (**a)[1] = 6; // CHECK: 6 @@ -73,7 +74,8 @@ void zoo(int **f, SSA *sa) { *(f + sa->i + 1) = t; *(sa->sa->i + *(f + sa->i + 1)) = 4; printf("%d\n", *(sa->sa->i + *(1 + sa->i + f))); -#pragma omp target map(sa, *(sa->sa->i + *(1 + sa->i + f))) +#pragma omp target map(*(sa->sa->i + *(1 + sa->i + f))) map(alloc : sa->sa) \ + map(to : sa->i) map(to : sa->sa->i) map(to : *(1 + sa->i + f)) { *(sa->sa->i + *(1 + sa->i + f)) = 7; } // CHECK: 7 printf("%d\n", *(sa->sa->i + *(1 + sa->i + f))); @@ -87,13 +89,13 @@ void xoo() { void yoo(int **x) { *x = (int *)malloc(2 * sizeof(int)); -#pragma omp target map(**x) +#pragma omp target map(**x) map(alloc : *x) { **x = 8; // CHECK: 8 printf("%d\n", **x); } -#pragma omp target map(*(*x + 1)) +#pragma omp target map(*(*x + 1)) map(alloc : *x) { *(*x + 1) = 9; // CHECK: 9 diff --git a/offload/test/mapping/target_has_device_addr.c b/offload/test/mapping/target_has_device_addr.c index e8bfff8..f238832 100644 --- a/offload/test/mapping/target_has_device_addr.c +++ b/offload/test/mapping/target_has_device_addr.c @@ -66,8 +66,9 @@ void zoo() { short **xpp = &xp[0]; x[1] = 111; -#pragma omp target data map(tofrom : xpp[1][1]) use_device_addr(xpp[1][1]) -#pragma omp target has_device_addr(xpp[1][1]) +#pragma omp target data map(tofrom : xpp[1][1]) map(xpp[1]) \ + use_device_addr(xpp[1]) +#pragma omp target has_device_addr(xpp[1]) { xpp[1][1] = 222; // CHECK: 222 diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp new file mode 100644 index 0000000..3b1a819 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp @@ -0,0 +1,85 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on an array-section. +// The corresponding data is mapped on a previous enter_data directive. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g, h[10]; +int *ph = &h[0]; + +struct S { + int *paa[10][10]; + + void f1(int i) { + paa[0][2] = &g; + + int *original_ph3 = &ph[3]; + int **original_paa02 = &paa[0][2]; + +#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5]) + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device()); + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device()); + + // CHECK-COUNT-4: 1 + printf("%d\n", mapped_ptr_ph3 != nullptr); + printf("%d\n", mapped_ptr_paa02 != nullptr); + printf("%d\n", original_ph3 != mapped_ptr_ph3); + printf("%d\n", original_paa02 != mapped_ptr_paa02); + +// (A) use_device_addr operand within mapped address range. +// CHECK: A: 1 +#pragma omp target data use_device_addr(ph[3 : 4]) + printf("A: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (B) use_device_addr operand in extended address range, but not +// mapped address range. +// CHECK: B: 1 +#pragma omp target data use_device_addr(ph[2]) + printf("B: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (C) use_device_addr/map: same base-array, different first-location. +// CHECK: C: 1 +#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1]) + printf("C: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (D) use_device_addr/map: different base-array/pointers. +// CHECK: D: 1 +#pragma omp target data map(ph) use_device_addr(ph[3 : 4]) + printf("D: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (E) use_device_addr operand within mapped range of previous map. +// CHECK: E: 1 +#pragma omp target data use_device_addr(paa[0]) + printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]); + +// (F) use_device_addr/map: different operands, same base-array. +// CHECK: F: 1 +#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2]) + printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]); + +// (G) use_device_addr/map: different base-array/pointers. +// CHECK: G: 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2]) + printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]); + +#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5]) + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp new file mode 100644 index 0000000..b9ebde4 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp @@ -0,0 +1,143 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on an array-section. +// The corresponding data is not previously mapped. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g, h[10]; +int *ph = &h[0]; + +struct S { + int *paa[10][10]; + + void f1(int i) { + paa[0][2] = &g; + + int *original_ph3 = &ph[3]; + int **original_paa02 = &paa[0][2]; + +// (A) No corresponding map, lookup should fail. +// CHECK: A: 1 1 1 +#pragma omp target data use_device_addr(ph[3 : 4]) + { + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device()); + printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3); + } + +// (B) use_device_addr/map: different operands, same base-pointer. +// use_device_addr operand within mapped address range. +// CHECK: B: 1 1 1 +#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1]) + { + int *mapped_ptr_ph4 = + (int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device()); + printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr, + mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4); + } + +// (C) use_device_addr/map: different base-pointers. +// No corresponding storage, lookup should fail. +// CHECK: C: 1 1 1 +#pragma omp target data map(ph) use_device_addr(ph[3 : 4]) + { + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device()); + printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3); + } + +// (D) use_device_addr/map: one of two maps with matching base-pointer. +// use_device_addr operand within mapped address range of second map, +// lookup should succeed. +// CHECK: D: 1 1 1 +#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4]) + { + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device()); + printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, + mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3); + } + +// (E) No corresponding map, lookup should fail +// CHECK: E: 1 1 1 +#pragma omp target data use_device_addr(paa[0]) + { + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device()); + printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr, + mapped_ptr_paa02 != original_paa02, + &paa[0][2] == (int **)nullptr + 2); + } + +// (F) use_device_addr/map: different operands, same base-array. +// use_device_addr within mapped address range. Lookup should succeed. +// CHECK: F: 1 1 1 +#pragma omp target data map(paa) use_device_addr(paa[0]) + { + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device()); + printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr, + mapped_ptr_paa02 != original_paa02, + &paa[0][2] == mapped_ptr_paa02); + } + +// (G) use_device_addr/map: different operands, same base-array. +// use_device_addr extends beyond existing mapping. Not spec compliant. +// But the lookup succeeds because we use the base-address for translation. +// CHECK: G: 1 1 1 +#pragma omp target data map(paa[0][4]) use_device_addr(paa[0]) + { + int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr( + original_paa02 + 2, omp_get_default_device()); + printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr, + mapped_ptr_paa04 != original_paa02 + 2, + &paa[0][4] == mapped_ptr_paa04); + } + + int *original_paa020 = &paa[0][2][0]; + int **original_paa0 = (int **)&paa[0]; + +// (H) use_device_addr/map: different base-pointers. +// No corresponding storage for use_device_addr opnd, lookup should fail. +// CHECK: H: 1 1 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0]) + { + int **mapped_ptr_paa020 = + (int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device()); + int **mapped_ptr_paa0 = + (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device()); + printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr, + mapped_ptr_paa0 == nullptr, &paa[0] == nullptr); + } + +// (I) use_device_addr/map: one map with different, one with same base-ptr. +// Lookup should succeed. +// CHECK: I: 1 1 1 +#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2]) + { + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device()); + printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr, + mapped_ptr_paa02 != original_paa02, + &paa[0][2] == mapped_ptr_paa02); + } + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp new file mode 100644 index 0000000..e9a1124 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp @@ -0,0 +1,98 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on an array-section on a reference. +// The corresponding data is mapped on a previous enter_data directive. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g_ptee; +int &g = g_ptee; + +int h_ptee[10]; +int (&h)[10] = h_ptee; + +int *ph_ptee = &h_ptee[0]; +int *&ph = ph_ptee; +int *paa_ptee[10][10]; + +struct S { + int *(&paa)[10][10] = paa_ptee; + + void f1(int i) { + paa[0][2] = &g; + + int *original_ph3 = &ph[3]; + int **original_paa02 = &paa[0][2]; + +#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5]) + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device()); + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device()); + + // CHECK-COUNT-4: 1 + printf("%d\n", mapped_ptr_ph3 != nullptr); + printf("%d\n", mapped_ptr_paa02 != nullptr); + printf("%d\n", original_ph3 != mapped_ptr_ph3); + printf("%d\n", original_paa02 != mapped_ptr_paa02); + +// (A) use_device_addr operand within mapped address range. +// EXPECTED: A: 1 +// CHECK: A: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data use_device_addr(ph[3 : 4]) + printf("A: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (B) use_device_addr operand in extended address range, but not +// mapped address range. +// EXPECTED: B: 1 +// CHECK: B: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data use_device_addr(ph[2]) + printf("B: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (C) use_device_addr/map: same base-array, different first-location. +// EXPECTED: C: 1 +// CHECK: C: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1]) + printf("C: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (D) use_device_addr/map: different base-array/pointers. +// EXPECTED: D: 1 +// CHECK: D: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) use_device_addr(ph[3 : 4]) + printf("D: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (E) use_device_addr operand within mapped range of previous map. +// CHECK: E: 1 +#pragma omp target data use_device_addr(paa[0]) + printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]); + +// (F) use_device_addr/map: different operands, same base-array. +// CHECK: F: 1 +#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2]) + printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]); + +// (G) use_device_addr/map: different base-array/pointers. +// CHECK: G: 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2]) + printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]); + +#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5]) + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp new file mode 100644 index 0000000..0090cdb --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp @@ -0,0 +1,158 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on an array-section on a reference. +// The corresponding data is not previously mapped. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g_ptee; +int &g = g_ptee; + +int h_ptee[10]; +int (&h)[10] = h_ptee; + +int *ph_ptee = &h_ptee[0]; +int *&ph = ph_ptee; +int *paa_ptee[10][10]; + +struct S { + int *(&paa)[10][10] = paa_ptee; + + void f1(int i) { + paa[0][2] = &g; + + int *original_ph3 = &ph[3]; + int **original_paa02 = &paa[0][2]; + +// (A) No corresponding map, lookup should fail. +// EXPECTED: A: 1 1 1 +// CHECK: A: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data use_device_addr(ph[3 : 4]) + { + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device()); + printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3); + } + +// (B) use_device_addr/map: different operands, same base-pointer. +// use_device_addr operand within mapped address range. +// EXPECTED: B: 1 1 1 +// CHECK: B: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1]) + { + int *mapped_ptr_ph4 = + (int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device()); + printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr, + mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4); + } + +// (C) use_device_addr/map: different base-pointers. +// No corresponding storage, lookup should fail. +// EXPECTED: C: 1 1 1 +// CHECK: C: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) use_device_addr(ph[3 : 4]) + { + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device()); + printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3); + } + +// (D) use_device_addr/map: one of two maps with matching base-pointer. +// use_device_addr operand within mapped address range of second map, +// lookup should succeed. +// EXPECTED: D: 1 1 1 +// CHECK: D: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4]) + { + int *mapped_ptr_ph3 = + (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device()); + printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, + mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3); + } + +// (E) No corresponding map, lookup should fail +// CHECK: E: 1 1 1 +#pragma omp target data use_device_addr(paa[0]) + { + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device()); + printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr, + mapped_ptr_paa02 != original_paa02, + &paa[0][2] == (int **)nullptr + 2); + } + +// (F) use_device_addr/map: different operands, same base-array. +// use_device_addr within mapped address range. Lookup should succeed. +// CHECK: F: 1 1 1 +#pragma omp target data map(paa) use_device_addr(paa[0]) + { + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device()); + printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr, + mapped_ptr_paa02 != original_paa02, + &paa[0][2] == mapped_ptr_paa02); + } + +// (G) use_device_addr/map: different operands, same base-array. +// use_device_addr extends beyond existing mapping. Not spec compliant. +// But the lookup succeeds because we use the base-address for translation. +// CHECK: G: 1 1 1 +#pragma omp target data map(paa[0][4]) use_device_addr(paa[0]) + { + int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr( + original_paa02 + 2, omp_get_default_device()); + printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr, + mapped_ptr_paa04 != original_paa02 + 2, + &paa[0][4] == mapped_ptr_paa04); + } + + int *original_paa020 = &paa[0][2][0]; + int **original_paa0 = (int **)&paa[0]; + +// (H) use_device_addr/map: different base-pointers. +// No corresponding storage for use_device_addr opnd, lookup should fail. +// CHECK: H: 1 1 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0]) + { + int **mapped_ptr_paa020 = + (int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device()); + int **mapped_ptr_paa0 = + (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device()); + printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr, + mapped_ptr_paa0 == nullptr, &paa[0] == nullptr); + } + +// (I) use_device_addr/map: one map with different, one with same base-ptr. +// Lookup should succeed. +// CHECK: I: 1 1 1 +#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2]) + { + int **mapped_ptr_paa02 = + (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device()); + printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr, + mapped_ptr_paa02 != original_paa02, + &paa[0][2] == mapped_ptr_paa02); + } + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp new file mode 100644 index 0000000..883297f --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_existing.cpp @@ -0,0 +1,93 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on a variable (not a section). +// The corresponding data is mapped on a previous enter_data directive. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g, h[10]; +int *ph = &h[0]; + +struct S { + int *paa[10][10]; + + void f1(int i) { + paa[0][2] = &g; + + void *original_addr_g = &g; + void *original_addr_h = &h; + void *original_addr_ph = &ph; + void *original_addr_paa = &paa; + +#pragma omp target enter data map(to : g, h, ph, paa) + void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device()); + void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device()); + void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device()); + void *mapped_ptr_paa = omp_get_mapped_ptr(&paa, omp_get_default_device()); + + // CHECK-COUNT-8: 1 + printf("%d\n", mapped_ptr_g != nullptr); + printf("%d\n", mapped_ptr_h != nullptr); + printf("%d\n", mapped_ptr_ph != nullptr); + printf("%d\n", mapped_ptr_paa != nullptr); + printf("%d\n", original_addr_g != mapped_ptr_g); + printf("%d\n", original_addr_h != mapped_ptr_h); + printf("%d\n", original_addr_ph != mapped_ptr_ph); + printf("%d\n", original_addr_paa != mapped_ptr_paa); + +// (A) +// CHECK: A: 1 +#pragma omp target data use_device_addr(g) + printf("A: %d\n", mapped_ptr_g == &g); + +// (B) +// CHECK: B: 1 +#pragma omp target data use_device_addr(h) + printf("B: %d\n", mapped_ptr_h == &h); + +// (C) +// CHECK: C: 1 +#pragma omp target data use_device_addr(ph) + printf("C: %d\n", mapped_ptr_ph == &ph); + +// (D) use_device_addr/map with different base-array/pointer. +// Address translation should happen for &ph, not &ph[0/1]. +// CHECK: D: 1 +#pragma omp target data map(ph[1 : 2]) use_device_addr(ph) + printf("D: %d\n", mapped_ptr_ph == &ph); + +// (E) +// CHECK: E: 1 +#pragma omp target data use_device_addr(paa) + printf("E: %d\n", mapped_ptr_paa == &paa); + +// (F) use_device_addr/map with same base-array, paa. +// Address translation should happen for &paa. +// CHECK: F: 1 +#pragma omp target data map(paa[0][2]) use_device_addr(paa) + printf("F: %d\n", mapped_ptr_paa == &paa); + +// (G) use_device_addr/map with different base-array/pointer. +// Address translation should happen for &paa. +// CHECK: G: 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa) + printf("G: %d\n", mapped_ptr_paa == &paa); + +#pragma omp target exit data map(release : g, h, ph, paa) + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp new file mode 100644 index 0000000..79c6f69 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp @@ -0,0 +1,159 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on a variable (not a section). +// The corresponding data is not previously mapped. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g, h[10]; +int *ph = &h[0]; + +struct S { + int *paa[10][10]; + + void f1(int i) { + paa[0][2] = &g; + + void *original_addr_g = &g; + void *original_addr_h = &h; + void *original_addr_ph = &ph; + void *original_addr_paa = &paa; + +// (A) No corresponding item, lookup should fail. +// CHECK: A: 1 1 1 +#pragma omp target data use_device_addr(g) + { + void *mapped_ptr_g = + omp_get_mapped_ptr(original_addr_g, omp_get_default_device()); + printf("A: %d %d %d\n", mapped_ptr_g == nullptr, + mapped_ptr_g != original_addr_g, (void *)&g == nullptr); + } + +// (B) Lookup should succeed. +// CHECK: B: 1 1 1 +#pragma omp target data map(g) use_device_addr(g) + { + void *mapped_ptr_g = + omp_get_mapped_ptr(original_addr_g, omp_get_default_device()); + printf("B: %d %d %d\n", mapped_ptr_g != nullptr, + mapped_ptr_g != original_addr_g, &g == mapped_ptr_g); + } + +// (C) No corresponding item, lookup should fail. +// CHECK: C: 1 1 1 +#pragma omp target data use_device_addr(h) + { + void *mapped_ptr_h = + omp_get_mapped_ptr(original_addr_h, omp_get_default_device()); + printf("C: %d %d %d\n", mapped_ptr_h == nullptr, + mapped_ptr_h != original_addr_h, (void *)&h == nullptr); + } + +// (D) Lookup should succeed. +// CHECK: D: 1 1 1 +#pragma omp target data map(h) use_device_addr(h) + { + void *mapped_ptr_h = + omp_get_mapped_ptr(original_addr_h, omp_get_default_device()); + printf("D: %d %d %d\n", mapped_ptr_h != nullptr, + mapped_ptr_h != original_addr_h, &h == mapped_ptr_h); + } + +// (E) No corresponding item, lookup should fail. +// CHECK: E: 1 1 1 +#pragma omp target data use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("E: %d %d %d\n", mapped_ptr_ph == nullptr, + mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr); + } + +// (F) Lookup should succeed. +// CHECK: F: 1 1 1 +#pragma omp target data map(ph) use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("F: %d %d %d\n", mapped_ptr_ph != nullptr, + mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph); + } + +// (G) Maps pointee only, but use_device_addr operand is pointer. +// Lookup should fail. +// CHECK: G: 1 1 1 +#pragma omp target data map(ph[0 : 1]) use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("G: %d %d %d\n", mapped_ptr_ph == nullptr, + mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr); + } + +// (H) Maps both pointee and pointer. Lookup for pointer should succeed. +// CHECK: H: 1 1 1 +#pragma omp target data map(ph[0 : 1]) map(ph) use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("H: %d %d %d\n", mapped_ptr_ph != nullptr, + mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph); + } + +// (I) No corresponding item, lookup should fail. +// CHECK: I: 1 1 1 +#pragma omp target data use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("I: %d %d %d\n", mapped_ptr_paa == nullptr, + mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr); + } + +// (J) Maps pointee only, but use_device_addr operand is pointer. +// Lookup should fail. +// CHECK: J: 1 1 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("J: %d %d %d\n", mapped_ptr_paa == nullptr, + mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr); + } + +// (K) Lookup should succeed. +// CHECK: K: 1 1 1 +#pragma omp target data map(paa) use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("K: %d %d %d\n", mapped_ptr_paa != nullptr, + mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa); + } + +// (L) Maps both pointee and pointer. Lookup for pointer should succeed. +// CHECK: L: 1 1 1 +#pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("L: %d %d %d\n", mapped_ptr_paa != nullptr, + mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa); + } + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp new file mode 100644 index 0000000..f018c65 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_existing.cpp @@ -0,0 +1,100 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on a reference variable. +// The corresponding data is mapped on a previous enter_data directive. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g_ptee; +int &g = g_ptee; + +int h_ptee[10]; +int (&h)[10] = h_ptee; + +int *ph_ptee = &h_ptee[0]; +int *&ph = ph_ptee; +int *paa_ptee[10][10]; + +struct S { + int *(&paa)[10][10] = paa_ptee; + + void f1(int i) { + paa[0][2] = &g; + + void *original_addr_g = &g; + void *original_addr_h = &h; + void *original_addr_ph = &ph; + void *original_addr_paa = &paa; + +#pragma omp target enter data map(to : g, h, ph, paa) + void *mapped_ptr_g = omp_get_mapped_ptr(&g, omp_get_default_device()); + void *mapped_ptr_h = omp_get_mapped_ptr(&h, omp_get_default_device()); + void *mapped_ptr_ph = omp_get_mapped_ptr(&ph, omp_get_default_device()); + void *mapped_ptr_paa = omp_get_mapped_ptr(&paa, omp_get_default_device()); + + // CHECK-COUNT-8: 1 + printf("%d\n", mapped_ptr_g != nullptr); + printf("%d\n", mapped_ptr_h != nullptr); + printf("%d\n", mapped_ptr_ph != nullptr); + printf("%d\n", mapped_ptr_paa != nullptr); + printf("%d\n", original_addr_g != mapped_ptr_g); + printf("%d\n", original_addr_h != mapped_ptr_h); + printf("%d\n", original_addr_ph != mapped_ptr_ph); + printf("%d\n", original_addr_paa != mapped_ptr_paa); + +// (A) +// CHECK: A: 1 +#pragma omp target data use_device_addr(g) + printf("A: %d\n", mapped_ptr_g == &g); + +// (B) +// CHECK: B: 1 +#pragma omp target data use_device_addr(h) + printf("B: %d\n", mapped_ptr_h == &h); + +// (C) +// CHECK: C: 1 +#pragma omp target data use_device_addr(ph) + printf("C: %d\n", mapped_ptr_ph == &ph); + +// (D) use_device_addr/map with different base-array/pointer. +// Address translation should happen for &ph, not &ph[0/1]. +// CHECK: D: 1 +#pragma omp target data map(ph[1 : 2]) use_device_addr(ph) + printf("D: %d\n", mapped_ptr_ph == &ph); + +// (E) +// CHECK: E: 1 +#pragma omp target data use_device_addr(paa) + printf("E: %d\n", mapped_ptr_paa == &paa); + +// (F) use_device_addr/map with same base-array, paa. +// Address translation should happen for &paa. +// CHECK: F: 1 +#pragma omp target data map(paa[0][2]) use_device_addr(paa) + printf("F: %d\n", mapped_ptr_paa == &paa); + +// (G) use_device_addr/map with different base-array/pointer. +// Address translation should happen for &paa. +// CHECK: G: 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa) + printf("G: %d\n", mapped_ptr_paa == &paa); + +#pragma omp target exit data map(release : g, h, ph, paa) + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp new file mode 100644 index 0000000..9360db4 --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp @@ -0,0 +1,166 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_addr on a reference variable. +// The corresponding data is not previously mapped. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int g_ptee; +int &g = g_ptee; + +int h_ptee[10]; +int (&h)[10] = h_ptee; + +int *ph_ptee = &h_ptee[0]; +int *&ph = ph_ptee; +int *paa_ptee[10][10]; + +struct S { + int *(&paa)[10][10] = paa_ptee; + + void f1(int i) { + paa[0][2] = &g; + + void *original_addr_g = &g; + void *original_addr_h = &h; + void *original_addr_ph = &ph; + void *original_addr_paa = &paa; + +// (A) No corresponding item, lookup should fail. +// CHECK: A: 1 1 1 +#pragma omp target data use_device_addr(g) + { + void *mapped_ptr_g = + omp_get_mapped_ptr(original_addr_g, omp_get_default_device()); + printf("A: %d %d %d\n", mapped_ptr_g == nullptr, + mapped_ptr_g != original_addr_g, (void *)&g == nullptr); + } + +// (B) Lookup should succeed. +// CHECK: B: 1 1 1 +#pragma omp target data map(g) use_device_addr(g) + { + void *mapped_ptr_g = + omp_get_mapped_ptr(original_addr_g, omp_get_default_device()); + printf("B: %d %d %d\n", mapped_ptr_g != nullptr, + mapped_ptr_g != original_addr_g, &g == mapped_ptr_g); + } + +// (C) No corresponding item, lookup should fail. +// CHECK: C: 1 1 1 +#pragma omp target data use_device_addr(h) + { + void *mapped_ptr_h = + omp_get_mapped_ptr(original_addr_h, omp_get_default_device()); + printf("C: %d %d %d\n", mapped_ptr_h == nullptr, + mapped_ptr_h != original_addr_h, (void *)&h == nullptr); + } + +// (D) Lookup should succeed. +// CHECK: D: 1 1 1 +#pragma omp target data map(h) use_device_addr(h) + { + void *mapped_ptr_h = + omp_get_mapped_ptr(original_addr_h, omp_get_default_device()); + printf("D: %d %d %d\n", mapped_ptr_h != nullptr, + mapped_ptr_h != original_addr_h, &h == mapped_ptr_h); + } + +// (E) No corresponding item, lookup should fail. +// CHECK: E: 1 1 1 +#pragma omp target data use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("E: %d %d %d\n", mapped_ptr_ph == nullptr, + mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr); + } + +// (F) Lookup should succeed. +// CHECK: F: 1 1 1 +#pragma omp target data map(ph) use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("F: %d %d %d\n", mapped_ptr_ph != nullptr, + mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph); + } + +// (G) Maps pointee only, but use_device_addr operand is pointer. +// Lookup should fail. +// CHECK: G: 1 1 1 +#pragma omp target data map(ph[0 : 1]) use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("G: %d %d %d\n", mapped_ptr_ph == nullptr, + mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr); + } + +// (H) Maps both pointee and pointer. Lookup for pointer should succeed. +// CHECK: H: 1 1 1 +#pragma omp target data map(ph[0 : 1]) map(ph) use_device_addr(ph) + { + void *mapped_ptr_ph = + omp_get_mapped_ptr(original_addr_ph, omp_get_default_device()); + printf("H: %d %d %d\n", mapped_ptr_ph != nullptr, + mapped_ptr_ph != original_addr_ph, &ph == mapped_ptr_ph); + } + +// (I) No corresponding item, lookup should fail. +// CHECK: I: 1 1 1 +#pragma omp target data use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("I: %d %d %d\n", mapped_ptr_paa == nullptr, + mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr); + } + +// (J) Maps pointee only, but use_device_addr operand is pointer. +// Lookup should fail. +// CHECK: J: 1 1 1 +#pragma omp target data map(paa[0][2][0]) use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("J: %d %d %d\n", mapped_ptr_paa == nullptr, + mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr); + } + +// (K) Lookup should succeed. +// CHECK: K: 1 1 1 +#pragma omp target data map(paa) use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("K: %d %d %d\n", mapped_ptr_paa != nullptr, + mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa); + } + +// (L) Maps both pointee and pointer. Lookup for pointer should succeed. +// CHECK: L: 1 1 1 +#pragma omp target data map(paa[0][2][0]) map(paa) use_device_addr(paa) + { + void *mapped_ptr_paa = + omp_get_mapped_ptr(original_addr_paa, omp_get_default_device()); + printf("L: %d %d %d\n", mapped_ptr_paa != nullptr, + mapped_ptr_paa != original_addr_paa, &paa == mapped_ptr_paa); + } + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/target_use_device_addr.c b/offload/test/mapping/use_device_addr/target_use_device_addr.c index 5c2bb8a..4a9dbe2 100644 --- a/offload/test/mapping/target_use_device_addr.c +++ b/offload/test/mapping/use_device_addr/target_use_device_addr.c @@ -12,7 +12,9 @@ int main() { printf("%d, %p\n", xp[1], &xp[1]); #pragma omp target data use_device_addr(xp[1 : 3]) map(tofrom : x) #pragma omp target is_device_ptr(xp) - { xp[1] = 222; } + { + xp[1] = 222; + } // CHECK: 222 printf("%d, %p\n", xp[1], &xp[1]); } diff --git a/offload/test/mapping/target_wrong_use_device_addr.c b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c index 7a5babd..28ec685 100644 --- a/offload/test/mapping/target_wrong_use_device_addr.c +++ b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c @@ -14,7 +14,7 @@ int main() { // CHECK: host addr=0x[[#%x,HOST_ADDR:]] fprintf(stderr, "host addr=%p\n", x); -#pragma omp target data map(to : x [0:10]) +#pragma omp target data map(to : x[0 : 10]) { // CHECK: omptarget device 0 info: variable x does not have a valid device // counterpart @@ -27,4 +27,3 @@ int main() { return 0; } - diff --git a/offload/test/mapping/array_section_use_device_ptr.c b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c index 86e2875..4cfcce2 100644 --- a/offload/test/mapping/array_section_use_device_ptr.c +++ b/offload/test/mapping/use_device_ptr/array_section_use_device_ptr.c @@ -20,7 +20,9 @@ int main() { float *A_dev = NULL; #pragma omp target data use_device_ptr(A) - { A_dev = A; } + { + A_dev = A; + } #pragma omp target exit data map(delete : A[FROM : LENGTH]) // CHECK: Success diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp new file mode 100644 index 0000000..a7745de --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_existing.cpp @@ -0,0 +1,100 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_ptr on a variable. +// The corresponding data is mapped on a previous enter_data directive. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int aa[10][10]; +int h[10]; +int *ph = &h[0]; + +struct S { + int (*paa)[10][10] = &aa; + + void f1(int i) { + paa--; + void *original_ph3 = &ph[3]; + void *original_paa102 = &paa[1][0][2]; + +#pragma omp target enter data map(to : ph[3 : 4], paa[1][0][2 : 5]) + void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device()); + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device()); + + // CHECK-COUNT-4: 1 + printf("%d\n", mapped_ptr_ph3 != nullptr); + printf("%d\n", mapped_ptr_paa102 != nullptr); + printf("%d\n", original_ph3 != mapped_ptr_ph3); + printf("%d\n", original_paa102 != mapped_ptr_paa102); + +// (A) Mapped data is within extended address range. Lookup should succeed. +// CHECK: A: 1 +#pragma omp target data use_device_ptr(ph) + printf("A: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (B) use_device_ptr/map on pointer, and pointee already exists. +// Lookup should succeed. +// CHECK: B: 1 +#pragma omp target data map(ph) use_device_ptr(ph) + printf("B: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (C) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// CHECK: C: 1 +#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph) + printf("C: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (D) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// CHECK: D: 1 +#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph) + printf("D: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (E) Mapped data is within extended address range. Lookup should succeed. +// Lookup should succeed. +// CHECK: E: 1 +#pragma omp target data use_device_ptr(paa) + printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +// (F) use_device_ptr/map on pointer, and pointee already exists. +// &paa[0] should be in extended address-range of the existing paa[1][...] +// Lookup should succeed. +// FIXME: However, it currently does not. Might need an RT fix. +// EXPECTED: F: 1 +// CHECK: F: 0 +#pragma omp target data map(paa) use_device_ptr(paa) + printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +// (G) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// CHECK: G: 1 +#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa) + printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +// (H) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// CHECK: H: 1 +#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa) + printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +#pragma omp target exit data map(release : ph[3 : 4], paa[1][0][2 : 5]) + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp new file mode 100644 index 0000000..fe3cdb5 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp @@ -0,0 +1,125 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_ptr on a variable. +// The corresponding data is not previously mapped. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int aa[10][10]; +int h[10]; +int *ph = &h[0]; + +struct S { + int (*paa)[10][10] = &aa; + + void f1(int i) { + paa--; + void *original_addr_ph3 = &ph[3]; + void *original_addr_paa102 = &paa[1][0][2]; + +// (A) No corresponding item, lookup should fail. +// CHECK: A: 1 1 1 +#pragma omp target data use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_addr_ph3, ph == nullptr); + } + +// (B) use_device_ptr/map on pointer, and pointee does not exist. +// Lookup should fail. +// CHECK: B: 1 1 1 +#pragma omp target data map(ph) use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_addr_ph3, ph == nullptr); + } + +// (C) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// CHECK: C: 1 1 1 +#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr, + mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3); + } + +// (D) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// CHECK: D: 1 1 1 +#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, + mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3); + } + +// (E) No corresponding item, lookup should fail. +// CHECK: E: 1 1 1 +#pragma omp target data use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr, + mapped_ptr_paa102 != original_addr_paa102, paa == nullptr); + } + +// (F) use_device_ptr/map on pointer, and pointee does not exist. +// Lookup should fail. +// CHECK: F: 1 1 1 +#pragma omp target data map(paa) use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr, + mapped_ptr_paa102 != original_addr_paa102, paa == nullptr); + } + +// (G) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// CHECK: G: 1 1 1 +#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr, + mapped_ptr_paa102 != original_addr_paa102, + &paa[1][0][2] == mapped_ptr_paa102); + } + +// (H) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// CHECK: H: 1 1 1 +#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr, + mapped_ptr_paa102 != original_addr_paa102, + &paa[1][0][2] == mapped_ptr_paa102); + } + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp new file mode 100644 index 0000000..66e65de --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_existing.cpp @@ -0,0 +1,111 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_ptr on a reference variable. +// The corresponding data is mapped on a previous enter_data directive. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int aa[10][10]; +int (*paa_ptee)[10][10] = &aa; + +int h[10]; +int *ph_ptee = &h[0]; +int *&ph = ph_ptee; + +struct S { + int (*&paa)[10][10] = paa_ptee; + + void f1(int i) { + paa--; + void *original_ph3 = &ph[3]; + void *original_paa102 = &paa[1][0][2]; + +#pragma omp target enter data map(to : ph[3 : 4], paa[1][0][2 : 5]) + void *mapped_ptr_ph3 = omp_get_mapped_ptr(&ph[3], omp_get_default_device()); + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(&paa[1][0][2], omp_get_default_device()); + + // CHECK-COUNT-4: 1 + printf("%d\n", mapped_ptr_ph3 != nullptr); + printf("%d\n", mapped_ptr_paa102 != nullptr); + printf("%d\n", original_ph3 != mapped_ptr_ph3); + printf("%d\n", original_paa102 != mapped_ptr_paa102); + +// (A) Mapped data is within extended address range. Lookup should succeed. +// EXPECTED: A: 1 +// CHECK: A: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data use_device_ptr(ph) + printf("A: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (B) use_device_ptr/map on pointer, and pointee already exists. +// Lookup should succeed. +// EXPECTED: B: 1 +// CHECK: B: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) use_device_ptr(ph) + printf("B: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (C) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// EXPECTED: C: 1 +// CHECK: C: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph) + printf("C: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (D) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// EXPECTED: D: 1 +// CHECK: D: 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph) + printf("D: %d\n", mapped_ptr_ph3 == &ph[3]); + +// (E) Mapped data is within extended address range. Lookup should succeed. +// Lookup should succeed. +// CHECK: E: 1 +#pragma omp target data use_device_ptr(paa) + printf("E: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +// (F) use_device_ptr/map on pointer, and pointee already exists. +// &paa[0] should be in extended address-range of the existing paa[1][...] +// Lookup should succeed. +// FIXME: However, it currently does not. Might need an RT fix. +// EXPECTED: F: 1 +// CHECK: F: 0 +#pragma omp target data map(paa) use_device_ptr(paa) + printf("F: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +// (G) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// CHECK: G: 1 +#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa) + printf("G: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +// (H) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// CHECK: H: 1 +#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa) + printf("H: %d\n", mapped_ptr_paa102 == &paa[1][0][2]); + +#pragma omp target exit data map(release : ph[3 : 4], paa[1][0][2 : 5]) + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp new file mode 100644 index 0000000..419ab3e --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp @@ -0,0 +1,136 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include <omp.h> +#include <stdio.h> + +// Test for various cases of use_device_ptr on a reference variable. +// The corresponding data is not previously mapped. + +// Note that this tests for the current behavior wherein if a lookup fails, +// the runtime returns nullptr, instead of the original host-address. +// That was compatible with OpenMP 5.0, where it was a user error if +// corresponding storage didn't exist, but with 5.1+, the runtime needs to +// return the host address, as it needs to assume that the host-address is +// device-accessible, as the user has guaranteed it. +// Once the runtime returns the original host-address when the lookup fails, the +// test will need to be updated. + +int aa[10][10]; +int (*paa_ptee)[10][10] = &aa; + +int h[10]; +int *ph_ptee = &h[0]; +int *&ph = ph_ptee; + +struct S { + int (*&paa)[10][10] = paa_ptee; + + void f1(int i) { + paa--; + void *original_addr_ph3 = &ph[3]; + void *original_addr_paa102 = &paa[1][0][2]; + +// (A) No corresponding item, lookup should fail. +// EXPECTED: A: 1 1 1 +// CHECK: A: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_addr_ph3, ph == nullptr); + } + +// (B) use_device_ptr/map on pointer, and pointee does not exist. +// Lookup should fail. +// EXPECTED: B: 1 1 1 +// CHECK: B: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr, + mapped_ptr_ph3 != original_addr_ph3, ph == nullptr); + } + +// (C) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// EXPECTED: C: 1 1 1 +// CHECK: C: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph[3 : 2]) use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("C: %d %d %d\n", mapped_ptr_ph3 != nullptr, + mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3); + } + +// (D) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// EXPECTED: D: 1 1 1 +// CHECK: D: 1 1 0 +// FIXME: ph is not being privatized in the region. +#pragma omp target data map(ph) map(ph[3 : 2]) use_device_ptr(ph) + { + void *mapped_ptr_ph3 = + omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device()); + printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr, + mapped_ptr_ph3 != original_addr_ph3, &ph[3] == mapped_ptr_ph3); + } + +// (E) No corresponding item, lookup should fail. +// CHECK: E: 1 1 1 +#pragma omp target data use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr, + mapped_ptr_paa102 != original_addr_paa102, paa == nullptr); + } + +// (F) use_device_ptr/map on pointer, and pointee does not exist. +// Lookup should fail. +// CHECK: F: 1 1 1 +#pragma omp target data map(paa) use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr, + mapped_ptr_paa102 != original_addr_paa102, paa == nullptr); + } + +// (G) map on pointee: base-pointer of map matches use_device_ptr operand. +// Lookup should succeed. +// CHECK: G: 1 1 1 +#pragma omp target data map(paa[1][0][2]) use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("G: %d %d %d\n", mapped_ptr_paa102 != nullptr, + mapped_ptr_paa102 != original_addr_paa102, + &paa[1][0][2] == mapped_ptr_paa102); + } + +// (H) map on pointer and pointee. Base-pointer of map on pointee matches +// use_device_ptr operand. +// Lookup should succeed. +// CHECK: H: 1 1 1 +#pragma omp target data map(paa) map(paa[1][0][2]) use_device_ptr(paa) + { + void *mapped_ptr_paa102 = + omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device()); + printf("H: %d %d %d\n", mapped_ptr_paa102 != nullptr, + mapped_ptr_paa102 != original_addr_paa102, + &paa[1][0][2] == mapped_ptr_paa102); + } + } +}; + +S s1; +int main() { s1.f1(1); } diff --git a/offload/test/offloading/fortran/declare-target-automap.f90 b/offload/test/offloading/fortran/declare-target-automap.f90 new file mode 100644 index 0000000..b9c2d34 --- /dev/null +++ b/offload/test/offloading/fortran/declare-target-automap.f90 @@ -0,0 +1,37 @@ +!Offloading test for AUTOMAP modifier in declare target enter +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic +program automap_program + use iso_c_binding, only: c_loc + use omp_lib, only: omp_get_default_device, omp_target_is_present + integer, parameter :: N = 10 + integer :: i + integer, allocatable, target :: automap_array(:) + !$omp declare target enter(automap:automap_array) + + ! false since the storage is not present even though the descriptor is present + write (*, *) omp_target_is_present(c_loc(automap_array), omp_get_default_device()) + ! CHECK: 0 + + allocate (automap_array(N)) + ! true since the storage should be allocated and reference count incremented by the allocate + write (*, *) omp_target_is_present(c_loc(automap_array), omp_get_default_device()) + ! CHECK: 1 + + ! since storage is present this should not be a runtime error + !$omp target teams loop + do i = 1, N + automap_array(i) = i + end do + + !$omp target update from(automap_array) + write (*, *) automap_array + ! CHECK: 1 2 3 4 5 6 7 8 9 10 + + deallocate (automap_array) + + ! automap_array should have it's storage unmapped on device here + write (*, *) omp_target_is_present(c_loc(automap_array), omp_get_default_device()) + ! CHECK: 0 +end program diff --git a/offload/test/offloading/mandatory_but_no_devices.c b/offload/test/offloading/mandatory_but_no_devices.c index ecdee72..df8a5f3 100644 --- a/offload/test/offloading/mandatory_but_no_devices.c +++ b/offload/test/offloading/mandatory_but_no_devices.c @@ -3,6 +3,47 @@ // device. This behavior is proposed for OpenMP 5.2 in OpenMP spec github // issue 2669. +// AMD Tests +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -DDIR=target +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -DDIR='target teams' +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -DDIR='target data map(X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa \ +// RUN: -DDIR='target enter data map(to:X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa \ +// RUN: -DDIR='target exit data map(from:X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa \ +// RUN: -DDIR='target update to(X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa \ +// RUN: -DDIR='target update from(X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory ROCR_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-amdgcn-amd-amdhsa 2>&1 | \ +// RUN: %fcheck-amdgcn-amd-amdhsa + +// Nvidia Tests // RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DDIR=target // RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ // RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ @@ -42,8 +83,6 @@ // RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ // RUN: %fcheck-nvptx64-nvidia-cuda -// REQUIRES: nvptx64-nvidia-cuda - #include <omp.h> #include <stdio.h> diff --git a/offload/test/offloading/memory_manager.cpp b/offload/test/offloading/memory_manager.cpp index fba1e4a..d6d8697 100644 --- a/offload/test/offloading/memory_manager.cpp +++ b/offload/test/offloading/memory_manager.cpp @@ -1,7 +1,5 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// REQUIRES: nvidiagpu - #include <omp.h> #include <cassert> diff --git a/offload/test/offloading/strided_multiple_update.c b/offload/test/offloading/strided_multiple_update.c new file mode 100644 index 0000000..a3e8d10 --- /dev/null +++ b/offload/test/offloading/strided_multiple_update.c @@ -0,0 +1,62 @@ +// This test checks that #pragma omp target update from(data1[0:3:4], +// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays +// from the device to the host. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 12; + double data1[len], data2[len]; + +// 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: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data1[i]); + printf("\ndata2: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data2[i]); + printf("\n\n"); + +#pragma omp target data map(to : data1[0 : len], data2[0 : len]) + { +// Modify arrays on device +#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:3:4] // indices 0,4,8 +// data2[0:2:5] // indices 0,5 +#pragma omp target update from(data1[0 : 3 : 4], data2[0 : 2 : 5]) + } + + printf("device array values after update from:\n"); + printf("data1: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data1[i]); + printf("\ndata2: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data2[i]); + printf("\n\n"); + + // CHECK: data1: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0 + // CHECK: data2: 0.0 10.0 20.0 30.0 40.0 50.0 60.0 70.0 80.0 90.0 100.0 110.0 + + // CHECK: data1: 0.0 1.0 2.0 3.0 8.0 5.0 6.0 7.0 16.0 9.0 10.0 11.0 + // CHECK: data2: 100.0 10.0 20.0 30.0 40.0 150.0 60.0 70.0 80.0 90.0 100.0 + // 110.0 +} diff --git a/offload/test/offloading/strided_partial_update.c b/offload/test/offloading/strided_partial_update.c new file mode 100644 index 0000000..15d477f --- /dev/null +++ b/offload/test/offloading/strided_partial_update.c @@ -0,0 +1,63 @@ +// This test checks that #pragma omp target update from(data[0:4:3]) correctly +// updates every third element (stride 3) from the device to the host, partially +// across the array + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 11; + double data[len]; + +#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]) + { +// Modify arrays on device +#pragma omp target + for (int i = 0; i < len; i++) + data[i] += i; + +#pragma omp target update from(data[0 : 4 : 3]) // indices 0,3,6,9 + } + + printf("device array values after update from:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 2.000000 + // CHECK: 3.000000 + // CHECK: 4.000000 + // CHECK: 5.000000 + // CHECK: 6.000000 + // CHECK: 7.000000 + // CHECK: 8.000000 + // CHECK: 9.000000 + // CHECK: 10.000000 + + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 2.000000 + // CHECK: 6.000000 + // CHECK: 4.000000 + // CHECK: 5.000000 + // CHECK: 12.000000 + // CHECK: 7.000000 + // CHECK: 8.000000 + // CHECK: 18.000000 + // CHECK: 10.000000 +} diff --git a/offload/test/offloading/strided_update.c b/offload/test/offloading/strided_update.c new file mode 100644 index 0000000..fe875b7 --- /dev/null +++ b/offload/test/offloading/strided_update.c @@ -0,0 +1,54 @@ +// This test checks that "update from" clause in OpenMP is supported when the +// elements are updated in a non-contiguous manner. This test checks that +// #pragma omp target update from(data[0:4:2]) correctly updates only every +// other element (stride 2) from the device to the host + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 8; + double data[len]; +#pragma omp target map(tofrom : len, 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 : len, data[0 : len]) + { +// Modify arrays on device +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += i; + } + +#pragma omp target update from(data[0 : 4 : 2]) + } + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 4.000000 + // CHECK: 3.000000 + // CHECK: 8.000000 + // CHECK: 5.000000 + // CHECK: 12.000000 + // CHECK: 7.000000 + // CHECK-NOT: 2.000000 + // CHECK-NOT: 6.000000 + // CHECK-NOT: 10.000000 + // CHECK-NOT: 14.000000 + + printf("from target array results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + return 0; +} diff --git a/offload/test/sanitizer/use_after_free_2.c b/offload/test/sanitizer/use_after_free_2.c index 02aa453..1c1e097 100644 --- a/offload/test/sanitizer/use_after_free_2.c +++ b/offload/test/sanitizer/use_after_free_2.c @@ -10,6 +10,9 @@ // UNSUPPORTED: s390x-ibm-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu-LTO +// If offload memory pooling is enabled for a large allocation, reuse error is +// not detected. UNSUPPORTED: large_allocation_memory_pool + #include <omp.h> int main() { diff --git a/offload/test/sanitizer/use_after_free_3.c b/offload/test/sanitizer/use_after_free_3.c new file mode 100644 index 0000000..9d88614 --- /dev/null +++ b/offload/test/sanitizer/use_after_free_3.c @@ -0,0 +1,37 @@ +// clang-format off +// RUN: %libomptarget-compileopt-generic +// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD=1024 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK +// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK-PASS +// clang-format on + +// If offload memory pooling is enabled for a large allocation, reuse error is +// not detected. Run the test w/ and w/o ENV var override on memory pooling +// threshold. REQUIRES: large_allocation_memory_pool + +#include <omp.h> +#include <stdio.h> + +int main() { + int N = (1 << 30); + char *A = (char *)malloc(N); + char *P; +#pragma omp target map(A[ : N]) map(from : P) + { + P = &A[N / 2]; + *P = 3; + } + // clang-format off +// CHECK: OFFLOAD ERROR: memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}} +// CHECK: Device pointer [[PTR]] points into prior host-issued allocation: +// CHECK: Last deallocation: +// CHECK: Last allocation of size 1073741824 +// clang-format on +#pragma omp target + { + *P = 5; + } + + // CHECK-PASS: PASS + printf("PASS\n"); + return 0; +} diff --git a/offload/test/tools/llvm-omp-device-info.c b/offload/test/tools/llvm-omp-device-info.c index 6f49730..1ce8d4a 100644 --- a/offload/test/tools/llvm-omp-device-info.c +++ b/offload/test/tools/llvm-omp-device-info.c @@ -2,5 +2,5 @@ // // Just check any device was found and something is printed // -// CHECK: Found {{[1-9].*}} devices: -// CHECK: Device 0: +// CHECK: Num Devices: {{[1-9].*}} +// CHECK: [{{[1-9A-Za-z].*}}] diff --git a/offload/test/tools/offload-tblgen/default_returns.td b/offload/test/tools/offload-tblgen/default_returns.td index e919492..41949db 100644 --- a/offload/test/tools/offload-tblgen/default_returns.td +++ b/offload/test/tools/offload-tblgen/default_returns.td @@ -6,13 +6,11 @@ include "APIDefs.td" -def : Handle { - let name = "ol_foo_handle_t"; +def ol_foo_handle_t : Handle { let desc = "Example handle type"; } -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ diff --git a/offload/test/tools/offload-tblgen/entry_points.td b/offload/test/tools/offload-tblgen/entry_points.td index c66d5b4..94ea820 100644 --- a/offload/test/tools/offload-tblgen/entry_points.td +++ b/offload/test/tools/offload-tblgen/entry_points.td @@ -4,8 +4,7 @@ include "APIDefs.td" -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ diff --git a/offload/test/tools/offload-tblgen/functions_basic.td b/offload/test/tools/offload-tblgen/functions_basic.td index dec9357..2802c78 100644 --- a/offload/test/tools/offload-tblgen/functions_basic.td +++ b/offload/test/tools/offload-tblgen/functions_basic.td @@ -6,8 +6,7 @@ include "APIDefs.td" -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ diff --git a/offload/test/tools/offload-tblgen/functions_code_loc.td b/offload/test/tools/offload-tblgen/functions_code_loc.td index aec2012..8d7aa00 100644 --- a/offload/test/tools/offload-tblgen/functions_code_loc.td +++ b/offload/test/tools/offload-tblgen/functions_code_loc.td @@ -7,8 +7,7 @@ include "APIDefs.td" -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ diff --git a/offload/test/tools/offload-tblgen/functions_ranged_param.td b/offload/test/tools/offload-tblgen/functions_ranged_param.td index d0996b2..1ce8b39 100644 --- a/offload/test/tools/offload-tblgen/functions_ranged_param.td +++ b/offload/test/tools/offload-tblgen/functions_ranged_param.td @@ -8,13 +8,11 @@ include "APIDefs.td" -def : Handle { - let name = "some_handle_t"; +def some_handle_t : Handle { let desc = "An example handle type"; } -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ diff --git a/offload/test/tools/offload-tblgen/print_enum.td b/offload/test/tools/offload-tblgen/print_enum.td index 97f8696..c7573a9 100644 --- a/offload/test/tools/offload-tblgen/print_enum.td +++ b/offload/test/tools/offload-tblgen/print_enum.td @@ -4,8 +4,7 @@ include "APIDefs.td" -def : Enum { - let name = "my_enum_t"; +def my_enum_t : Enum { let desc = "An example enum"; let etors =[ Etor<"VALUE_ONE", "The first enum value">, diff --git a/offload/test/tools/offload-tblgen/print_function.td b/offload/test/tools/offload-tblgen/print_function.td index ce1fe4c..74b39f1 100644 --- a/offload/test/tools/offload-tblgen/print_function.td +++ b/offload/test/tools/offload-tblgen/print_function.td @@ -5,13 +5,11 @@ include "APIDefs.td" -def : Handle { - let name = "ol_foo_handle_t"; +def ol_foo_handle_t : Handle { let desc = "Example handle type"; } -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ diff --git a/offload/test/tools/offload-tblgen/type_tagged_enum.td b/offload/test/tools/offload-tblgen/type_tagged_enum.td index 95964e3..b32531a 100644 --- a/offload/test/tools/offload-tblgen/type_tagged_enum.td +++ b/offload/test/tools/offload-tblgen/type_tagged_enum.td @@ -9,13 +9,11 @@ include "APIDefs.td" -def : Handle { - let name = "some_handle_t"; +def some_handle_t: Handle { let desc = "An example handle type"; } -def : Enum { - let name = "my_type_tagged_enum_t"; +def my_type_tagged_enum_t : Enum { let desc = "Example type tagged enum"; let is_typed = 1; let etors = [ @@ -34,8 +32,7 @@ def : Enum { // CHECK-API-NEXT: [some_handle_t] Value three. // CHECK-API-NEXT: MY_TYPE_TAGGED_ENUM_VALUE_THREE = 2, -def : Function { - let name = "FunctionA"; +def FunctionA : Function { let desc = "Function A description"; let details = [ "Function A detailed information" ]; let params = [ |