aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c
diff options
context:
space:
mode:
Diffstat (limited to 'libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c')
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c100
1 files changed, 100 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c
new file mode 100644
index 00000000000..ce9632cf1c0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c
@@ -0,0 +1,100 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* This code uses nvptx inline assembly guarded with acc_on_device, which is
+ not optimized away at -O0, and then confuses the target assembler.
+ { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
+
+#include <assert.h>
+#include <openacc.h>
+
+#define N 100
+
+#define GANG_ID(I) \
+ (acc_on_device (acc_device_nvidia) \
+ ? ({unsigned __r; \
+ __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (__r)); \
+ __r; }) : (I))
+
+int
+test_static(int *a, int num_gangs, int sarg)
+{
+ int i, j;
+
+ if (sarg == 0)
+ sarg = 1;
+
+ for (i = 0; i < N / sarg; i++)
+ for (j = 0; j < sarg; j++)
+ assert (a[i*sarg+j] == i % num_gangs);
+}
+
+int
+test_nonstatic(int *a, int gangs)
+{
+ int i, j;
+
+ for (i = 0; i < N; i+=gangs)
+ for (j = 0; j < gangs; j++)
+ assert (a[i+j] == i/gangs);
+}
+
+int
+main ()
+{
+ int a[N];
+ int i, x;
+
+#pragma acc parallel loop gang (static:*) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_nonstatic (a, 10);
+
+#pragma acc parallel loop gang (static:1) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_static (a, 10, 1);
+
+#pragma acc parallel loop gang (static:2) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_static (a, 10, 2);
+
+#pragma acc parallel loop gang (static:5) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_static (a, 10, 5);
+
+#pragma acc parallel loop gang (static:20) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_static (a, 10, 20);
+
+ /* Non-static gang. */
+#pragma acc parallel loop gang num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_nonstatic (a, 10);
+
+ /* Static arguments with a variable expression. */
+
+ x = 20;
+#pragma acc parallel loop gang (static:0+x) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_static (a, 10, 20);
+
+ x = 20;
+#pragma acc parallel loop gang (static:x) num_gangs (10)
+ for (i = 0; i < 100; i++)
+ a[i] = GANG_ID (i);
+
+ test_static (a, 10, 20);
+
+ return 0;
+}