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 Doerferttemplate <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 Doerferttemplate <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 Doerfertint 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