diff options
| -rw-r--r-- | libgomp/testsuite/libgomp.c/target-32.c | 9 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/thread-limit-2.c | 3 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/usleep.h | 24 |
3 files changed, 31 insertions, 5 deletions
diff --git a/libgomp/testsuite/libgomp.c/target-32.c b/libgomp/testsuite/libgomp.c/target-32.c index 233877b..1444523 100644 --- a/libgomp/testsuite/libgomp.c/target-32.c +++ b/libgomp/testsuite/libgomp.c/target-32.c @@ -1,5 +1,6 @@ #include <stdlib.h> #include <unistd.h> +#include "usleep.h" int main () { @@ -18,28 +19,28 @@ int main () #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[3]) { - usleep (1000); + tgt_usleep (1000); #pragma omp atomic update b |= 4; } #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[4]) { - usleep (5000); + tgt_usleep (5000); #pragma omp atomic update b |= 1; } #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[5]) { - usleep (5000); + tgt_usleep (5000); #pragma omp atomic update c |= 8; } #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[6]) { - usleep (1000); + tgt_usleep (1000); #pragma omp atomic update c |= 2; } diff --git a/libgomp/testsuite/libgomp.c/thread-limit-2.c b/libgomp/testsuite/libgomp.c/thread-limit-2.c index 1a97fb6..dc247a7 100644 --- a/libgomp/testsuite/libgomp.c/thread-limit-2.c +++ b/libgomp/testsuite/libgomp.c/thread-limit-2.c @@ -4,6 +4,7 @@ #include <stdlib.h> #include <unistd.h> #include <omp.h> +#include "usleep.h" int main () @@ -48,7 +49,7 @@ main () v = ++cnt; if (v > 6) abort (); - usleep (10000); + tgt_usleep (10000); #pragma omp atomic --cnt; } diff --git a/libgomp/testsuite/libgomp.c/usleep.h b/libgomp/testsuite/libgomp.c/usleep.h new file mode 100644 index 0000000..c01aaa0 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/usleep.h @@ -0,0 +1,24 @@ +#include <unistd.h> + +int +nvptx_usleep (useconds_t d) +{ + /* This function serves as a replacement for usleep in + this test case. It does not even attempt to be functionally + equivalent - we just want some sort of delay. */ + int i; + int N = d * 2000; + for (i = 0; i < N; i++) + asm volatile ("" : : : "memory"); + return 0; +} + +#pragma omp declare variant (nvptx_usleep) match(construct={target},device={arch(nvptx)}) +#pragma omp declare variant (usleep) match(user={condition(1)}) +int +tgt_usleep (useconds_t d) +{ + return 0; +} + +#pragma omp declare target to (nvptx_usleep, tgt_usleep) |