From 1467100fc72562a59f70cdd4e05f6c810d1fadcc Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 19 May 2021 11:58:49 +0200 Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/private-atomic-1.c' [PR83812] ... to at least document/test/XFAIL nvptx offloading: PR83812 "operation not supported on global/shared address space". libgomp/ PR target/83812 * testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c: New. --- .../libgomp.oacc-c-c++-common/private-atomic-1.c | 37 ++++++++++++++++++++++ 1 file changed, 37 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c new file mode 100644 index 0000000..77197d8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c @@ -0,0 +1,37 @@ +// 'atomic' access of thread-private variable + +#include + +int main (void) +{ + int res; + + res = 0; +#pragma acc parallel reduction(+: res) + { +#pragma acc loop vector reduction(+: res) + for (int i = 0; i < 2322; i++) + { + int v = -222; + +#pragma acc loop seq + for (int j = 0; j < 121; ++j) + { +#pragma acc atomic update + ++v; + /* nvptx offloading: PR83812 "operation not supported on global/shared address space". + { dg-output "(\n|\r\n|\r)libgomp: cuStreamSynchronize error: operation not supported on global/shared address space(\n|\r\n|\r)$" { target openacc_nvidia_accel_selected } } + Scan for what we expect in the "XFAILed" case (without actually XFAILing). + { dg-shouldfail "XFAILed" { openacc_nvidia_accel_selected } } + ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all. + { dg-final { if { [dg-process-target { xfail openacc_nvidia_accel_selected }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } + ... so that we still get an XFAIL visible in the log. */ + } + + res += (v == -222 + 121); + } + } + assert (res == 2322); + + return 0; +} -- 2.7.4