libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c

   1 /* Tests of reduction on loop directive.  */
   2
   3 #include <assert.h>
   4
   5
   6 /* Test of reduction on loop directive (gangs, non-private reduction
   7    variable).  */
   8
   9 void g_np_1()
  10 {
  11   int i, arr[1024], res = 0, hres = 0;
  12
  13   for (i = 0; i < 1024; i++)
  14     arr[i] = i;
  15
  16   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
  17   {
  18     #pragma acc loop gang reduction(+:res)
  19     for (i = 0; i < 1024; i++)
  20       res += arr[i];
  21   }
  22
  23   for (i = 0; i < 1024; i++)
  24     hres += arr[i];
  25
  26   assert (res == hres);
  27
  28   res = hres = 1;
  29
  30   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
  31   {
  32     #pragma acc loop gang reduction(*:res)
  33     for (i = 0; i < 12; i++)
  34       res *= arr[i];
  35   }
  36
  37   for (i = 0; i < 12; i++)
  38     hres *= arr[i];
  39
  40   assert (res == hres);
  41 }
  42
  43
  44 /* Test of reduction on loop directive (gangs and vectors, non-private
  45    reduction variable).  */
  46
  47 void gv_np_1()
  48 {
  49   int i, arr[1024], res = 0, hres = 0;
  50
  51   for (i = 0; i < 1024; i++)
  52     arr[i] = i;
  53
  54   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
  55   {
  56     #pragma acc loop gang vector reduction(+:res)
  57     for (i = 0; i < 1024; i++)
  58       res += arr[i];
  59   }
  60
  61   for (i = 0; i < 1024; i++)
  62     hres += arr[i];
  63
  64   assert (res == hres);
  65 }
  66
  67
  68 /* Test of reduction on loop directive (gangs and workers, non-private
  69    reduction variable).  */
  70
  71 void gw_np_1()
  72 {
  73   int i, arr[1024], res = 0, hres = 0;
  74
  75   for (i = 0; i < 1024; i++)
  76     arr[i] = i;
  77
  78   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
  79   {
  80     #pragma acc loop gang worker reduction(+:res)
  81     for (i = 0; i < 1024; i++)
  82       res += arr[i];
  83   }
  84
  85   for (i = 0; i < 1024; i++)
  86     hres += arr[i];
  87
  88   assert (res == hres);
  89 }
  90
  91
  92 /* Test of reduction on loop directive (gangs, workers and vectors, non-private
  93    reduction variable).  */
  94
  95 void gwv_np_1()
  96 {
  97   int i, arr[1024], res = 0, hres = 0;
  98
  99   for (i = 0; i < 1024; i++)
 100     arr[i] = i;
 101
 102   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
 103   {
 104     #pragma acc loop gang worker vector reduction(+:res)
 105     for (i = 0; i < 1024; i++)
 106       res += arr[i];
 107   }
 108
 109   for (i = 0; i < 1024; i++)
 110     hres += arr[i];
 111
 112   assert (res == hres);
 113 }
 114
 115
 116 /* Test of reduction on loop directive (gangs, workers and vectors, non-private
 117    reduction variable: separate gang and worker/vector loops).  */
 118
 119 void gwv_np_2()
 120 {
 121   int i, j, arr[32768], res = 0, hres = 0;
 122
 123   for (i = 0; i < 32768; i++)
 124     arr[i] = i;
 125
 126   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
 127   {
 128     #pragma acc loop gang reduction(+:res)
 129     for (j = 0; j < 32; j++)
 130       {
 131         #pragma acc loop worker vector reduction(+:res)
 132         for (i = 0; i < 1024; i++)
 133           res += arr[j * 1024 + i];
 134       }
 135     /* "res" is non-private, and is not available until after the parallel
 136        region.  */
 137   }
 138
 139   for (i = 0; i < 32768; i++)
 140     hres += arr[i];
 141
 142   assert (res == hres);
 143 }
 144
 145
 146 /* Test of reduction on loop directive (gangs, workers and vectors, non-private
 147    reduction variable: separate gang and worker/vector loops).  */
 148
 149 void gwv_np_3()
 150 {
 151   int i, j;
 152   double arr[32768], res = 0, hres = 0;
 153
 154   for (i = 0; i < 32768; i++)
 155     arr[i] = i;
 156
 157   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 158                        copyin(arr)
 159   {
 160     #pragma acc loop gang reduction(+:res)
 161     for (j = 0; j < 32; j++)
 162       {
 163         #pragma acc loop worker vector reduction(+:res)
 164         for (i = 0; i < 1024; i++)
 165           res += arr[j * 1024 + i];
 166       }
 167   }
 168
 169   for (i = 0; i < 32768; i++)
 170     hres += arr[i];
 171
 172   assert (res == hres);
 173 }
 174
 175
 176 /* Test of reduction on loop directive (gangs, workers and vectors, multiple
 177    non-private reduction variables, float type).  */
 178
 179 void gwv_np_4()
 180 {
 181   int i, j;
 182   float arr[32768];
 183   float res = 0, mres = 0, hres = 0, hmres = 0;
 184
 185   for (i = 0; i < 32768; i++)
 186     arr[i] = i;
 187
 188   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
 189   {
 190     #pragma acc loop gang reduction(+:res) reduction(max:mres)
 191     for (j = 0; j < 32; j++)
 192       {
 193         #pragma acc loop worker vector reduction(+:res) reduction(max:mres)
 194         for (i = 0; i < 1024; i++)
 195           {
 196             res += arr[j * 1024 + i];
 197             if (arr[j * 1024 + i] > mres)
 198               mres = arr[j * 1024 + i];
 199           }
 200
 201         #pragma acc loop worker vector reduction(+:res) reduction(max:mres)
 202         for (i = 0; i < 1024; i++)
 203           {
 204             res += arr[j * 1024 + (1023 - i)];
 205             if (arr[j * 1024 + (1023 - i)] > mres)
 206               mres = arr[j * 1024 + (1023 - i)];
 207           }
 208       }
 209   }
 210
 211   for (j = 0; j < 32; j++)
 212     for (i = 0; i < 1024; i++)
 213       {
 214         hres += arr[j * 1024 + i];
 215         hres += arr[j * 1024 + (1023 - i)];
 216         if (arr[j * 1024 + i] > hmres)
 217           hmres = arr[j * 1024 + i];
 218         if (arr[j * 1024 + (1023 - i)] > hmres)
 219           hmres = arr[j * 1024 + (1023 - i)];
 220       }
 221
 222   assert (res == hres);
 223   assert (mres == hmres);
 224 }
 225
 226
 227 /* Test of reduction on loop directive (vectors, private reduction
 228    variable).  */
 229
 230 void v_p_1()
 231 {
 232   int i, j, arr[1024], out[32], res = 0, hres = 0;
 233
 234   for (i = 0; i < 1024; i++)
 235     arr[i] = i;
 236
 237   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 238                        private(res) copyout(out)
 239   {
 240     #pragma acc loop gang
 241     for (j = 0; j < 32; j++)
 242       {
 243         res = 0;
 244
 245         #pragma acc loop vector reduction(+:res)
 246         for (i = 0; i < 32; i++)
 247           res += arr[j * 32 + i];
 248
 249         out[j] = res;
 250       }
 251   }
 252
 253   for (j = 0; j < 32; j++)
 254     {
 255       hres = 0;
 256
 257       for (i = 0; i < 32; i++)
 258         hres += arr[j * 32 + i];
 259
 260       assert (out[j] == hres);
 261     }
 262 }
 263
 264
 265 /* Test of reduction on loop directive (vector reduction in
 266    gang-partitioned/worker-partitioned mode, private reduction variable).  */
 267
 268 void v_p_2()
 269 {
 270   int i, j, k;
 271   double ina[1024], inb[1024], out[1024], acc;
 272
 273   for (j = 0; j < 32; j++)
 274     for (i = 0; i < 32; i++)
 275       {
 276         ina[j * 32 + i] = (i == j) ? 2.0 : 0.0;
 277         inb[j * 32 + i] = (double) (i + j);
 278       }
 279
 280   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 281                        private(acc) copyin(ina, inb) copyout(out)
 282   {
 283     #pragma acc loop gang worker
 284     for (k = 0; k < 32; k++)
 285       for (j = 0; j < 32; j++)
 286         {
 287           acc = 0;
 288
 289           #pragma acc loop vector reduction(+:acc)
 290           for (i = 0; i < 32; i++)
 291             acc += ina[k * 32 + i] * inb[i * 32 + j];
 292
 293           out[k * 32 + j] = acc;
 294         }
 295   }
 296
 297   for (j = 0; j < 32; j++)
 298     for (i = 0; i < 32; i++)
 299       assert (out[j * 32 + i] == (i + j) * 2);
 300 }
 301
 302
 303 /* Test of reduction on loop directive (workers, private reduction
 304    variable).  */
 305
 306 void w_p_1()
 307 {
 308   int i, j, arr[1024], out[32], res = 0, hres = 0;
 309
 310   for (i = 0; i < 1024; i++)
 311     arr[i] = i;
 312
 313   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 314                        private(res) copyout(out)
 315   {
 316     #pragma acc loop gang
 317     for (j = 0; j < 32; j++)
 318       {
 319         res = 0;
 320
 321         #pragma acc loop worker reduction(+:res)
 322         for (i = 0; i < 32; i++)
 323           res += arr[j * 32 + i];
 324
 325         out[j] = res;
 326       }
 327   }
 328
 329   for (j = 0; j < 32; j++)
 330     {
 331       hres = 0;
 332
 333       for (i = 0; i < 32; i++)
 334         hres += arr[j * 32 + i];
 335
 336       assert (out[j] == hres);
 337     }
 338 }
 339
 340
 341 /* Test of reduction on loop directive (workers and vectors, private reduction
 342    variable).  */
 343
 344 void wv_p_1()
 345 {
 346   int i, j, arr[1024], out[32], res = 0, hres = 0;
 347
 348   for (i = 0; i < 1024; i++)
 349     arr[i] = i;
 350
 351   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 352                        private(res) copyout(out)
 353   {
 354     #pragma acc loop gang
 355     for (j = 0; j < 32; j++)
 356       {
 357         res = 0;
 358
 359         #pragma acc loop worker vector reduction(+:res)
 360         for (i = 0; i < 32; i++)
 361           res += arr[j * 32 + i];
 362
 363         out[j] = res;
 364       }
 365   }
 366
 367   for (j = 0; j < 32; j++)
 368     {
 369       hres = 0;
 370
 371       for (i = 0; i < 32; i++)
 372         hres += arr[j * 32 + i];
 373
 374       assert (out[j] == hres);
 375     }
 376 }
 377
 378
 379 /* Test of reduction on loop directive (workers and vectors, private reduction
 380    variable).  */
 381
 382 void wv_p_2()
 383 {
 384   int i, j, arr[32768], out[32], res = 0, hres = 0;
 385
 386   for (i = 0; i < 32768; i++)
 387     arr[i] = i;
 388
 389   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 390                        private(res) copyout(out)
 391   {
 392     #pragma acc loop gang
 393     for (j = 0; j < 32; j++)
 394       {
 395         res = j;
 396
 397         #pragma acc loop worker reduction(+:res)
 398         for (i = 0; i < 1024; i++)
 399           res += arr[j * 1024 + i];
 400
 401         #pragma acc loop vector reduction(+:res)
 402         for (i = 1023; i >= 0; i--)
 403           res += arr[j * 1024 + i];
 404
 405         out[j] = res;
 406       }
 407   }
 408
 409   for (j = 0; j < 32; j++)
 410     {
 411       hres = j;
 412
 413       for (i = 0; i < 1024; i++)
 414         hres += arr[j * 1024 + i] * 2;
 415
 416       assert (out[j] == hres);
 417     }
 418 }
 419
 420
 421 /* Test of reduction on loop directive (workers and vectors, private reduction
 422    variable: gang-redundant mode).  */
 423
 424 void wv_p_3()
 425 {
 426   int i, arr[1024], out[32], res = 0, hres = 0;
 427
 428   for (i = 0; i < 1024; i++)
 429     arr[i] = i ^ 33;
 430
 431   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 432                        private(res) copyin(arr) copyout(out)
 433   {
 434     /* Private variables aren't initialized by default in openacc.  */
 435     res = 0;
 436
 437     /* "res" should be available at the end of the following loop (and should
 438        have the same value redundantly in each gang).  */
 439     #pragma acc loop worker vector reduction(+:res)
 440     for (i = 0; i < 1024; i++)
 441       res += arr[i];
 442
 443     #pragma acc loop gang (static: 1)
 444     for (i = 0; i < 32; i++)
 445       out[i] = res;
 446   }
 447
 448   for (i = 0; i < 1024; i++)
 449     hres += arr[i];
 450
 451   for (i = 0; i < 32; i++)
 452     assert (out[i] == hres);
 453 }
 454
 455
 456 int main()
 457 {
 458   g_np_1();
 459   gv_np_1();
 460   gw_np_1();
 461   gwv_np_1();
 462   gwv_np_2();
 463   gwv_np_3();
 464   gwv_np_4();
 465   v_p_1();
 466   v_p_2();
 467   w_p_1();
 468   wv_p_1();
 469   wv_p_2();
 470   wv_p_3();
 471
 472   return 0;
 473 }