Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[ESIMD] E2E test for writeable simd_view subscript operator #415

Merged
merged 1 commit into from
Sep 8, 2021
Merged
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
100 changes: 100 additions & 0 deletions SYCL/ESIMD/api/simd_view_subscript_operator.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,100 @@
//==----- simd_view_subscript_operator.cpp - DPC++ ESIMD on-device test ---==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// REQUIRES: gpu
// UNSUPPORTED: cuda || rocm
// RUN: %clangxx -fsycl %s -o %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
//
// The test checks that it's possible to write through the simd_view subscript
// operator. E.g.:
// simd<int, 4> v = 1;
// auto v1 = v.select<2, 1>(0);
// v1[0] = 0; // v1[0] returns writable simd_view

#include "../esimd_test_utils.hpp"

#include <CL/sycl.hpp>
#include <sycl/ext/intel/experimental/esimd.hpp>

#include <iostream>

using namespace cl::sycl;
using namespace sycl::ext::intel::experimental::esimd;

int main(int argc, char **argv) {
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());

auto dev = q.get_device();
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

constexpr unsigned VL = 16;

int A[VL];
int B[VL];
int gold[VL];

for (unsigned i = 0; i < VL; ++i) {
A[i] = -i;
B[i] = i;
gold[i] = B[i];
}

// some random indices to overwrite elements in B with elements from A.
std::array<int, 5> indicesToCopy = {2, 5, 9, 10, 13};

try {
buffer<int, 1> bufA(A, range<1>(VL));
buffer<int, 1> bufB(B, range<1>(VL));
range<1> glob_range{1};

auto e = q.submit([&](handler &cgh) {
auto PA = bufA.get_access<access::mode::read>(cgh);
auto PB = bufB.template get_access<access::mode::read_write>(cgh);
cgh.parallel_for<class Test>(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
using namespace sycl::ext::intel::experimental::esimd;
simd<int, VL> va;
va.copy_from(PA, 0);
simd<int, VL> vb;
vb.copy_from(PB, 0);
auto view_va = va.select<VL, 1>(0);
auto view_vb = vb.select<VL, 1>(0);
for (auto idx : indicesToCopy)
view_vb[idx] = view_va[idx];
vb.copy_to(PB, 0);
});
});
q.wait_and_throw();
} catch (cl::sycl::exception const &e) {
std::cout << "SYCL exception caught: " << e.what() << '\n';
return e.get_cl_code();
}

int err_cnt = 0;

for (auto i : indicesToCopy)
gold[i] = A[i];

for (unsigned i = 0; i < VL; ++i) {
int val = B[i];

if (val != gold[i]) {
if (++err_cnt < 10) {
std::cout << "failed at index " << i << ": " << val << " != " << gold[i]
<< " (gold)\n";
}
}
}
if (err_cnt > 0) {
std::cout << " pass rate: " << ((float)(VL - err_cnt) / (float)VL) * 100.0f
<< "% (" << (VL - err_cnt) << "/" << VL << ")\n";
}

std::cout << (err_cnt > 0 ? " FAILED\n" : " Passed\n");

return err_cnt > 0 ? 1 : 0;
}