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

Commit 7f0e3f0

Browse files
v-klochkovv-klochkov
authored
Revert "[SYCL][ESIMD]Add tests for support for different types for lsc functions (#1305)"
This reverts commit eb4ecde.
1 parent 8b8ecdf commit 7f0e3f0

24 files changed

+208
-410
lines changed

‎SYCL/ESIMD/lsc/Inputs/common.hpp

Lines changed: 3 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -7,19 +7,12 @@
77
//===----------------------------------------------------------------------===//
88

99
#include <stdlib.h>
10-
#include <sycl/bit_cast.hpp>
1110

1211
template <int case_num> class KernelID;
1312

1413
template <typename T> T get_rand() {
15-
using Tuint = std::conditional_t<
16-
sizeof(T) == 1, uint8_t,
17-
std::conditional_t<
18-
sizeof(T) == 2, uint16_t,
19-
std::conditional_t<sizeof(T) == 4, uint32_t,
20-
std::conditional_t<sizeof(T) == 8, uint64_t, T>>>>;
21-
Tuint v = rand();
22-
if constexpr (sizeof(Tuint) > 4)
14+
T v = rand();
15+
if constexpr (sizeof(T) > 4)
2316
v = (v << 32) | rand();
24-
return sycl::bit_cast<T>(v);
17+
return v;
2518
}

‎SYCL/ESIMD/lsc/Inputs/lsc_surf_load.hpp

Lines changed: 10 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -38,8 +38,8 @@ bool test(uint32_t pmask = 0xffffffff) {
3838
}
3939

4040
uint16_t Size = Groups * Threads * VL * VS;
41-
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
42-
Tuint vmask = (Tuint)-1;
41+
42+
T vmask = (T)-1;
4343
if constexpr (DS == lsc_data_size::u8u32)
4444
vmask = (T)0xff;
4545
if constexpr (DS == lsc_data_size::u16u32)
@@ -123,24 +123,20 @@ bool test(uint32_t pmask = 0xffffffff) {
123123

124124
if constexpr (transpose) {
125125
for (int i = 0; i < Size; i++) {
126-
Tuint e = sycl::bit_cast<Tuint>(in[i]);
127-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
128-
if (out_val != e) {
126+
T e = in[i];
127+
if (out[i] != e) {
129128
passed = false;
130-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
131-
<< " vs etalon = 0x" << e << std::dec << std::endl;
129+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
130+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
132131
}
133132
}
134133
} else {
135134
for (int i = 0; i < Size; i++) {
136-
Tuint in_val = sycl::bit_cast<Tuint>(in[i]);
137-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
138-
Tuint e = (pmask >> ((i / VS) % VL)) & 1 ? in_val & vmask
139-
: sycl::bit_cast<Tuint>(old_val);
140-
if (out_val != e) {
135+
T e = (pmask >> ((i / VS) % VL)) & 1 ? in[i] & vmask : old_val;
136+
if (out[i] != e) {
141137
passed = false;
142-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
143-
<< " vs etalon = 0x" << e << std::dec << std::endl;
138+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
139+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
144140
}
145141
}
146142
}

‎SYCL/ESIMD/lsc/Inputs/lsc_surf_store.hpp

Lines changed: 12 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -37,8 +37,8 @@ bool test(uint32_t pmask = 0xffffffff) {
3737
}
3838

3939
uint16_t Size = Groups * Threads * VL * VS;
40-
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
41-
Tuint vmask = (Tuint)-1;
40+
41+
T vmask = (T)-1;
4242
if constexpr (DS == lsc_data_size::u8u32)
4343
vmask = (T)0xff;
4444
if constexpr (DS == lsc_data_size::u16u32)
@@ -103,28 +103,22 @@ bool test(uint32_t pmask = 0xffffffff) {
103103

104104
if constexpr (transpose) {
105105
for (int i = 0; i < Size; i++) {
106-
T expected_value = new_val + i;
107-
Tuint e = sycl::bit_cast<Tuint>(expected_value);
108-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
109-
if (out_val != e) {
106+
T e = new_val + i;
107+
if (out[i] != e) {
110108
passed = false;
111-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
112-
<< " vs etalon = 0x" << e << std::dec << std::endl;
109+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
110+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
113111
}
114112
}
115113
} else {
116114
for (int i = 0; i < Size; i++) {
117-
T expected_value = new_val + i;
118-
Tuint in_val = sycl::bit_cast<Tuint>(expected_value);
119-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
120-
Tuint e =
121-
(pmask >> ((i / VS) % VL)) & 1
122-
? (in_val & vmask) | (sycl::bit_cast<Tuint>(old_val) & ~vmask)
123-
: sycl::bit_cast<Tuint>(old_val);
124-
if (out_val != e) {
115+
T e = (pmask >> ((i / VS) % VL)) & 1
116+
? ((new_val + i) & vmask) | (old_val & ~vmask)
117+
: old_val;
118+
if (out[i] != e) {
125119
passed = false;
126-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
127-
<< " vs etalon = 0x" << e << std::dec << std::endl;
120+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
121+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
128122
}
129123
}
130124
}

‎SYCL/ESIMD/lsc/Inputs/lsc_usm_load.hpp

Lines changed: 10 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -38,8 +38,8 @@ bool test(uint32_t pmask = 0xffffffff) {
3838
}
3939

4040
uint16_t Size = Groups * Threads * VL * VS;
41-
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
42-
Tuint vmask = (Tuint)-1;
41+
42+
T vmask = (T)-1;
4343
if constexpr (DS == lsc_data_size::u8u32)
4444
vmask = (T)0xff;
4545
if constexpr (DS == lsc_data_size::u16u32)
@@ -124,24 +124,20 @@ bool test(uint32_t pmask = 0xffffffff) {
124124

125125
if constexpr (transpose) {
126126
for (int i = 0; i < Size; i++) {
127-
Tuint e = sycl::bit_cast<Tuint>(in[i]);
128-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
129-
if (out_val != e) {
127+
T e = in[i];
128+
if (out[i] != e) {
130129
passed = false;
131-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
132-
<< " vs etalon = 0x" << e << std::dec << std::endl;
130+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
131+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
133132
}
134133
}
135134
} else {
136135
for (int i = 0; i < Size; i++) {
137-
Tuint in_val = sycl::bit_cast<Tuint>(in[i]);
138-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
139-
Tuint e = (pmask >> ((i / VS) % VL)) & 1 ? in_val & vmask
140-
: sycl::bit_cast<Tuint>(old_val);
141-
if (out_val != e) {
136+
T e = (pmask >> ((i / VS) % VL)) & 1 ? in[i] & vmask : old_val;
137+
if (out[i] != e) {
142138
passed = false;
143-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
144-
<< " vs etalon = 0x" << e << std::dec << std::endl;
139+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
140+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
145141
}
146142
}
147143
}

‎SYCL/ESIMD/lsc/Inputs/lsc_usm_store.hpp

Lines changed: 11 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -37,9 +37,8 @@ bool test(uint32_t pmask = 0xffffffff) {
3737
}
3838

3939
uint16_t Size = Groups * Threads * VL * VS;
40-
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
4140

42-
Tuint vmask = (Tuint)-1;
41+
T vmask = (T)-1;
4342
if constexpr (DS == lsc_data_size::u8u32)
4443
vmask = (T)0xff;
4544
if constexpr (DS == lsc_data_size::u16u32)
@@ -104,29 +103,22 @@ bool test(uint32_t pmask = 0xffffffff) {
104103

105104
if constexpr (transpose) {
106105
for (int i = 0; i < Size; i++) {
107-
T expected_value = new_val + i;
108-
Tuint e = sycl::bit_cast<Tuint>(expected_value);
109-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
110-
if (out_val != e) {
106+
T e = new_val + i;
107+
if (out[i] != e) {
111108
passed = false;
112-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
113-
<< " vs etalon = 0x" << e << std::dec << std::endl;
109+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
110+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
114111
}
115112
}
116113
} else {
117114
for (int i = 0; i < Size; i++) {
118-
T expected_value = new_val + i;
119-
Tuint in_val = sycl::bit_cast<Tuint>(expected_value);
120-
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
121-
122-
Tuint e =
123-
(pmask >> ((i / VS) % VL)) & 1
124-
? (in_val & vmask) | (sycl::bit_cast<Tuint>(old_val) & ~vmask)
125-
: sycl::bit_cast<Tuint>(old_val);
126-
if (out_val != e) {
115+
T e = (pmask >> ((i / VS) % VL)) & 1
116+
? ((new_val + i) & vmask) | (old_val & ~vmask)
117+
: old_val;
118+
if (out[i] != e) {
127119
passed = false;
128-
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
129-
<< " vs etalon = 0x" << e << std::dec << std::endl;
120+
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
121+
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
130122
}
131123
}
132124
}

‎SYCL/ESIMD/lsc/lsc_surf_load_u32.cpp

Lines changed: 13 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -14,34 +14,25 @@
1414

1515
constexpr uint32_t seed = 199;
1616

17-
template <int TestCastNum, typename T> bool tests() {
17+
int main(void) {
18+
srand(seed);
1819
bool passed = true;
1920

2021
// non transpose
21-
passed &= test<TestCastNum, T, 1, 4, 32, 1, false>(rand());
22-
passed &= test<TestCastNum + 1, T, 1, 4, 32, 2, false>(rand());
23-
passed &= test<TestCastNum + 2, T, 1, 4, 16, 2, false>(rand());
24-
passed &= test<TestCastNum + 3, T, 1, 4, 4, 1, false>(rand());
25-
passed &= test<TestCastNum + 4, T, 1, 1, 1, 1, false>(1);
26-
passed &= test<TestCastNum + 5, T, 2, 1, 1, 1, false>(1);
22+
passed &= test<0, uint32_t, 1, 4, 32, 1, false>(rand());
23+
passed &= test<1, uint32_t, 1, 4, 32, 2, false>(rand());
24+
passed &= test<2, uint32_t, 1, 4, 16, 2, false>(rand());
25+
passed &= test<3, uint32_t, 1, 4, 4, 1, false>(rand());
26+
passed &= test<4, uint32_t, 1, 1, 1, 1, false>(1);
27+
passed &= test<5, uint32_t, 2, 1, 1, 1, false>(1);
2728

28-
// passed &= test<TestCastNum+6, T, 1, 4, 8, 2, false>(rand());
29-
// passed &= test<TestCastNum+7, T, 1, 4, 8, 3, false>(rand());
29+
// passed &= test<6, uint32_t, 1, 4, 8, 2, false>(rand());
30+
// passed &= test<7, uint32_t, 1, 4, 8, 3, false>(rand());
3031

3132
// transpose
32-
passed &= test<TestCastNum + 8, T, 1, 4, 1, 32, true>();
33-
passed &= test<TestCastNum + 9, T, 2, 2, 1, 16, true>();
34-
passed &= test<TestCastNum + 10, T, 4, 4, 1, 4, true>();
35-
return passed;
36-
}
37-
38-
int main(void) {
39-
srand(seed);
40-
bool passed = true;
41-
42-
passed &= tests<0, uint32_t>();
43-
passed &= tests<11, float>();
44-
passed &= tests<22, sycl::ext::intel::experimental::esimd::tfloat32>();
33+
passed &= test<8, uint32_t, 1, 4, 1, 32, true>();
34+
passed &= test<9, uint32_t, 2, 2, 1, 16, true>();
35+
passed &= test<10, uint32_t, 4, 4, 1, 4, true>();
4536

4637
std::cout << (passed ? "Passed\n" : "FAILED\n");
4738
return passed ? 0 : 1;

‎SYCL/ESIMD/lsc/lsc_surf_load_u32_stateless.cpp

Lines changed: 0 additions & 13 deletions
This file was deleted.

‎SYCL/ESIMD/lsc/lsc_surf_load_u64.cpp

Lines changed: 14 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -13,32 +13,26 @@
1313
#include "Inputs/lsc_surf_load.hpp"
1414

1515
constexpr uint32_t seed = 198;
16-
template <int TestCastNum, typename T> bool tests() {
16+
17+
int main(void) {
18+
srand(seed);
1719
bool passed = true;
1820

1921
// non transpose
20-
passed &= test<TestCastNum, T, 1, 4, 32, 1, false>(rand());
21-
passed &= test<TestCastNum + 1, T, 1, 4, 32, 2, false>(rand());
22-
passed &= test<TestCastNum + 2, T, 1, 4, 16, 2, false>(rand());
23-
passed &= test<TestCastNum + 3, T, 1, 4, 4, 1, false>(rand());
24-
passed &= test<TestCastNum + 4, T, 1, 1, 1, 1, false>(1);
25-
passed &= test<TestCastNum + 5, T, 2, 1, 1, 1, false>(1);
22+
passed &= test<0, uint64_t, 1, 4, 32, 1, false>(rand());
23+
passed &= test<1, uint64_t, 1, 4, 32, 2, false>(rand());
24+
passed &= test<2, uint64_t, 1, 4, 16, 2, false>(rand());
25+
passed &= test<3, uint64_t, 1, 4, 4, 1, false>(rand());
26+
passed &= test<4, uint64_t, 1, 1, 1, 1, false>(1);
27+
passed &= test<5, uint64_t, 2, 1, 1, 1, false>(1);
2628

27-
// passed &= test<TestCastNum+6, T, 1, 4, 8, 2, false>(rand());
28-
// passed &= test<TestCastNum+7, T, 1, 4, 8, 3, false>(rand());
29+
// passed &= test<6, uint64_t, 1, 4, 8, 2, false>(rand());
30+
// passed &= test<7, uint64_t, 1, 4, 8, 3, false>(rand());
2931

3032
// transpose
31-
passed &= test<TestCastNum + 8, T, 1, 4, 1, 32, true>();
32-
passed &= test<TestCastNum + 9, T, 2, 2, 1, 16, true>();
33-
passed &= test<TestCastNum + 10, T, 4, 4, 1, 4, true>();
34-
return passed;
35-
}
36-
int main(void) {
37-
srand(seed);
38-
bool passed = true;
39-
40-
passed &= tests<0, uint64_t>();
41-
passed &= tests<11, double>();
33+
passed &= test<8, uint64_t, 1, 4, 1, 32, true>();
34+
passed &= test<9, uint64_t, 2, 2, 1, 16, true>();
35+
passed &= test<10, uint64_t, 4, 4, 1, 4, true>();
4236

4337
std::cout << (passed ? "Passed\n" : "FAILED\n");
4438
return passed ? 0 : 1;

‎SYCL/ESIMD/lsc/lsc_surf_load_u64_stateless.cpp

Lines changed: 0 additions & 13 deletions
This file was deleted.

‎SYCL/ESIMD/lsc/lsc_surf_load_u8_u16.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,6 @@ int main(void) {
2929

3030
passed &= tests<0, uint8_t>();
3131
passed &= tests<3, uint16_t>();
32-
passed &= tests<6, sycl::ext::oneapi::experimental::bfloat16>();
33-
passed &= tests<9, half>();
3432

3533
std::cout << (passed ? "Passed\n" : "FAILED\n");
3634
return passed ? 0 : 1;

‎SYCL/ESIMD/lsc/lsc_surf_load_u8_u16_stateless.cpp

Lines changed: 0 additions & 13 deletions
This file was deleted.

0 commit comments

Comments
 (0)