caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[OpenMP] Allow more tests to run on AMDGPU

Browse Source 
This basically works around the printf issue to increase test coverage.

Differential Revision: https://reviews.llvm.org/D146838
This commit is contained in:
Johannes Doerfert
2023-03-24 11:43:38 -07:00
parent 151b58d802
commit 747af24155
5 changed files with 65 additions and 30 deletions
Download Patch File Download Diff File Expand all files Collapse all files

17
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
View File

@@ -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

19
openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
View File

@@ -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);

21
openmp/libomptarget/test/mapping/lambda_by_value.cpp
View File

@@ -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;

29
openmp/libomptarget/test/mapping/ompx_hold/struct.c
View File

@@ -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

9
openmp/libomptarget/test/offloading/host_as_target.c
View File

@@ -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)