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

Commit eb4ecde

Browse files
fineg74fineg74
authored
[SYCL][ESIMD]Add tests for support for different types for lsc functions (#1305)
* Add tests for support for different types for lsc functions
1 parent 61a5926 commit eb4ecde

24 files changed

+410
-208
lines changed

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

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

99
#include <stdlib.h>
10+
#include <sycl/bit_cast.hpp>
1011

1112
template <int case_num> class KernelID;
1213

1314
template <typename T> T get_rand() {
14-
T v = rand();
15-
if constexpr (sizeof(T) > 4)
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)
1623
v = (v << 32) | rand();
17-
return v;
24+
return sycl::bit_cast<T>(v);
1825
}

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

Lines changed: 14 additions & 10 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-
42-
T vmask = (T)-1;
41+
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
42+
Tuint vmask = (Tuint)-1;
4343
if constexpr (DS == lsc_data_size::u8u32)
4444
vmask = (T)0xff;
4545
if constexpr (DS == lsc_data_size::u16u32)
@@ -123,20 +123,24 @@ bool test(uint32_t pmask = 0xffffffff) {
123123

124124
if constexpr (transpose) {
125125
for (int i = 0; i < Size; i++) {
126-
T e = in[i];
127-
if (out[i] != e) {
126+
Tuint e = sycl::bit_cast<Tuint>(in[i]);
127+
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
128+
if (out_val != e) {
128129
passed = false;
129-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
130-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
130+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
131+
<< " vs etalon = 0x" << e << std::dec << std::endl;
131132
}
132133
}
133134
} else {
134135
for (int i = 0; i < Size; i++) {
135-
T e = (pmask >> ((i / VS) % VL)) & 1 ? in[i] & vmask : old_val;
136-
if (out[i] != e) {
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) {
137141
passed = false;
138-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
139-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
142+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
143+
<< " vs etalon = 0x" << e << std::dec << std::endl;
140144
}
141145
}
142146
}

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

Lines changed: 18 additions & 12 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-
41-
T vmask = (T)-1;
40+
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
41+
Tuint vmask = (Tuint)-1;
4242
if constexpr (DS == lsc_data_size::u8u32)
4343
vmask = (T)0xff;
4444
if constexpr (DS == lsc_data_size::u16u32)
@@ -103,22 +103,28 @@ bool test(uint32_t pmask = 0xffffffff) {
103103

104104
if constexpr (transpose) {
105105
for (int i = 0; i < Size; i++) {
106-
T e = new_val + i;
107-
if (out[i] != e) {
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) {
108110
passed = false;
109-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
110-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
111+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
112+
<< " vs etalon = 0x" << e << std::dec << std::endl;
111113
}
112114
}
113115
} else {
114116
for (int i = 0; i < Size; i++) {
115-
T e = (pmask >> ((i / VS) % VL)) & 1
116-
? ((new_val + i) & vmask) | (old_val & ~vmask)
117-
: old_val;
118-
if (out[i] != e) {
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) {
119125
passed = false;
120-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
121-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
126+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
127+
<< " vs etalon = 0x" << e << std::dec << std::endl;
122128
}
123129
}
124130
}

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

Lines changed: 14 additions & 10 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-
42-
T vmask = (T)-1;
41+
using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
42+
Tuint vmask = (Tuint)-1;
4343
if constexpr (DS == lsc_data_size::u8u32)
4444
vmask = (T)0xff;
4545
if constexpr (DS == lsc_data_size::u16u32)
@@ -124,20 +124,24 @@ bool test(uint32_t pmask = 0xffffffff) {
124124

125125
if constexpr (transpose) {
126126
for (int i = 0; i < Size; i++) {
127-
T e = in[i];
128-
if (out[i] != e) {
127+
Tuint e = sycl::bit_cast<Tuint>(in[i]);
128+
Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
129+
if (out_val != e) {
129130
passed = false;
130-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
131-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
131+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
132+
<< " vs etalon = 0x" << e << std::dec << std::endl;
132133
}
133134
}
134135
} else {
135136
for (int i = 0; i < Size; i++) {
136-
T e = (pmask >> ((i / VS) % VL)) & 1 ? in[i] & vmask : old_val;
137-
if (out[i] != e) {
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) {
138142
passed = false;
139-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
140-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
143+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
144+
<< " vs etalon = 0x" << e << std::dec << std::endl;
141145
}
142146
}
143147
}

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

Lines changed: 19 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -37,8 +37,9 @@ 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)>;
4041

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

104105
if constexpr (transpose) {
105106
for (int i = 0; i < Size; i++) {
106-
T e = new_val + i;
107-
if (out[i] != e) {
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) {
108111
passed = false;
109-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
110-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
112+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
113+
<< " vs etalon = 0x" << e << std::dec << std::endl;
111114
}
112115
}
113116
} else {
114117
for (int i = 0; i < Size; i++) {
115-
T e = (pmask >> ((i / VS) % VL)) & 1
116-
? ((new_val + i) & vmask) | (old_val & ~vmask)
117-
: old_val;
118-
if (out[i] != e) {
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) {
119127
passed = false;
120-
std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
121-
<< " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
128+
std::cout << "out[" << i << "] = 0x" << std::hex << out_val
129+
<< " vs etalon = 0x" << e << std::dec << std::endl;
122130
}
123131
}
124132
}

‎SYCL/ESIMD/lsc/lsc_surf_load_u32.cpp

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

1515
constexpr uint32_t seed = 199;
1616

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

2120
// non transpose
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);
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);
2827

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

3231
// transpose
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>();
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>();
3645

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

‎SYCL/ESIMD/lsc/lsc_surf_load_u32_stateless.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
//==----- lsc_surf_load_u32_stateless.cpp - DPC++ ESIMD on-device test --==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu-intel-pvc || esimd_emulator
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl -fsycl-esimd-force-stateless-mem %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include "lsc_surf_load_u32.cpp"

‎SYCL/ESIMD/lsc/lsc_surf_load_u64.cpp

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

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

2119
// non transpose
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);
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);
2826

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

3230
// transpose
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>();
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>();
3642

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

‎SYCL/ESIMD/lsc/lsc_surf_load_u64_stateless.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
//==----- lsc_surf_load_u64_stateless.cpp - DPC++ ESIMD on-device test --==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu-intel-pvc || esimd_emulator
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl -fsycl-esimd-force-stateless-mem %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include "lsc_surf_load_u64.cpp"

‎SYCL/ESIMD/lsc/lsc_surf_load_u8_u16.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,8 @@ 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>();
3234

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

‎SYCL/ESIMD/lsc/lsc_surf_load_u8_u16_stateless.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
//==----- lsc_surf_load_u8_u16_stateless.cpp - DPC++ ESIMD on-device test --==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu-intel-pvc || esimd_emulator
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl -fsycl-esimd-force-stateless-mem %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include "lsc_surf_load_u8_u16.cpp"

0 commit comments

Comments
 (0)