reduction-5.c 4.6 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193
  1. /* { dg-additional-options "-foffload-options=nvptx-none=-latomic" { target { offload_target_nvptx } } } */
  2. /* C / C++'s logical AND and OR operators take any scalar argument
  3. which compares (un)equal to 0 - the result 1 or 0 and of type int.
  4. In this testcase, the int result is again converted to a floating-poing
  5. or complex type.
  6. While having a floating-point/complex array element with || and && can make
  7. sense, having a non-integer/non-bool reduction variable is odd but valid.
  8. Test: FP reduction variable + FP array - as reduction-1.c but with target */
  9. #define N 1024
  10. _Complex float rcf[N];
  11. _Complex double rcd[N];
  12. float rf[N];
  13. double rd[N];
  14. int
  15. reduction_or ()
  16. {
  17. float orf = 0;
  18. double ord = 0;
  19. _Complex float orfc = 0;
  20. _Complex double ordc = 0;
  21. #pragma omp target parallel reduction(||: orf) map(orf)
  22. for (int i=0; i < N; ++i)
  23. orf = orf || rf[i];
  24. #pragma omp target parallel for reduction(||: ord) map(ord)
  25. for (int i=0; i < N; ++i)
  26. ord = ord || rcd[i];
  27. #pragma omp target parallel for simd reduction(||: orfc) map(orfc)
  28. for (int i=0; i < N; ++i)
  29. orfc = orfc || rcf[i];
  30. #pragma omp target parallel loop reduction(||: ordc) map(ordc)
  31. for (int i=0; i < N; ++i)
  32. ordc = ordc || rcd[i];
  33. return orf + ord + __real__ orfc + __real__ ordc;
  34. }
  35. int
  36. reduction_or_teams ()
  37. {
  38. float orf = 0;
  39. double ord = 0;
  40. _Complex float orfc = 0;
  41. _Complex double ordc = 0;
  42. #pragma omp target teams distribute parallel for reduction(||: orf) map(orf)
  43. for (int i=0; i < N; ++i)
  44. orf = orf || rf[i];
  45. #pragma omp target teams distribute parallel for simd reduction(||: ord) map(ord)
  46. for (int i=0; i < N; ++i)
  47. ord = ord || rcd[i];
  48. #pragma omp target teams distribute parallel for reduction(||: orfc) map(orfc)
  49. for (int i=0; i < N; ++i)
  50. orfc = orfc || rcf[i];
  51. #pragma omp target teams distribute parallel for simd reduction(||: ordc) map(ordc)
  52. for (int i=0; i < N; ++i)
  53. ordc = ordc || rcd[i];
  54. return orf + ord + __real__ orfc + __real__ ordc;
  55. }
  56. int
  57. reduction_and ()
  58. {
  59. float andf = 1;
  60. double andd = 1;
  61. _Complex float andfc = 1;
  62. _Complex double anddc = 1;
  63. #pragma omp target parallel reduction(&&: andf) map(andf)
  64. for (int i=0; i < N; ++i)
  65. andf = andf && rf[i];
  66. #pragma omp target parallel for reduction(&&: andd) map(andd)
  67. for (int i=0; i < N; ++i)
  68. andd = andd && rcd[i];
  69. #pragma omp target parallel for simd reduction(&&: andfc) map(andfc)
  70. for (int i=0; i < N; ++i)
  71. andfc = andfc && rcf[i];
  72. #pragma omp target parallel loop reduction(&&: anddc) map(anddc)
  73. for (int i=0; i < N; ++i)
  74. anddc = anddc && rcd[i];
  75. return andf + andd + __real__ andfc + __real__ anddc;
  76. }
  77. int
  78. reduction_and_teams ()
  79. {
  80. float andf = 1;
  81. double andd = 1;
  82. _Complex float andfc = 1;
  83. _Complex double anddc = 1;
  84. #pragma omp target teams distribute parallel for reduction(&&: andf) map(andf)
  85. for (int i=0; i < N; ++i)
  86. andf = andf && rf[i];
  87. #pragma omp target teams distribute parallel for simd reduction(&&: andd) map(andd)
  88. for (int i=0; i < N; ++i)
  89. andd = andd && rcd[i];
  90. #pragma omp target teams distribute parallel for reduction(&&: andfc) map(andfc)
  91. for (int i=0; i < N; ++i)
  92. andfc = andfc && rcf[i];
  93. #pragma omp target teams distribute parallel for simd reduction(&&: anddc) map(anddc)
  94. for (int i=0; i < N; ++i)
  95. anddc = anddc && rcd[i];
  96. return andf + andd + __real__ andfc + __real__ anddc;
  97. }
  98. int
  99. main ()
  100. {
  101. for (int i = 0; i < N; ++i)
  102. {
  103. rf[i] = 0;
  104. rd[i] = 0;
  105. rcf[i] = 0;
  106. rcd[i] = 0;
  107. }
  108. if (reduction_or () != 0)
  109. __builtin_abort ();
  110. if (reduction_or_teams () != 0)
  111. __builtin_abort ();
  112. if (reduction_and () != 0)
  113. __builtin_abort ();
  114. if (reduction_and_teams () != 0)
  115. __builtin_abort ();
  116. rf[10] = 1.0;
  117. rd[15] = 1.0;
  118. rcf[10] = 1.0;
  119. rcd[15] = 1.0i;
  120. if (reduction_or () != 4)
  121. __builtin_abort ();
  122. if (reduction_or_teams () != 4)
  123. __builtin_abort ();
  124. if (reduction_and () != 0)
  125. __builtin_abort ();
  126. if (reduction_and_teams () != 0)
  127. __builtin_abort ();
  128. for (int i = 0; i < N; ++i)
  129. {
  130. rf[i] = 1;
  131. rd[i] = 1;
  132. rcf[i] = 1;
  133. rcd[i] = 1;
  134. }
  135. if (reduction_or () != 4)
  136. __builtin_abort ();
  137. if (reduction_or_teams () != 4)
  138. __builtin_abort ();
  139. if (reduction_and () != 4)
  140. __builtin_abort ();
  141. if (reduction_and_teams () != 4)
  142. __builtin_abort ();
  143. rf[10] = 0.0;
  144. rd[15] = 0.0;
  145. rcf[10] = 0.0;
  146. rcd[15] = 0.0;
  147. if (reduction_or () != 4)
  148. __builtin_abort ();
  149. if (reduction_or_teams () != 4)
  150. __builtin_abort ();
  151. if (reduction_and () != 0)
  152. __builtin_abort ();
  153. if (reduction_and_teams () != 0)
  154. __builtin_abort ();
  155. return 0;
  156. }