target-in-reduction-2.c 4.2 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173
  1. struct S { int a, b, c[2]; };
  2. #pragma omp declare reduction (+: struct S : (omp_out.a += omp_in.a, omp_out.b += omp_in.b)) \
  3. initializer (omp_priv = { 0, 0, { 0, 0 } })
  4. void
  5. foo (struct S x, struct S *y, int n, int v)
  6. {
  7. struct S z[3] = { { 45, 47, {} }, { 46, 48, {} }, { 47, 49, {} } };
  8. struct S u[n], w[n];
  9. int i;
  10. for (i = 0; i < n; i++)
  11. {
  12. w[i].a = u[i].a = n + i;
  13. w[i].b = u[i].b = n - i;
  14. w[i].c[0] = u[i].c[0] = 0;
  15. w[i].c[1] = u[i].c[1] = 0;
  16. }
  17. #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v])
  18. {
  19. #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v])
  20. {
  21. x.a++;
  22. x.b++;
  23. y[0].a += 2;
  24. y[0].b += 12;
  25. y[1].a += 3;
  26. y[1].b += 13;
  27. z[1].a += 4;
  28. z[1].b += 14;
  29. u[0].a += 5;
  30. u[0].b += 15;
  31. w[1].a += 6;
  32. w[1].b += 16;
  33. }
  34. #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) map(tofrom: x.a, x.b, x.c[:2])
  35. {
  36. x.a += 4;
  37. x.b += 14;
  38. y[0].a += 5;
  39. y[0].b += 15;
  40. y[1].a += 6;
  41. y[1].b += 16;
  42. z[2].a += 7;
  43. z[2].b += 17;
  44. u[1].a += 8;
  45. u[1].b += 18;
  46. w[2].a += 7;
  47. w[2].b += 17;
  48. }
  49. #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2])
  50. {
  51. x.a += 9;
  52. x.b += 19;
  53. y[0].a += 10;
  54. y[0].b += 20;
  55. y[1].a += 11;
  56. y[1].b += 21;
  57. z[1].a += 12;
  58. z[1].b += 22;
  59. u[2].a += 13;
  60. u[2].b += 23;
  61. w[1].a += 14;
  62. w[1].b += 24;
  63. }
  64. }
  65. if (x.a != 56 || y[0].a != 60 || y[1].a != 64)
  66. __builtin_abort ();
  67. if (x.b != 86 || y[0].b != 100 || y[1].b != 104)
  68. __builtin_abort ();
  69. if (z[0].a != 45 || z[1].a != 62 || z[2].a != 54)
  70. __builtin_abort ();
  71. if (z[0].b != 47 || z[1].b != 84 || z[2].b != 66)
  72. __builtin_abort ();
  73. if (u[0].a != 8 || u[1].a != 12 || u[2].a != 18)
  74. __builtin_abort ();
  75. if (u[0].b != 18 || u[1].b != 20 || u[2].b != 24)
  76. __builtin_abort ();
  77. if (w[0].a != 3 || w[1].a != 24 || w[2].a != 12)
  78. __builtin_abort ();
  79. if (w[0].b != 3 || w[1].b != 42 || w[2].b != 18)
  80. __builtin_abort ();
  81. }
  82. void
  83. bar (struct S x, struct S *y, int n, int v)
  84. {
  85. struct S z[3] = { { 45, 47, {} }, { 46, 48, {} }, { 47, 49, {} } };
  86. struct S u[n], w[n];
  87. int i;
  88. for (i = 0; i < n; i++)
  89. {
  90. w[i].a = u[i].a = n + i;
  91. w[i].b = u[i].b = n - i;
  92. w[i].c[0] = u[i].c[0] = 0;
  93. w[i].c[1] = u[i].c[1] = 0;
  94. }
  95. #pragma omp parallel master
  96. #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v])
  97. {
  98. #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v])
  99. {
  100. x.a++;
  101. x.b++;
  102. y[0].a += 2;
  103. y[0].b += 12;
  104. y[1].a += 3;
  105. y[1].b += 13;
  106. z[1].a += 4;
  107. z[1].b += 14;
  108. u[0].a += 5;
  109. u[0].b += 15;
  110. w[1].a += 6;
  111. w[1].b += 16;
  112. }
  113. #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) map(tofrom: x.a, x.b, x.c[:2])
  114. {
  115. x.a += 4;
  116. x.b += 14;
  117. y[0].a += 5;
  118. y[0].b += 15;
  119. y[1].a += 6;
  120. y[1].b += 16;
  121. z[2].a += 7;
  122. z[2].b += 17;
  123. u[1].a += 8;
  124. u[1].b += 18;
  125. w[2].a += 7;
  126. w[2].b += 17;
  127. }
  128. #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2])
  129. {
  130. x.a += 9;
  131. x.b += 19;
  132. y[0].a += 10;
  133. y[0].b += 20;
  134. y[1].a += 11;
  135. y[1].b += 21;
  136. z[1].a += 12;
  137. z[1].b += 22;
  138. u[2].a += 13;
  139. u[2].b += 23;
  140. w[1].a += 14;
  141. w[1].b += 24;
  142. }
  143. }
  144. if (x.a != 56 || y[0].a != 77 || y[1].a != 84)
  145. __builtin_abort ();
  146. if (x.b != 86 || y[0].b != 147 || y[1].b != 154)
  147. __builtin_abort ();
  148. if (z[0].a != 45 || z[1].a != 62 || z[2].a != 54)
  149. __builtin_abort ();
  150. if (z[0].b != 47 || z[1].b != 84 || z[2].b != 66)
  151. __builtin_abort ();
  152. if (u[0].a != 8 || u[1].a != 12 || u[2].a != 18)
  153. __builtin_abort ();
  154. if (u[0].b != 18 || u[1].b != 20 || u[2].b != 24)
  155. __builtin_abort ();
  156. if (w[0].a != 3 || w[1].a != 24 || w[2].a != 12)
  157. __builtin_abort ();
  158. if (w[0].b != 3 || w[1].b != 42 || w[2].b != 18)
  159. __builtin_abort ();
  160. }
  161. int
  162. main ()
  163. {
  164. struct S x = { 42, 52 };
  165. struct S y[2] = { { 43, 53 }, { 44, 54 } };
  166. #pragma omp parallel master
  167. foo (x, y, 3, 2);
  168. bar (x, y, 3, 2);
  169. return 0;
  170. }