78ff5b55fd
Since we can now build the DeviceRTL with SPIR-V, redo the `XFAIL/UNSUPPORTED` specifications for the tests we see passing/failing on the Level Zero backend with the DeviceRTL being used. The tests marked `UNSUPPORTED` hang or sporadically fail and those are tracked in https://github.com/llvm/llvm-project/issues/182119. This change will allow us to enable CI testing with the DeviceRTL. Here are the full test results with this change applied, running only the `spirv64-intel` `check-offload` tests: ``` Total Discovered Tests: 453 Unsupported : 206 (45.47%) Passed : 141 (31.13%) Expectedly Failed: 106 (23.40%) ``` 31% is not a bad start. --------- Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
73 lines
2.2 KiB
C++
73 lines
2.2 KiB
C++
// RUN: %libomptarget-compilexx-and-run-generic
|
|
// RUN: %libomptarget-compileoptxx-and-run-generic
|
|
|
|
// FIXME: This is a bug in host offload, this should run fine.
|
|
// REQUIRES: gpu
|
|
// https://github.com/llvm/llvm-project/issues/182119
|
|
// UNSUPPORTED: intelgpu
|
|
|
|
// This test validates that the OpenMP target reductions to find a maximum work
|
|
// as intended for a few common data types.
|
|
|
|
#include <algorithm>
|
|
#include <cassert>
|
|
#include <limits>
|
|
#include <vector>
|
|
|
|
template <class Tp> void test_max_idx_reduction() {
|
|
const Tp length = 1000;
|
|
const Tp nmaximas = 8;
|
|
std::vector<float> a(length, 3.0f);
|
|
const Tp step = length / nmaximas;
|
|
for (Tp i = 0; i < nmaximas; i++) {
|
|
a[i * step] += 1.0f;
|
|
}
|
|
for (Tp i = nmaximas; i > 0; i--) {
|
|
Tp idx = 0;
|
|
float *b = a.data();
|
|
#pragma omp target teams distribute parallel for reduction(max : idx) \
|
|
map(always, to : b[0 : length])
|
|
for (Tp j = 0; j < length - 1; j++) {
|
|
if (b[j] > b[j + 1]) {
|
|
idx = std::max(idx, j);
|
|
}
|
|
}
|
|
assert(idx == (i - 1) * step &&
|
|
"#pragma omp target teams distribute parallel for "
|
|
"reduction(max:<identifier list>) does not work as intended.");
|
|
a[idx] -= 1.0f;
|
|
}
|
|
}
|
|
|
|
template <class Tp> void test_max_val_reduction() {
|
|
const int length = 1000;
|
|
const int half = length / 2;
|
|
std::vector<Tp> a(length, (Tp)3);
|
|
a[half] += (Tp)1;
|
|
Tp max_val = std::numeric_limits<Tp>::lowest();
|
|
Tp *b = a.data();
|
|
#pragma omp target teams distribute parallel for reduction(max : max_val) \
|
|
map(always, to : b[0 : length])
|
|
for (int i = 0; i < length; i++) {
|
|
max_val = std::max(max_val, b[i]);
|
|
}
|
|
assert(std::abs(((double)a[half + 1]) - ((double)max_val) + 1.0) < 1e-6 &&
|
|
"#pragma omp target teams distribute parallel for "
|
|
"reduction(max:<identifier list>) does not work as intended.");
|
|
}
|
|
|
|
int main() {
|
|
// Reducing over indices
|
|
test_max_idx_reduction<int>();
|
|
test_max_idx_reduction<unsigned int>();
|
|
test_max_idx_reduction<long>();
|
|
|
|
// Reducing over values
|
|
test_max_val_reduction<int>();
|
|
test_max_val_reduction<unsigned int>();
|
|
test_max_val_reduction<long>();
|
|
test_max_val_reduction<float>();
|
|
test_max_val_reduction<double>();
|
|
return 0;
|
|
}
|