Skip to content

Commit fa6e6ec

Browse files
fineg74fineg74
authored
[SYCL][ESIMD] Fix lsc predicate test that failed on PVC hardware (intel#1315) (intel#1305)
The removed sub-tests were incorrect as they stored garbage (random return from lsc_block_load(with preducate=0) to result and then checking it with something meaningful expected values.
1 parent 7aa8eb5 commit fa6e6ec

File tree

2 files changed

+42
-82
lines changed

2 files changed

+42
-82
lines changed

‎SYCL/ESIMD/lsc/lsc_predicate.cpp

Lines changed: 29 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -14,8 +14,6 @@
1414

1515
#include "../esimd_test_utils.hpp"
1616

17-
#include <algorithm>
18-
#include <cmath>
1917
#include <numeric>
2018
#include <sycl/ext/intel/esimd.hpp>
2119
#include <sycl/sycl.hpp>
@@ -28,25 +26,17 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
2826
auto size = size_t{128};
2927

3028
auto vec_0 = std::vector<int>(size);
31-
auto vec_1 = std::vector<int>(size);
3229
auto vec_2 = std::vector<int>(size);
33-
auto vec_3 = std::vector<int>(size);
3430

3531
std::iota(vec_0.begin(), vec_0.end(), 0);
36-
std::iota(vec_1.begin(), vec_1.end(), 0);
3732
std::iota(vec_2.begin(), vec_2.end(), 0);
38-
std::iota(vec_3.begin(), vec_3.end(), 0);
3933
auto buf_0 = buffer{vec_0};
40-
auto buf_1 = buffer{vec_1};
4134
auto buf_2 = buffer{vec_2};
42-
auto buf_3 = buffer{vec_3};
4335

4436
try {
4537
q.submit([&](handler &h) {
4638
auto access_0 = buf_0.template get_access<access::mode::read_write>(h);
47-
auto access_1 = buf_1.template get_access<access::mode::read_write>(h);
4839
auto access_2 = buf_2.template get_access<access::mode::read_write>(h);
49-
auto access_3 = buf_3.template get_access<access::mode::read_write>(h);
5040

5141
h.parallel_for(
5242
range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL {
@@ -61,38 +51,34 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
6151
lsc_block_store<int, SIMDSize>(access_0, offset, data_0 * 2,
6252
pred_enable);
6353

64-
auto data_1 =
65-
lsc_block_load<int, SIMDSize>(access_1, offset, pred_disable);
66-
lsc_block_store<int, SIMDSize>(access_1, offset, data_1 * 2,
67-
pred_enable);
68-
6954
auto data_2 =
7055
lsc_block_load<int, SIMDSize>(access_2, offset, pred_enable);
7156
lsc_block_store<int, SIMDSize>(access_2, offset, data_2 * 2,
7257
pred_disable);
73-
74-
auto data_3 =
75-
lsc_block_load<int, SIMDSize>(access_3, offset, pred_disable);
76-
lsc_block_store<int, SIMDSize>(access_3, offset, data_3 * 2,
77-
pred_disable);
7858
});
7959
});
8060
q.wait();
8161
buf_0.template get_access<access::mode::read_write>();
82-
buf_1.template get_access<access::mode::read_write>();
8362
buf_2.template get_access<access::mode::read_write>();
84-
buf_3.template get_access<access::mode::read_write>();
8563
} catch (sycl::exception e) {
8664
std::cout << "SYCL exception caught: " << e.what();
8765
return 1;
8866
}
8967

9068
auto error = 0;
9169
for (auto i = 0; i != size; ++i) {
92-
error += vec_0[i] != 2 * i;
93-
error += vec_1[i] > 0;
94-
error += vec_2[i] != i;
95-
error += vec_3[i] != i;
70+
if (vec_0[i] != 2 * i) {
71+
++error;
72+
std::cout << " Accessor Test 1 out[" << i << "] = 0x" << std::hex
73+
<< vec_0[i] << " vs etalon = 0x" << 2 * i << std::dec
74+
<< std::endl;
75+
}
76+
77+
if (vec_2[i] != i) {
78+
++error;
79+
std::cout << " Accessor Test 2 out[" << i << "] = 0x" << std::hex
80+
<< vec_2[i] << " vs etalon = 0x" << i << std::dec << std::endl;
81+
}
9682
}
9783
std::cout << "Accessor lsc predicate test ";
9884
std::cout << (error != 0 ? "FAILED" : "passed") << std::endl;
@@ -104,80 +90,67 @@ template <unsigned SIMDSize> int testUSM(queue q) {
10490

10591
auto *vec_0 = malloc_shared<int>(size, q);
10692
auto *vec_1 = malloc_shared<int>(size, q);
107-
auto *vec_2 = malloc_shared<int>(size, q);
108-
auto *vec_3 = malloc_shared<int>(size, q);
10993
std::iota(vec_0, vec_0 + size, 0);
11094
std::iota(vec_1, vec_1 + size, 0);
111-
std::iota(vec_2, vec_2 + size, 0);
112-
std::iota(vec_3, vec_3 + size, 0);
11395

11496
try {
11597
q.submit([&](handler &h) {
11698
h.parallel_for(
11799
range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL {
118100
auto offset = id[0] * SIMDSize;
119-
auto offsets = simd<uint32_t, SIMDSize>(id * SIMDSize * sizeof(int),
120-
sizeof(int));
101+
121102
auto pred_enable = simd_mask<1>(1);
122103
auto pred_disable = simd_mask<1>(0);
123104

124105
auto data_0 =
125106
lsc_block_load<int, SIMDSize>(vec_0 + offset, pred_enable);
126107
lsc_block_store<int, SIMDSize>(vec_0 + offset, data_0 * 2,
127108
pred_enable);
128-
129109
auto data_1 =
130-
lsc_block_load<int, SIMDSize>(vec_1 + offset, pred_disable);
110+
lsc_block_load<int, SIMDSize>(vec_1 + offset, pred_enable);
131111
lsc_block_store<int, SIMDSize>(vec_1 + offset, data_1 * 2,
132-
pred_enable);
133-
134-
auto data_2 =
135-
lsc_block_load<int, SIMDSize>(vec_2 + offset, pred_enable);
136-
lsc_block_store<int, SIMDSize>(vec_2 + offset, data_2 * 2,
137-
pred_disable);
138-
auto data_3 =
139-
lsc_block_load<int, SIMDSize>(vec_3 + offset, pred_disable);
140-
lsc_block_store<int, SIMDSize>(vec_3 + offset, data_3 * 2,
141112
pred_disable);
142113
});
143114
});
144-
q.wait();
115+
q.wait_and_throw();
145116
} catch (sycl::exception e) {
146117
std::cout << "SYCL exception caught: " << e.what();
147118
sycl::free(vec_0, q);
148119
sycl::free(vec_1, q);
149-
sycl::free(vec_2, q);
150-
sycl::free(vec_3, q);
151120
return 1;
152121
}
153122

154123
int error = 0;
155124
for (auto i = 0; i != size; ++i) {
156-
error += vec_0[i] != 2 * i;
157-
error += vec_1[i] > 0;
158-
error += vec_2[i] != i;
159-
error += vec_3[i] != i;
125+
if (vec_0[i] != 2 * i) {
126+
++error;
127+
std::cout << " USM Test 1 out[" << i << "] = 0x" << std::hex << vec_0[i]
128+
<< " vs etalon = 0x" << 2 * i << std::dec << std::endl;
129+
}
130+
131+
if (vec_1[i] != i) {
132+
++error;
133+
std::cout << " USM Test 2 out[" << i << "] = 0x" << std::hex << vec_1[i]
134+
<< " vs etalon = 0x" << i << std::dec << std::endl;
135+
}
160136
}
161137
sycl::free(vec_0, q);
162138
sycl::free(vec_1, q);
163-
sycl::free(vec_2, q);
164-
sycl::free(vec_3, q);
165139
std::cout << "USM lsc predicate test ";
166140
std::cout << (error != 0 ? "FAILED" : "passed") << std::endl;
167141
return error;
168142
}
169143

170144
int main() {
145+
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
171146

172-
auto q =
173-
queue{esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()};
174147
auto device = q.get_device();
175148
std::cout << "Device name: " << device.get_info<info::device::name>()
176149
<< std::endl;
177150

178151
int error = testUSM<8>(q);
179-
error = testUSM<16>(q);
180-
error = testUSM<32>(q);
152+
error += testUSM<16>(q);
153+
error += testUSM<32>(q);
181154

182155
error += testAccessor<8>(q);
183156
error += testAccessor<16>(q);

‎SYCL/ESIMD/lsc/lsc_predicate_stateless.cpp

Lines changed: 13 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,6 @@
1616

1717
#include "../esimd_test_utils.hpp"
1818

19-
#include <algorithm>
20-
#include <cmath>
2119
#include <numeric>
2220
#include <sycl/ext/intel/esimd.hpp>
2321
#include <sycl/sycl.hpp>
@@ -30,25 +28,17 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
3028
auto size = size_t{128};
3129

3230
auto vec_0 = std::vector<int>(size);
33-
auto vec_1 = std::vector<int>(size);
3431
auto vec_2 = std::vector<int>(size);
35-
auto vec_3 = std::vector<int>(size);
3632

3733
std::iota(vec_0.begin(), vec_0.end(), 0);
38-
std::iota(vec_1.begin(), vec_1.end(), 0);
3934
std::iota(vec_2.begin(), vec_2.end(), 0);
40-
std::iota(vec_3.begin(), vec_3.end(), 0);
4135
auto buf_0 = buffer{vec_0};
42-
auto buf_1 = buffer{vec_1};
4336
auto buf_2 = buffer{vec_2};
44-
auto buf_3 = buffer{vec_3};
4537

4638
try {
4739
q.submit([&](handler &h) {
4840
auto access_0 = buf_0.template get_access<access::mode::read_write>(h);
49-
auto access_1 = buf_1.template get_access<access::mode::read_write>(h);
5041
auto access_2 = buf_2.template get_access<access::mode::read_write>(h);
51-
auto access_3 = buf_3.template get_access<access::mode::read_write>(h);
5242

5343
h.parallel_for(
5444
range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL {
@@ -63,39 +53,36 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
6353
lsc_block_store<int, SIMDSize>(access_0, offset, data_0 * 2,
6454
pred_enable);
6555

66-
auto data_1 =
67-
lsc_block_load<int, SIMDSize>(access_1, offset, pred_disable);
68-
lsc_block_store<int, SIMDSize>(access_1, offset, data_1 * 2,
69-
pred_enable);
70-
7156
auto data_2 =
7257
lsc_block_load<int, SIMDSize>(access_2, offset, pred_enable);
7358
lsc_block_store<int, SIMDSize>(access_2, offset, data_2 * 2,
7459
pred_disable);
75-
76-
auto data_3 =
77-
lsc_block_load<int, SIMDSize>(access_3, offset, pred_disable);
78-
lsc_block_store<int, SIMDSize>(access_3, offset, data_3 * 2,
79-
pred_disable);
8060
});
8161
});
8262
q.wait();
8363
buf_0.template get_access<access::mode::read_write>();
84-
buf_1.template get_access<access::mode::read_write>();
8564
buf_2.template get_access<access::mode::read_write>();
86-
buf_3.template get_access<access::mode::read_write>();
8765
} catch (sycl::exception e) {
8866
std::cout << "SYCL exception caught: " << e.what();
8967
return 1;
9068
}
9169

9270
auto error = 0;
9371
for (auto i = 0; i != size; ++i) {
94-
error += vec_0[i] != 2 * i;
95-
error += vec_1[i] > 0;
96-
error += vec_2[i] != i;
97-
error += vec_3[i] != i;
72+
if (vec_0[i] != 2 * i) {
73+
++error;
74+
std::cout << " Accessor Test 1 out[" << i << "] = 0x" << std::hex
75+
<< vec_0[i] << " vs etalon = 0x" << 2 * i << std::dec
76+
<< std::endl;
77+
}
78+
79+
if (vec_2[i] != i) {
80+
++error;
81+
std::cout << " Accessor Test 2 out[" << i << "] = 0x" << std::hex
82+
<< vec_2[i] << " vs etalon = 0x" << i << std::dec << std::endl;
83+
}
9884
}
85+
9986
std::cout << "Accessor lsc predicate test ";
10087
std::cout << (error != 0 ? "FAILED" : "passed") << std::endl;
10188
return error;

0 commit comments

Comments
 (0)