diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp --- a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp +++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp @@ -1,8 +1,5 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// Wrong results on amdgpu -// XFAIL: amdgcn-amd-amdhsa - #include #include @@ -45,17 +42,23 @@ spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0); // CHECK: 111 222 777 20.00000 1 + 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) +#pragma omp target map(tofrom: spp[0][0]) firstprivate(p) \ + map(from: spp00fa, spp00fca, spp00fb_r) { - printf("%d %d %d\n", spp[0][0].f.a, spp[0][0].f.c.a, - spp[0][0].f.b == reinterpret_cast(p) ? 1 : 0); - // CHECK: 222 777 0 + spp00fa = spp[0][0].f.a; + spp00fca = spp[0][0].f.c.a; + spp00fb_r = spp[0][0].f.b == reinterpret_cast(p) ? 1 : 0; + printf("%d %d %d\n", spp00fa, spp00fca, spp00fb_r); + // XCHECK: 222 777 0 spp[0][0].e = 333; spp[0][0].f.a = 444; spp[0][0].f.c.a = 555; spp[0][0].f.b[1] = 40; } + printf("%d %d %d\n", spp00fa, spp00fca, spp00fb_r); + // CHECK: 222 777 0 printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a, spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0); // CHECK: 333 222 777 40.00000 1 diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp --- a/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp +++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp @@ -1,8 +1,5 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// Wrong results on amdgpu -// XFAIL: amdgcn-amd-amdhsa - #include #include @@ -42,19 +39,25 @@ spp[0][0].g == &y[0] ? 1 : 0); // CHECK: 111 222 20.00000 1 30 1 + 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) +#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1) \ + map(from: spp00fa, spp00fb_r, spp00fg1, spp00fg_r) { - printf("%d %d %d %d\n", spp[0][0].f.a, - spp[0][0].f.b == reinterpret_cast(p) ? 1 : 0, spp[0][0].g[1], - spp[0][0].g == reinterpret_cast(p1) ? 1 : 0); - // CHECK: 222 0 30 0 + spp00fa = spp[0][0].f.a; + spp00fb_r = spp[0][0].f.b == reinterpret_cast(p) ? 1 : 0; + spp00fg1 = spp[0][0].g[1]; + spp00fg_r = spp[0][0].g == reinterpret_cast(p1) ? 1 : 0; + printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r); + // XCHECK: 222 0 30 0 spp[0][0].e = 333; spp[0][0].f.a = 444; spp[0][0].f.b[1] = 40; spp[0][0].g[1] = 50; } + printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r); + // CHECK: 222 0 30 0 printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1], spp[0][0].g == &y[0] ? 1 : 0); diff --git a/openmp/libomptarget/test/mapping/lambda_by_value.cpp b/openmp/libomptarget/test/mapping/lambda_by_value.cpp --- a/openmp/libomptarget/test/mapping/lambda_by_value.cpp +++ b/openmp/libomptarget/test/mapping/lambda_by_value.cpp @@ -1,8 +1,5 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// Wrong results on amdgpu -// XFAIL: amdgcn-amd-amdhsa - #include #include @@ -11,6 +8,13 @@ // CHECK: tgt : [[V2]] [[PX_TGT]] 1 // CHECK: out : [[V2]] [[V2]] [[PX]] [[PY]] +#pragma omp begin declare target +int a = -1, *c; +long b = -1; +const long *d; +int e = -1, *f, g = -1; +#pragma omp end declare target + int main() { int x[10]; long y[8]; @@ -18,18 +22,27 @@ y[1] = 222; auto lambda = [&x, y]() { + a = x[1]; + b = y[1]; + c = &x[0]; + d = &y[0]; printf("lambda: %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]); x[1] = y[1]; }; - printf("before: %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]); intptr_t xp = (intptr_t)&x[0]; #pragma omp target firstprivate(xp) { lambda(); + e = x[1]; + f = &x[0]; + g = (&x[0] != (int *)xp); printf("tgt : %d %p %d\n", x[1], &x[0], (&x[0] != (int *)xp)); } +#pragma omp target update from(a, b, c, d, e, f, g) + printf("lambda: %d %ld %p %p\n", a, b, c, d); + printf("tgt : %d %p %d\n", e, f, g); printf("out : %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]); return 0; diff --git a/openmp/libomptarget/test/mapping/ompx_hold/struct.c b/openmp/libomptarget/test/mapping/ompx_hold/struct.c --- a/openmp/libomptarget/test/mapping/ompx_hold/struct.c +++ b/openmp/libomptarget/test/mapping/ompx_hold/struct.c @@ -1,20 +1,35 @@ // RUN: %libomptarget-compile-generic -fopenmp-extensions // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace -// Wrong results on amdgpu -// XFAIL: amdgcn-amd-amdhsa - #include #include +#pragma omp begin declare target +char *N1, *N2; +int V1, V2; +#pragma omp declare target + #define CHECK_PRESENCE(Var1, Var2, Var3) \ printf(" presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3, \ omp_target_is_present(&(Var1), omp_get_default_device()), \ omp_target_is_present(&(Var2), omp_get_default_device()), \ omp_target_is_present(&(Var3), omp_get_default_device())) +#define CHECK_VALUES_HELPER(N1, N2, Var1, Var2) \ + printf(" values of %s, %s: %d, %d\n", N1, N2, (Var1), (Var2)) + +#define CHECK_VALUES_DELAYED(Var1, Var2) \ + N1 = #Var1; \ + N2 = #Var2; \ + V1 = (Var1); \ + V2 = (Var2); + +#define CHECK_DELAYED_VALUS() \ + _Pragma("omp target update from(N1, N2, V1, V2)") \ + CHECK_VALUES_HELPER(N1, N2, V1, V2) + #define CHECK_VALUES(Var1, Var2) \ - printf(" values of %s, %s: %d, %d\n", #Var1, #Var2, (Var1), (Var2)) + CHECK_VALUES_HELPER(#Var1, #Var2, (Var1), (Var2)) int main() { struct S { @@ -132,8 +147,9 @@ #pragma omp target map(to : s.i, s.j) { // No transfer here even though parent's DynRefCount=1. // CHECK-NEXT: values of s.i, s.j: 21, 31 - CHECK_VALUES(s.i, s.j); + CHECK_VALUES_DELAYED(s.i, s.j); } + CHECK_DELAYED_VALUS(); } // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0 // CHECK-NEXT: values of s.i, s.j: 21, 31 @@ -162,8 +178,9 @@ #pragma omp target map(ompx_hold, to : s.i, s.j) { // No transfer here even though parent's HoldRefCount=1. // CHECK-NEXT: values of s.i, s.j: 21, 31 - CHECK_VALUES(s.i, s.j); + CHECK_VALUES_DELAYED(s.i, s.j); } + CHECK_DELAYED_VALUS(); } // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0 // CHECK-NEXT: values of s.i, s.j: 21, 31 diff --git a/openmp/libomptarget/test/offloading/host_as_target.c b/openmp/libomptarget/test/offloading/host_as_target.c --- a/openmp/libomptarget/test/offloading/host_as_target.c +++ b/openmp/libomptarget/test/offloading/host_as_target.c @@ -7,16 +7,15 @@ // RUN: %libomptarget-compile-run-and-check-generic -// amdgpu does not have a working printf definition -// XFAIL: amdgcn-amd-amdhsa - #include #include static void check(char *X, int Dev) { printf(" host X = %c\n", *X); -#pragma omp target device(Dev) - printf("device X = %c\n", *X); + char DV = -1; +#pragma omp target device(Dev) map(from : DV) + DV = *X; + printf("device X = %c\n", DV); } #define CHECK_DATA() check(&X, DevDefault)