libgomp/testsuite/libgomp.oacc-c-c++-common/gang-static-2.c

   1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
   2 /* This code uses nvptx inline assembly guarded with acc_on_device, which is
   3    not optimized away at -O0, and then confuses the target assembler.
   4    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
   5
   6 #include <assert.h>
   7 #include <openacc.h>
   8
   9 #define N 100
  10
  11 #define GANG_ID(I)                                              \
  12   (acc_on_device (acc_device_nvidia)                            \
  13    ? ({unsigned __r;                                            \
  14        __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (__r)); \
  15        __r; }) : (I))
  16
  17 void
  18 test_static(int *a, int num_gangs, int sarg)
  19 {
  20   int i, j;
  21
  22   if (sarg == 0)
  23     sarg = 1;
  24
  25   for (i = 0; i < N / sarg; i++)
  26     for (j = 0; j < sarg; j++)
  27       assert (a[i*sarg+j] == i % num_gangs);
  28 }
  29
  30 void
  31 test_nonstatic(int *a, int gangs)
  32 {
  33   int i, j;
  34
  35   for (i = 0; i < N; i+=gangs)
  36     for (j = 0; j < gangs; j++)
  37       assert (a[i+j] == i/gangs);
  38 }
  39
  40 int
  41 main ()
  42 {
  43   int a[N];
  44   int i, x;
  45
  46 #pragma acc parallel loop gang (static:*) num_gangs (10)
  47   for (i = 0; i < 100; i++)
  48     a[i] = GANG_ID (i);
  49
  50   test_nonstatic (a, 10);
  51
  52 #pragma acc parallel loop gang (static:1) num_gangs (10)
  53   for (i = 0; i < 100; i++)
  54     a[i] = GANG_ID (i);
  55
  56   test_static (a, 10, 1);
  57
  58 #pragma acc parallel loop gang (static:2) num_gangs (10)
  59   for (i = 0; i < 100; i++)
  60     a[i] = GANG_ID (i);
  61
  62   test_static (a, 10, 2);
  63
  64 #pragma acc parallel loop gang (static:5) num_gangs (10)
  65   for (i = 0; i < 100; i++)
  66     a[i] = GANG_ID (i);
  67
  68   test_static (a, 10, 5);
  69
  70 #pragma acc parallel loop gang (static:20) num_gangs (10)
  71   for (i = 0; i < 100; i++)
  72     a[i] = GANG_ID (i);
  73
  74   test_static (a, 10, 20);
  75
  76   /* Non-static gang.  */
  77 #pragma acc parallel loop gang num_gangs (10)
  78   for (i = 0; i < 100; i++)
  79     a[i] = GANG_ID (i);
  80
  81   test_nonstatic (a, 10);
  82
  83   /* Static arguments with a variable expression.  */
  84
  85   x = 20;
  86 #pragma acc parallel loop gang (static:0+x) num_gangs (10)
  87   for (i = 0; i < 100; i++)
  88     a[i] = GANG_ID (i);
  89
  90   test_static (a, 10, 20);
  91
  92   x = 20;
  93 #pragma acc parallel loop gang (static:x) num_gangs (10)
  94   for (i = 0; i < 100; i++)
  95     a[i] = GANG_ID (i);
  96
  97   test_static (a, 10, 20);
  98
  99   return 0;
 100 }