[Openmp-commits] [openmp] 747af24 - [OpenMP] Allow more tests to run on AMDGPU
Johannes Doerfert via Openmp-commits
openmp-commits at lists.llvm.org
Mon Mar 27 21:30:48 PDT 2023
Author: Johannes Doerfert
Date: 2023-03-27T21:30:22-07:00
New Revision: 747af2415519b10259cdd12c700b0aa2f8946c97
URL: https://github.com/llvm/llvm-project/commit/747af2415519b10259cdd12c700b0aa2f8946c97
DIFF: https://github.com/llvm/llvm-project/commit/747af2415519b10259cdd12c700b0aa2f8946c97.diff
LOG: [OpenMP] Allow more tests to run on AMDGPU
This basically works around the printf issue to increase test coverage.
Differential Revision: https://reviews.llvm.org/D146838
Added:
Modified:
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
openmp/libomptarget/test/mapping/lambda_by_value.cpp
openmp/libomptarget/test/mapping/ompx_hold/struct.c
openmp/libomptarget/test/offloading/host_as_target.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
index e6a44f2b99336..c6c5657ae6166 100644
--- 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 <cstdio>
#include <cstdlib>
@@ -45,17 +42,23 @@ int main() {
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<void *>(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<void *>(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
index c0f09ac61b235..a9e3f05e0f5fd 100644
--- 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 <cstdio>
#include <cstdlib>
@@ -42,19 +39,25 @@ int main() {
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<void *>(p) ? 1 : 0, spp[0][0].g[1],
- spp[0][0].g == reinterpret_cast<void *>(p1) ? 1 : 0);
- // CHECK: 222 0 30 0
+ spp00fa = spp[0][0].f.a;
+ spp00fb_r = spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0;
+ spp00fg1 = spp[0][0].g[1];
+ spp00fg_r = spp[0][0].g == reinterpret_cast<void *>(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
index 383e39ef27d44..5516dedd72a98 100644
--- 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 <stdint.h>
#include <stdio.h>
@@ -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 @@ int main() {
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
index 16936f9b8933d..f61ba44dc2a3d 100644
--- 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 <omp.h>
#include <stdio.h>
+#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 @@ int main() {
#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 @@ int main() {
#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
index 22869e86226b9..8ab991b723388 100644
--- 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 <omp.h>
#include <stdio.h>
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)
More information about the Openmp-commits
mailing list