xref: /llvm-project/offload/test/offloading/parallel_target_teams_reduction_min.cpp (revision 8823448807f3b1a1362d1417e062d763734e02f5)
1330d8983SJohannes Doerfert // RUN: %libomptarget-compilexx-and-run-generic
2330d8983SJohannes Doerfert // RUN: %libomptarget-compileoptxx-and-run-generic
3330d8983SJohannes Doerfert 
4330d8983SJohannes Doerfert // FIXME: This is a bug in host offload, this should run fine.
5*88234488SEthan Luis McDonough // REQUIRES: gpu
6330d8983SJohannes Doerfert 
7330d8983SJohannes Doerfert // This test validates that the OpenMP target reductions to find a minimum work
8330d8983SJohannes Doerfert // as indended for a few common data types.
9330d8983SJohannes Doerfert 
10330d8983SJohannes Doerfert #include <algorithm>
11330d8983SJohannes Doerfert #include <cassert>
12330d8983SJohannes Doerfert #include <limits>
13330d8983SJohannes Doerfert #include <vector>
14330d8983SJohannes Doerfert 
test_min_idx_reduction()15330d8983SJohannes Doerfert template <class Tp> void test_min_idx_reduction() {
16330d8983SJohannes Doerfert   const Tp length = 1000;
17330d8983SJohannes Doerfert   const Tp nminimas = 8;
18330d8983SJohannes Doerfert   std::vector<float> a(length, 3.0f);
19330d8983SJohannes Doerfert   const Tp step = length / nminimas;
20330d8983SJohannes Doerfert   for (Tp i = 0; i < nminimas; i++) {
21330d8983SJohannes Doerfert     a[i * step] -= 1.0f;
22330d8983SJohannes Doerfert   }
23330d8983SJohannes Doerfert   for (Tp i = 0; i < nminimas; i++) {
24330d8983SJohannes Doerfert     Tp idx = a.size();
25330d8983SJohannes Doerfert     float *b = a.data();
26330d8983SJohannes Doerfert #pragma omp target teams distribute parallel for reduction(min : idx)          \
27330d8983SJohannes Doerfert     map(always, to : b[0 : length])
28330d8983SJohannes Doerfert     for (Tp j = 0; j < length - 1; j++) {
29330d8983SJohannes Doerfert       if (b[j] < b[j + 1]) {
30330d8983SJohannes Doerfert         idx = std::min(idx, j);
31330d8983SJohannes Doerfert       }
32330d8983SJohannes Doerfert     }
33330d8983SJohannes Doerfert     assert(idx == i * step &&
34330d8983SJohannes Doerfert            "#pragma omp target teams distribute parallel for "
35330d8983SJohannes Doerfert            "reduction(min:<identifier list>) does not work as intended.");
36330d8983SJohannes Doerfert     a[idx] += 1.0f;
37330d8983SJohannes Doerfert   }
38330d8983SJohannes Doerfert }
39330d8983SJohannes Doerfert 
test_min_val_reduction()40330d8983SJohannes Doerfert template <class Tp> void test_min_val_reduction() {
41330d8983SJohannes Doerfert   const int length = 1000;
42330d8983SJohannes Doerfert   const int half = length / 2;
43330d8983SJohannes Doerfert   std::vector<Tp> a(length, (Tp)3);
44330d8983SJohannes Doerfert   a[half] -= (Tp)1;
45330d8983SJohannes Doerfert   Tp min_val = std::numeric_limits<Tp>::max();
46330d8983SJohannes Doerfert   Tp *b = a.data();
47330d8983SJohannes Doerfert #pragma omp target teams distribute parallel for reduction(min : min_val)      \
48330d8983SJohannes Doerfert     map(always, to : b[0 : length])
49330d8983SJohannes Doerfert   for (int i = 0; i < length; i++) {
50330d8983SJohannes Doerfert     min_val = std::min(min_val, b[i]);
51330d8983SJohannes Doerfert   }
52330d8983SJohannes Doerfert   assert(std::abs(((double)a[half + 1]) - ((double)min_val) - 1.0) < 1e-6 &&
53330d8983SJohannes Doerfert          "#pragma omp target teams distribute parallel for "
54330d8983SJohannes Doerfert          "reduction(min:<identifier list>) does not work as intended.");
55330d8983SJohannes Doerfert }
56330d8983SJohannes Doerfert 
main()57330d8983SJohannes Doerfert int main() {
58330d8983SJohannes Doerfert   // Reducing over indices
59330d8983SJohannes Doerfert   test_min_idx_reduction<int>();
60330d8983SJohannes Doerfert   test_min_idx_reduction<unsigned int>();
61330d8983SJohannes Doerfert   test_min_idx_reduction<long>();
62330d8983SJohannes Doerfert 
63330d8983SJohannes Doerfert   // Reducing over values
64330d8983SJohannes Doerfert   test_min_val_reduction<int>();
65330d8983SJohannes Doerfert   test_min_val_reduction<unsigned int>();
66330d8983SJohannes Doerfert   test_min_val_reduction<long>();
67330d8983SJohannes Doerfert   test_min_val_reduction<float>();
68330d8983SJohannes Doerfert   test_min_val_reduction<double>();
69330d8983SJohannes Doerfert   return 0;
70330d8983SJohannes Doerfert }
71