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

Commit 6e7b3ce

Browse files
sndmitrievsndmitriev
authored
[SYCL][ESIMD] Add tests for dpas API (#866)
Signed-off-by: Sergey Dmitriev <[email protected]>
1 parent 78c3d9b commit 6e7b3ce

File tree

3 files changed

+225
-0
lines changed

3 files changed

+225
-0
lines changed

‎SYCL/ESIMD/dpas/dpas_test1.cpp

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
//==--------------- dpas_test1.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-dg2
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include "../esimd_test_utils.hpp"
14+
15+
#include <CL/sycl.hpp>
16+
#include <iostream>
17+
#include <sycl/ext/intel/experimental/esimd.hpp>
18+
19+
using namespace cl::sycl;
20+
21+
int main(void) {
22+
constexpr unsigned Size = 64;
23+
constexpr unsigned VL = 16;
24+
constexpr unsigned GroupSize = 1;
25+
26+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
27+
28+
auto dev = q.get_device();
29+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
30+
31+
int *C = malloc_shared<int>(Size, q);
32+
memset(C, 0, Size * sizeof(int));
33+
34+
// We need that many task groups
35+
range<1> GroupRange{GroupSize};
36+
37+
// We need that many tasks in each group
38+
range<1> TaskRange{GroupSize};
39+
nd_range<1> Range(GroupRange, TaskRange);
40+
41+
q.submit([&](handler &cgh) {
42+
cgh.parallel_for<class Test>(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
43+
using namespace sycl::ext::intel::experimental::esimd;
44+
45+
simd<char, Size * 2> va(0);
46+
auto ma = va.bit_cast_view<char, 8, 16>();
47+
ma.select<2, 1, 4, 4>(0, 0) = 4;
48+
49+
simd<char, 8 * 16> vb(0);
50+
auto mb = vb.bit_cast_view<char, 8, 16>();
51+
mb.select<8, 1, 1, 1>(0, 0) = 4;
52+
53+
simd<int, Size> vc(0);
54+
vc =
55+
dpas<argument_type::S2, argument_type::S2, 8, 8, int, int, int, Size,
56+
32, 32>(vc, ma.bit_cast_view<int>(), mb.bit_cast_view<int>());
57+
58+
for (int i = 0; i < Size; i += VL) {
59+
simd<int, VL> output = vc.select<VL, 1>(i);
60+
output.copy_to(C + i);
61+
}
62+
});
63+
}).wait();
64+
65+
int err_cnt = 0;
66+
for (unsigned i = 0; i < Size && err_cnt < 10; ++i)
67+
if (C[i] != 1) {
68+
err_cnt++;
69+
std::cerr << "Failed at index " << i << ", " << C[i] << " != 1\n";
70+
}
71+
72+
free(C, q);
73+
std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
74+
return err_cnt > 0 ? 1 : 0;
75+
}

‎SYCL/ESIMD/dpas/dpas_test2.cpp

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
//==---------------- dpas_test2.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
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl -DESIMD_XE_HPC %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include "../esimd_test_utils.hpp"
14+
15+
#include <CL/sycl.hpp>
16+
#include <iostream>
17+
#include <sycl/ext/intel/experimental/esimd.hpp>
18+
19+
using namespace cl::sycl;
20+
21+
int main(void) {
22+
constexpr unsigned Size = 128;
23+
constexpr unsigned VL = 16;
24+
constexpr unsigned GroupSize = 1;
25+
26+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
27+
28+
auto dev = q.get_device();
29+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
30+
31+
int *C = malloc_shared<int>(Size, q);
32+
memset(C, 0, Size * sizeof(int));
33+
34+
// We need that many task groups
35+
range<1> GroupRange{GroupSize};
36+
37+
// We need that many tasks in each group
38+
range<1> TaskRange{GroupSize};
39+
nd_range<1> Range(GroupRange, TaskRange);
40+
41+
q.submit([&](handler &cgh) {
42+
cgh.parallel_for<class Test>(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
43+
using namespace sycl::ext::intel::experimental::esimd;
44+
45+
simd<char, Size * 2> va(0);
46+
auto ma = va.bit_cast_view<char, 8, 32>();
47+
ma.select<2, 1, 8, 4>(0, 0) = 4;
48+
49+
simd<char, Size> vb(0);
50+
auto mb = vb.bit_cast_view<char, 8, 16>();
51+
mb.select<8, 1, 1, 1>(0, 0) = 4;
52+
53+
simd<int, Size> vc(0);
54+
vc =
55+
dpas<argument_type::S2, argument_type::S2, 8, 8, int, int, int, Size,
56+
64, 32>(vc, ma.bit_cast_view<int>(), mb.bit_cast_view<int>());
57+
58+
for (int i = 0; i < Size; i += VL) {
59+
simd<int, VL> output = vc.select<VL, 1>(i);
60+
output.copy_to(C + i);
61+
}
62+
});
63+
}).wait();
64+
65+
int err_cnt = 0;
66+
for (unsigned i = 0; i < Size && err_cnt < 10; ++i)
67+
if (C[i] != 1) {
68+
err_cnt++;
69+
std::cerr << "Failed at index " << i << ", " << C[i] << " != 1\n";
70+
}
71+
72+
free(C, q);
73+
std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
74+
return err_cnt > 0 ? 1 : 0;
75+
}

‎SYCL/ESIMD/dpas/dpasw_test.cpp

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
//==------------- dpasw_test.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-dg2
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include "../esimd_test_utils.hpp"
14+
15+
#include <CL/sycl.hpp>
16+
#include <iostream>
17+
#include <sycl/ext/intel/experimental/esimd.hpp>
18+
19+
using namespace cl::sycl;
20+
21+
int main(void) {
22+
constexpr unsigned Size = 64;
23+
constexpr unsigned VL = 16;
24+
constexpr unsigned GroupSize = 2;
25+
26+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
27+
28+
auto dev = q.get_device();
29+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
30+
31+
int *C = malloc_shared<int>(Size, q);
32+
memset(C, 0, Size * sizeof(int));
33+
34+
// We need that many task groups
35+
range<1> GroupRange{GroupSize};
36+
37+
// We need that many tasks in each group
38+
range<1> TaskRange{GroupSize};
39+
nd_range<1> Range(GroupRange, TaskRange);
40+
41+
q.submit([&](handler &cgh) {
42+
cgh.parallel_for<class Test>(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
43+
using namespace sycl::ext::intel::experimental::esimd;
44+
45+
simd<char, Size * 2> va(0);
46+
auto ma = va.bit_cast_view<char, 8, 16>();
47+
ma.select<2, 1, 4, 4>(0, 0) = 4;
48+
49+
simd<char, 8 * 8> vb(0);
50+
auto mb = vb.bit_cast_view<char, 8, 8>();
51+
mb.select<4, 2, 1, 1>(0, 0) = 4;
52+
53+
simd<int, Size> vc(0);
54+
vc = dpasw<argument_type::S2, argument_type::S2, 8, 8, int, int, int,
55+
Size, 32, 16>(vc, ma.bit_cast_view<int>(),
56+
mb.bit_cast_view<int>());
57+
58+
for (int i = 0; i < Size; i += VL) {
59+
simd<int, VL> output = vc.select<VL, 1>(i);
60+
output.copy_to(C + i);
61+
}
62+
});
63+
}).wait();
64+
65+
int err_cnt = 0;
66+
for (unsigned i = 0; i < Size && err_cnt < 10; ++i)
67+
if (C[i] != 1) {
68+
err_cnt++;
69+
std::cerr << "Failed at index " << i << ", " << C[i] << " != 1\n";
70+
}
71+
72+
free(C, q);
73+
std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
74+
return err_cnt > 0 ? 1 : 0;
75+
}

0 commit comments

Comments
 (0)