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

Commit c108aa3

Browse files
fineg74fineg74
committed
Revert previous changes
1 parent c2e1182 commit c108aa3

File tree

2 files changed

+27
-35
lines changed

2 files changed

+27
-35
lines changed

‎SYCL/ESIMD/lsc/lsc_predicate.cpp

Lines changed: 15 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -29,26 +29,22 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
2929
auto vec_1 = std::vector<int>(size);
3030
auto vec_2 = std::vector<int>(size);
3131
auto vec_3 = std::vector<int>(size);
32-
auto vec_input = std::vector<int>(size);
3332

3433
std::iota(vec_0.begin(), vec_0.end(), 0);
3534
std::iota(vec_1.begin(), vec_1.end(), 0);
3635
std::iota(vec_2.begin(), vec_2.end(), 0);
3736
std::iota(vec_3.begin(), vec_3.end(), 0);
38-
std::iota(vec_input.begin(), vec_input.end(), 0);
3937
auto buf_0 = buffer{vec_0};
4038
auto buf_1 = buffer{vec_1};
4139
auto buf_2 = buffer{vec_2};
4240
auto buf_3 = buffer{vec_3};
43-
auto buf_input = buffer{vec_input};
4441

4542
try {
4643
q.submit([&](handler &h) {
47-
auto access_0 = buf_0.template get_access<access::mode::write>(h);
48-
auto access_1 = buf_1.template get_access<access::mode::write>(h);
49-
auto access_2 = buf_2.template get_access<access::mode::write>(h);
50-
auto access_3 = buf_3.template get_access<access::mode::write>(h);
51-
auto access_input = buf_input.template get_access<access::mode::read>(h);
44+
auto access_0 = buf_0.template get_access<access::mode::read_write>(h);
45+
auto access_1 = buf_1.template get_access<access::mode::read_write>(h);
46+
auto access_2 = buf_2.template get_access<access::mode::read_write>(h);
47+
auto access_3 = buf_3.template get_access<access::mode::read_write>(h);
5248

5349
h.parallel_for(
5450
range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL {
@@ -57,28 +53,28 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
5753
auto pred_enable = simd_mask<1>(1);
5854
auto pred_disable = simd_mask<1>(0);
5955

60-
auto data_0 = lsc_block_load<int, SIMDSize>(access_input, offset,
61-
pred_enable);
56+
auto data_0 =
57+
lsc_block_load<int, SIMDSize>(access_0, offset, pred_enable);
6258
lsc_block_store<int, SIMDSize>(access_0, offset, data_0 * 2,
6359
pred_enable);
6460

65-
auto data_1 = lsc_block_load<int, SIMDSize>(access_input, offset,
66-
pred_disable);
61+
auto data_1 =
62+
lsc_block_load<int, SIMDSize>(access_1, offset, pred_disable);
6763
lsc_block_store<int, SIMDSize>(access_1, offset, data_1 * 2,
6864
pred_enable);
6965

70-
auto data_2 = lsc_block_load<int, SIMDSize>(access_input, offset,
71-
pred_enable);
66+
auto data_2 =
67+
lsc_block_load<int, SIMDSize>(access_2, offset, pred_enable);
7268
lsc_block_store<int, SIMDSize>(access_2, offset, data_2 * 2,
7369
pred_disable);
7470

75-
auto data_3 = lsc_block_load<int, SIMDSize>(access_input, offset,
76-
pred_disable);
71+
auto data_3 =
72+
lsc_block_load<int, SIMDSize>(access_3, offset, pred_disable);
7773
lsc_block_store<int, SIMDSize>(access_3, offset, data_3 * 2,
7874
pred_disable);
7975
});
8076
});
81-
q.wait();
77+
q.wait_and_throw();
8278
} catch (sycl::exception e) {
8379
std::cout << "SYCL exception caught: " << e.what();
8480
return 1;
@@ -126,7 +122,7 @@ template <unsigned SIMDSize> int testUSM(queue q) {
126122
std::iota(vec_1, vec_1 + size, 0);
127123
std::iota(vec_2, vec_2 + size, 0);
128124
std::iota(vec_3, vec_3 + size, 0);
129-
std::iota(vec_input, vec_input + size, 0);
125+
std::iota(vec_input, vec_3 + size, 0);
130126

131127
try {
132128
q.submit([&](handler &h) {
@@ -157,7 +153,7 @@ template <unsigned SIMDSize> int testUSM(queue q) {
157153
pred_disable);
158154
});
159155
});
160-
q.wait();
156+
q.wait_and_throw();
161157
} catch (sycl::exception e) {
162158
std::cout << "SYCL exception caught: " << e.what();
163159
sycl::free(vec_0, q);

‎SYCL/ESIMD/lsc/lsc_predicate_stateless.cpp

Lines changed: 12 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -31,26 +31,22 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
3131
auto vec_1 = std::vector<int>(size);
3232
auto vec_2 = std::vector<int>(size);
3333
auto vec_3 = std::vector<int>(size);
34-
auto vec_input = std::vector<int>(size);
3534

3635
std::iota(vec_0.begin(), vec_0.end(), 0);
3736
std::iota(vec_1.begin(), vec_1.end(), 0);
3837
std::iota(vec_2.begin(), vec_2.end(), 0);
3938
std::iota(vec_3.begin(), vec_3.end(), 0);
40-
std::iota(vec_input.begin(), vec_input.end(), 0);
4139
auto buf_0 = buffer{vec_0};
4240
auto buf_1 = buffer{vec_1};
4341
auto buf_2 = buffer{vec_2};
4442
auto buf_3 = buffer{vec_3};
45-
auto buf_input = buffer{vec_input};
4643

4744
try {
4845
q.submit([&](handler &h) {
49-
auto access_0 = buf_0.template get_access<access::mode::write>(h);
50-
auto access_1 = buf_1.template get_access<access::mode::write>(h);
51-
auto access_2 = buf_2.template get_access<access::mode::write>(h);
52-
auto access_3 = buf_3.template get_access<access::mode::write>(h);
53-
auto access_input = buf_input.template get_access<access::mode::read>(h);
46+
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);
48+
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);
5450

5551
h.parallel_for(
5652
range<1>{size / SIMDSize}, [=](id<1> id) SYCL_ESIMD_KERNEL {
@@ -60,23 +56,23 @@ template <unsigned SIMDSize> int testAccessor(queue q) {
6056
auto pred_enable = simd_mask<1>(1);
6157
auto pred_disable = simd_mask<1>(0);
6258

63-
auto data_0 = lsc_block_load<int, SIMDSize>(access_input, offset,
64-
pred_enable);
59+
auto data_0 =
60+
lsc_block_load<int, SIMDSize>(access_0, offset, pred_enable);
6561
lsc_block_store<int, SIMDSize>(access_0, offset, data_0 * 2,
6662
pred_enable);
6763

68-
auto data_1 = lsc_block_load<int, SIMDSize>(access_input, offset,
69-
pred_disable);
64+
auto data_1 =
65+
lsc_block_load<int, SIMDSize>(access_1, offset, pred_disable);
7066
lsc_block_store<int, SIMDSize>(access_1, offset, data_1 * 2,
7167
pred_enable);
7268

73-
auto data_2 = lsc_block_load<int, SIMDSize>(access_input, offset,
74-
pred_enable);
69+
auto data_2 =
70+
lsc_block_load<int, SIMDSize>(access_2, offset, pred_enable);
7571
lsc_block_store<int, SIMDSize>(access_2, offset, data_2 * 2,
7672
pred_disable);
7773

78-
auto data_3 = lsc_block_load<int, SIMDSize>(access_input, offset,
79-
pred_disable);
74+
auto data_3 =
75+
lsc_block_load<int, SIMDSize>(access_3, offset, pred_disable);
8076
lsc_block_store<int, SIMDSize>(access_3, offset, data_3 * 2,
8177
pred_disable);
8278
});

0 commit comments

Comments
 (0)