Skip to content

Commit b5de4bf

Browse files
authored
[ESIMD] Introduce a test for a fix for copy_to handling of char buffers (intel/llvm-test-suite#1057)
1 parent c17d0a8 commit b5de4bf

File tree

1 file changed

+179
-0
lines changed

1 file changed

+179
-0
lines changed
+179
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,179 @@
1+
// REQUIRES: gpu
2+
// UNSUPPORTED: cuda || hip
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
6+
//==- copyto_char_test.cpp - Test for using copy_to to copy char buffers -==//
7+
//
8+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
9+
// See https://llvm.org/LICENSE.txt for license information.
10+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
11+
//
12+
//===----------------------------------------------------------------------===//
13+
14+
#include <sycl/ext/intel/esimd.hpp>
15+
#include <sycl/sycl.hpp>
16+
17+
#include <iostream>
18+
#include <memory>
19+
20+
using namespace sycl::ext::intel::esimd;
21+
22+
template <typename DataT>
23+
using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
24+
template <typename DataT>
25+
using shared_vector = std::vector<DataT, shared_allocator<DataT>>;
26+
27+
using DataT = char;
28+
29+
template <int NumElems, bool IsAcc, int ResultOffset = 0> int test_to_copy() {
30+
sycl::queue queue;
31+
constexpr int NumSelectedElems = NumElems / 3;
32+
constexpr int Stride = 2;
33+
constexpr int Offset = 6;
34+
35+
shared_allocator<DataT> allocator(queue);
36+
shared_vector<DataT> result(NumElems + ResultOffset, allocator);
37+
shared_vector<DataT> initial_ref_data(NumElems, allocator);
38+
shared_vector<DataT> ref_data_for_fill(NumSelectedElems, allocator);
39+
40+
constexpr size_t value_for_increase_ref_data_for_fill = 50;
41+
for (size_t i = 0; i < NumElems; ++i) {
42+
initial_ref_data[i] = i + 1;
43+
}
44+
// We should have different values in the first reference data and in the
45+
// second reference data.
46+
for (size_t i = 0; i < NumSelectedElems; ++i) {
47+
ref_data_for_fill[i] =
48+
initial_ref_data[i] + value_for_increase_ref_data_for_fill;
49+
}
50+
if constexpr (IsAcc) {
51+
sycl::buffer<DataT> output_buf(result.data() + ResultOffset,
52+
result.size() - ResultOffset);
53+
queue.submit([&](sycl::handler &cgh) {
54+
DataT *init_ref_ptr = initial_ref_data.data();
55+
DataT *ref_data_for_fill_ptr = ref_data_for_fill.data();
56+
auto acc =
57+
output_buf.template get_access<sycl::access::mode::read_write>(cgh);
58+
59+
cgh.single_task([=]() SYCL_ESIMD_KERNEL {
60+
simd<DataT, NumElems> src_simd_obj;
61+
src_simd_obj.copy_from(init_ref_ptr);
62+
simd<DataT, NumSelectedElems> dst_simd_obj;
63+
dst_simd_obj.copy_from(ref_data_for_fill_ptr);
64+
src_simd_obj.template select<NumSelectedElems, Stride>(Offset) =
65+
dst_simd_obj;
66+
src_simd_obj.copy_to(acc, 0);
67+
});
68+
});
69+
} else {
70+
queue.submit([&](sycl::handler &cgh) {
71+
DataT *init_ref_ptr = initial_ref_data.data();
72+
DataT *ref_data_for_fill_ptr = ref_data_for_fill.data();
73+
DataT *const out_ptr = result.data() + ResultOffset;
74+
75+
cgh.single_task([=]() SYCL_ESIMD_KERNEL {
76+
simd<DataT, NumElems> src_simd_obj;
77+
src_simd_obj.copy_from(init_ref_ptr);
78+
simd<DataT, NumSelectedElems> dst_simd_obj;
79+
dst_simd_obj.copy_from(ref_data_for_fill_ptr);
80+
src_simd_obj.template select<NumSelectedElems, Stride>(Offset) =
81+
dst_simd_obj;
82+
src_simd_obj.copy_to(out_ptr);
83+
});
84+
});
85+
}
86+
queue.wait_and_throw();
87+
88+
std::vector<size_t> selected_indices;
89+
// Collect the indexess that has been selected.
90+
for (size_t i = Offset; i < Stride * NumSelectedElems + Offset; i += Stride) {
91+
selected_indices.push_back(i);
92+
}
93+
94+
// Push the largest value to avoid the following error: can't dereference
95+
// out of range vector iterator.
96+
selected_indices.push_back(std::numeric_limits<size_t>::max());
97+
auto selected_indices_ptr = selected_indices.begin();
98+
99+
// Verify that values, that do not was selected has initial values.
100+
for (size_t i = 0; i < NumElems; ++i) {
101+
// If current index is less than selected index verify that this element
102+
// hasn't been selected and changed.
103+
if (i < *selected_indices_ptr) {
104+
DataT expected = initial_ref_data[i];
105+
106+
DataT retrieved = result[i + ResultOffset];
107+
108+
if (expected != retrieved) {
109+
std::cout << "Test failed, retrieved value: "
110+
<< static_cast<int>(retrieved)
111+
<< ", but expected: " << static_cast<int>(expected)
112+
<< ", at index: " << i << std::endl;
113+
return 1;
114+
}
115+
} else {
116+
117+
DataT expected = ref_data_for_fill[(i - Offset) / Stride];
118+
119+
DataT retrieved = result[i + ResultOffset];
120+
121+
if (expected != retrieved) {
122+
std::cout << "Test failed, retrieved value: "
123+
<< static_cast<int>(retrieved)
124+
<< ", but expected: " << static_cast<int>(expected)
125+
<< ", at index: " << i << std::endl;
126+
return 1;
127+
}
128+
selected_indices_ptr++;
129+
}
130+
}
131+
132+
return 0;
133+
}
134+
135+
int main() {
136+
int test_result = 0;
137+
138+
test_result |= test_to_copy<16, false>();
139+
test_result |= test_to_copy<32 + 8, false>();
140+
test_result |= test_to_copy<32 + 10, false>();
141+
test_result |= test_to_copy<10, false>();
142+
test_result |= test_to_copy<8, false>();
143+
test_result |= test_to_copy<7, false>();
144+
test_result |= test_to_copy<15, false>();
145+
test_result |= test_to_copy<32, false>();
146+
147+
test_result |= test_to_copy<16, true>();
148+
test_result |= test_to_copy<32 + 8, true>();
149+
test_result |= test_to_copy<32 + 10, true>();
150+
test_result |= test_to_copy<10, true>();
151+
test_result |= test_to_copy<8, true>();
152+
test_result |= test_to_copy<7, true>();
153+
test_result |= test_to_copy<15, true>();
154+
test_result |= test_to_copy<32, true>();
155+
156+
test_result |= test_to_copy<16, false, 1>();
157+
test_result |= test_to_copy<32 + 8, false, 1>();
158+
test_result |= test_to_copy<32 + 10, false, 1>();
159+
test_result |= test_to_copy<10, false, 1>();
160+
test_result |= test_to_copy<8, false, 1>();
161+
test_result |= test_to_copy<7, false, 1>();
162+
test_result |= test_to_copy<15, false, 1>();
163+
test_result |= test_to_copy<32, false, 1>();
164+
165+
test_result |= test_to_copy<16, true, 1>();
166+
test_result |= test_to_copy<32 + 8, true, 1>();
167+
test_result |= test_to_copy<32 + 10, true, 1>();
168+
test_result |= test_to_copy<10, true, 1>();
169+
test_result |= test_to_copy<8, true, 1>();
170+
test_result |= test_to_copy<7, true, 1>();
171+
test_result |= test_to_copy<15, true, 1>();
172+
test_result |= test_to_copy<32, true, 1>();
173+
174+
if (!test_result) {
175+
std::cout << "Test passed" << std::endl;
176+
}
177+
178+
return test_result;
179+
}

0 commit comments

Comments
 (0)