|
| 1 | +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-unnamed-lambda %s -o %t.out |
| 2 | +// RUN: %GPU_RUN_PLACEHOLDER %t.out |
| 3 | +// REQUIRES: gpu |
| 4 | +#include <iostream> |
| 5 | +#include <sycl.hpp> |
| 6 | +#include <vector> |
| 7 | + |
| 8 | +using namespace sycl; |
| 9 | + |
| 10 | +template <typename T> |
| 11 | +using shared_allocator = sycl::usm_allocator<T, sycl::usm::alloc::shared>; |
| 12 | + |
| 13 | +template <typename T> using shared_vector = std::vector<T, shared_allocator<T>>; |
| 14 | + |
| 15 | +template <typename T> bool are_bitwise_equal(T lhs, T rhs) { |
| 16 | + constexpr size_t size{sizeof(T)}; |
| 17 | + |
| 18 | + // Such type-punning is OK from the point of strict aliasing rules |
| 19 | + const auto &lhs_bytes = reinterpret_cast<const unsigned char(&)[size]>(lhs); |
| 20 | + const auto &rhs_bytes = reinterpret_cast<const unsigned char(&)[size]>(rhs); |
| 21 | + |
| 22 | + bool result{true}; |
| 23 | + for (size_t i = 0; i < size; ++i) { |
| 24 | + result &= lhs_bytes[i] == rhs_bytes[i]; |
| 25 | + } |
| 26 | + return result; |
| 27 | +} |
| 28 | + |
| 29 | +template <typename T> bool test() { |
| 30 | + sycl::queue queue{}; |
| 31 | + |
| 32 | + constexpr int NumElems{32}; |
| 33 | + bool pass{true}; |
| 34 | + |
| 35 | + static const T inexact = static_cast<T>(0.1); |
| 36 | + |
| 37 | + shared_vector<T> result_source{NumElems, shared_allocator<T>{queue}}; |
| 38 | + shared_vector<T> input{NumElems, shared_allocator<T>{queue}}; |
| 39 | + |
| 40 | + for (size_t i = 0; i < NumElems; ++i) { |
| 41 | + input[i] = inexact * i; |
| 42 | + } |
| 43 | + |
| 44 | + queue.submit([&](sycl::handler &cgh) { |
| 45 | + auto out_source = result_source.data(); |
| 46 | + auto in = input.data(); |
| 47 | + |
| 48 | + cgh.single_task<>([=]() { |
| 49 | + for (size_t i = 0; i < NumElems; ++i) { |
| 50 | + auto source = in[i]; |
| 51 | + ++source; |
| 52 | + out_source[i] = source; |
| 53 | + } |
| 54 | + }); |
| 55 | + }); |
| 56 | + queue.wait_and_throw(); |
| 57 | + |
| 58 | + for (size_t i = 0; i < NumElems; ++i) { |
| 59 | + T expected_value = input[i] + 1; |
| 60 | + |
| 61 | + if (!are_bitwise_equal(expected_value, result_source[i])) { |
| 62 | + pass = false; |
| 63 | + std::cout << "Sample failed retrieved value: " << result_source[i] |
| 64 | + << ", but expected: " << expected_value << ", at index: " << i |
| 65 | + << std::endl; |
| 66 | + } |
| 67 | + } |
| 68 | + return pass; |
| 69 | +} |
| 70 | + |
| 71 | +int main(int argc, char **argv) { |
| 72 | + bool passed = true; |
| 73 | + passed &= test<float>(); |
| 74 | + passed &= test<sycl::half>(); |
| 75 | + return passed ? 0 : 1; |
| 76 | +} |
0 commit comments