libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-many.c

   1 /* To avoid 'error: shared-memory region overflow':
   2    { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=64" { target openacc_radeon_accel_selected } }
   3 */
   4
   5 #include <assert.h>
   6 #include <stdio.h>
   7
   8 #define LOCAL(n) double n = input;
   9 #define LOCALS(n) LOCAL(n##1) LOCAL(n##2) LOCAL(n##3) LOCAL(n##4) \
  10                   LOCAL(n##5) LOCAL(n##6) LOCAL(n##7) LOCAL(n##8)
  11 #define LOCALS2(n) LOCALS(n##a) LOCALS(n##b) LOCALS(n##c) LOCALS(n##d) \
  12                    LOCALS(n##e) LOCALS(n##f) LOCALS(n##g) LOCALS(n##h)
  13
  14 #define USE(n) n
  15 #define USES(n,OP) USE(n##1) OP USE(n##2) OP USE(n##3) OP USE (n##4) OP \
  16                    USE(n##5) OP USE(n##6) OP USE(n##7) OP USE (n##8)
  17 #define USES2(n,OP) USES(n##a,OP) OP USES(n##b,OP) OP USES(n##c,OP) OP \
  18                     USES(n##d,OP) OP USES(n##e,OP) OP USES(n##f,OP) OP \
  19                     USES(n##g,OP) OP USES(n##h,OP)
  20
  21 int main (void)
  22 {
  23   int ret;
  24   int input = 1;
  25
  26   #pragma acc parallel num_gangs(1) num_workers(32) copyout(ret)
  27   {
  28     int w = 0;
  29     LOCALS2(h);
  30
  31     #pragma acc loop worker reduction(+:w)
  32     for (int i = 0; i < 32; i++)
  33       {
  34         int u = USES2(h,+);
  35         w += u;
  36       }
  37
  38     printf ("w=%d\n", w);
  39     /* { dg-output "w=2048(\n|\r\n|\r)" } */
  40
  41     LOCALS2(i);
  42
  43     #pragma acc loop worker reduction(+:w)
  44     for (int i = 0; i < 32; i++)
  45       {
  46         int u = USES2(i,+);
  47         w += u;
  48       }
  49
  50     printf ("w=%d\n", w);
  51     /* { dg-output "w=4096(\n|\r\n|\r)" } */
  52
  53     LOCALS2(j);
  54     LOCALS2(k);
  55
  56     #pragma acc loop worker reduction(+:w)
  57     for (int i = 0; i < 32; i++)
  58       {
  59         int u = USES2(j,+);
  60         w += u;
  61       }
  62
  63     printf ("w=%d\n", w);
  64     /* { dg-output "w=6144(\n|\r\n|\r)" } */
  65
  66     #pragma acc loop worker reduction(+:w)
  67     for (int i = 0; i < 32; i++)
  68       {
  69         int u = USES2(k,+);
  70         w += u;
  71       }
  72
  73     ret = (w == 64 * 32 * 4);
  74     printf ("w=%d\n", w);
  75     /* { dg-output "w=8192(\n|\r\n|\r)" } */
  76   }
  77
  78   assert (ret);
  79
  80   return 0;
  81 }