Skip to content

Commit f286ab7

Browse files
yuxianchyuxianch
authored
Merge pull request intel#49 from bb-sycl/syclos
Auto pulldown and update tc files for syclos branch on 20210629
2 parents 3880c3a + 0471f0b commit f286ab7

11 files changed

+274
-5
lines changed

‎SYCL/ESIMD/api/simd_binop_integer_promotion.cpp

Lines changed: 118 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,118 @@
1+
//==---- simd_binop_integer_promotion.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
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
//
13+
// The test checks that ESIMD binary operation APIs honor integer promotion
14+
// rules. E.g.:
15+
// simd <short, 32> a;
16+
// simd <short, 32> b;
17+
// a + b; // yield simd <int, 32>
18+
19+
#include "../esimd_test_utils.hpp"
20+
21+
#include <CL/sycl.hpp>
22+
#include <sycl/ext/intel/experimental/esimd.hpp>
23+
24+
#include <iostream>
25+
#include <limits>
26+
27+
using namespace cl::sycl;
28+
using namespace sycl::ext::intel::experimental::esimd;
29+
30+
template <typename T> struct KernelName;
31+
32+
// The main test routine.
33+
template <typename T> bool test(queue q) {
34+
constexpr unsigned VL = 16;
35+
36+
T A[VL];
37+
T B[VL];
38+
using T_promoted = decltype(T{} + T{});
39+
T_promoted C[VL];
40+
41+
std::cout << "Testing " << typeid(T).name() << " + " << typeid(T).name()
42+
<< " => " << typeid(T_promoted).name() << " ...\n";
43+
44+
for (unsigned i = 0; i < VL; ++i) {
45+
A[i] = i;
46+
B[i] = 1;
47+
}
48+
T maxNum = std::numeric_limits<T>::max();
49+
// test overflow in one of the lanes
50+
A[VL / 2] = maxNum;
51+
B[VL / 2] = maxNum;
52+
53+
try {
54+
buffer<T, 1> bufA(A, range<1>(VL));
55+
buffer<T, 1> bufB(B, range<1>(VL));
56+
buffer<T_promoted, 1> bufC(C, range<1>(VL));
57+
range<1> glob_range{1};
58+
59+
auto e = q.submit([&](handler &cgh) {
60+
auto PA = bufA.template get_access<access::mode::read>(cgh);
61+
auto PB = bufB.template get_access<access::mode::read>(cgh);
62+
auto PC = bufC.template get_access<access::mode::write>(cgh);
63+
cgh.parallel_for<KernelName<T>>(
64+
glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
65+
using namespace sycl::ext::intel::experimental::esimd;
66+
unsigned int offset = i * VL * sizeof(T);
67+
simd<T, VL> va;
68+
va.copy_from(PA, offset);
69+
simd<T, VL> vb;
70+
vb.copy_from(PB, offset);
71+
auto vc = va + vb;
72+
unsigned int offsetC = i * VL * sizeof(T_promoted);
73+
vc.copy_to(PC, offsetC);
74+
});
75+
});
76+
q.wait_and_throw();
77+
} catch (cl::sycl::exception const &e) {
78+
std::cout << "SYCL exception caught: " << e.what() << '\n';
79+
return e.get_cl_code();
80+
}
81+
82+
int err_cnt = 0;
83+
84+
for (unsigned i = 0; i < VL; ++i) {
85+
T_promoted gold = (T_promoted)(A[i] + B[i]);
86+
T_promoted val = C[i];
87+
88+
if (val != gold) {
89+
if (++err_cnt < 10) {
90+
std::cout << "failed at index " << i << ": " << val << " != " << gold
91+
<< " (gold)\n";
92+
}
93+
}
94+
}
95+
if (err_cnt > 0) {
96+
std::cout << " pass rate: " << ((float)(VL - err_cnt) / (float)VL) * 100.0f
97+
<< "% (" << (VL - err_cnt) << "/" << VL << ")\n";
98+
}
99+
100+
std::cout << (err_cnt > 0 ? " FAILED\n" : " Passed\n");
101+
return err_cnt > 0 ? false : true;
102+
}
103+
104+
int main(int argc, char **argv) {
105+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
106+
107+
auto dev = q.get_device();
108+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
109+
110+
bool passed = true;
111+
passed &= test<unsigned short>(q);
112+
passed &= test<short>(q);
113+
passed &= test<unsigned char>(q);
114+
passed &= test<char>(q);
115+
116+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
117+
return passed ? 0 : 1;
118+
}

‎SYCL/ESIMD/api/simd_subscript_operator.cpp

Lines changed: 98 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,98 @@
1+
//==----- simd_subscript_operator.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
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx -fsycl %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
//
13+
// The test checks that it's possible to write through the simd subscript
14+
// operator. E.g.:
15+
// simd<int, 4> v = 1;
16+
// v[1] = 0; // v[1] returns writable simd_view
17+
18+
#include "../esimd_test_utils.hpp"
19+
20+
#include <CL/sycl.hpp>
21+
#include <sycl/ext/intel/experimental/esimd.hpp>
22+
23+
#include <iostream>
24+
25+
using namespace cl::sycl;
26+
using namespace sycl::ext::intel::experimental::esimd;
27+
28+
int main(int argc, char **argv) {
29+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
30+
31+
auto dev = q.get_device();
32+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
33+
34+
constexpr unsigned VL = 16;
35+
36+
int A[VL];
37+
int B[VL];
38+
int gold[VL];
39+
40+
for (unsigned i = 0; i < VL; ++i) {
41+
A[i] = -i;
42+
B[i] = i;
43+
gold[i] = B[i];
44+
}
45+
46+
// some random indices to overwrite elements in B with elements from A.
47+
std::array<int, 5> indicesToCopy = {2, 5, 9, 10, 13};
48+
49+
try {
50+
buffer<int, 1> bufA(A, range<1>(VL));
51+
buffer<int, 1> bufB(B, range<1>(VL));
52+
range<1> glob_range{1};
53+
54+
auto e = q.submit([&](handler &cgh) {
55+
auto PA = bufA.get_access<access::mode::read>(cgh);
56+
auto PB = bufB.template get_access<access::mode::read_write>(cgh);
57+
cgh.parallel_for<class Test>(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL {
58+
using namespace sycl::ext::intel::experimental::esimd;
59+
unsigned int offset = i * VL * sizeof(int);
60+
simd<int, VL> va;
61+
va.copy_from(PA, offset);
62+
simd<int, VL> vb;
63+
vb.copy_from(PB, offset);
64+
for (auto idx : indicesToCopy)
65+
vb[idx] = va[idx];
66+
vb.copy_to(PB, offset);
67+
});
68+
});
69+
q.wait_and_throw();
70+
} catch (cl::sycl::exception const &e) {
71+
std::cout << "SYCL exception caught: " << e.what() << '\n';
72+
return e.get_cl_code();
73+
}
74+
75+
int err_cnt = 0;
76+
77+
for (auto i : indicesToCopy)
78+
gold[i] = A[i];
79+
80+
for (unsigned i = 0; i < VL; ++i) {
81+
int val = B[i];
82+
83+
if (val != gold[i]) {
84+
if (++err_cnt < 10) {
85+
std::cout << "failed at index " << i << ": " << val << " != " << gold[i]
86+
<< " (gold)\n";
87+
}
88+
}
89+
}
90+
if (err_cnt > 0) {
91+
std::cout << " pass rate: " << ((float)(VL - err_cnt) / (float)VL) * 100.0f
92+
<< "% (" << (VL - err_cnt) << "/" << VL << ")\n";
93+
}
94+
95+
std::cout << (err_cnt > 0 ? " FAILED\n" : " Passed\n");
96+
97+
return err_cnt > 0 ? 1 : 0;
98+
}

‎SYCL/ESIMD/fp_call_from_func.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,9 @@
1212
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1313
// UNSUPPORTED: cuda
1414
//
15+
// The test hangs after driver update to 21.23.20043
16+
// REQUIRES: TEMPORARY_DISABLE
17+
//
1518
// The test checks that ESIMD kernels support use of function pointers from
1619
// within other functions.
1720

‎SYCL/OnlineCompiler/online_compiler_L0.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,8 @@
66
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %th.out
77
// RUN: %HOST_RUN_PLACEHOLDER %th.out
88

9+
// The test regressed with GPU 21.23.20043. The fix is coming in next driver.
10+
// XFAIL: linux
911
// This test checks INTEL feature class online_compiler for Level-Zero.
1012
// All Level-Zero specific code is kept here and the common part that can be
1113
// re-used by other backends is kept in online_compiler_common.hpp file.

‎SYCL/Plugin/level-zero-event-leak.cpp

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
// REQUIRES: level_zero, level_zero_dev_kit
2+
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
3+
// RUN: env SYCL_DEVICE_FILTER=level_zero ZE_DEBUG=4 %GPU_RUN_PLACEHOLDER %t.out
4+
//
5+
// CHECK-NOT: ---> LEAK
6+
7+
// The test is to check that there are no leaks reported with the embedded
8+
// ZE_DEBUG=4 testing capability. Example of a leak reported is this:
9+
//
10+
// clang-format off
11+
// ZE_DEBUG=4: check balance of create/destroy calls
12+
// ----------------------------------------------------------
13+
// zeContextCreate = 1 \---> zeContextDestroy = 1
14+
// zeCommandQueueCreate = 1 \---> zeCommandQueueDestroy = 0 ---> LEAK = 1
15+
// zeModuleCreate = 1 \---> zeModuleDestroy = 0 ---> LEAK = 1
16+
// zeKernelCreate = 1 \---> zeKernelDestroy = 0 ---> LEAK = 1
17+
// zeEventPoolCreate = 1 \---> zeEventPoolDestroy = 1
18+
// zeCommandListCreateImmediate = 1 |
19+
// zeCommandListCreate = 2 \---> zeCommandListDestroy = 2 ---> LEAK = 1
20+
// zeEventCreate = 129 \---> zeEventDestroy = 1 ---> LEAK = 128
21+
// zeFenceCreate = 2 \---> zeFenceDestroy = 0 ---> LEAK = 2
22+
// zeImageCreate = 0 \---> zeImageDestroy = 0
23+
// zeSamplerCreate = 0 \---> zeSamplerDestroy = 0
24+
// zeMemAllocDevice = 0 |
25+
// zeMemAllocHost = 0 |
26+
// zeMemAllocShared = 0 \---> zeMemFree = 0
27+
//
28+
// clang-format on
29+
//
30+
// NOTE: The 1000 value below is to be larger than the "128" heuristic in
31+
// queue_impl::addSharedEvent.
32+
33+
#include <CL/sycl.hpp>
34+
using namespace cl;
35+
int main(int argc, char **argv) {
36+
sycl::queue Q;
37+
const unsigned n_chunk = 1000;
38+
for (int i = 0; i < n_chunk; i++)
39+
Q.single_task([=]() {});
40+
Q.wait();
41+
return 0;
42+
}

‎SYCL/USM/badmalloc.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,5 @@
11
// UNSUPPORTED: windows
22
//
3-
// XFAIL: level_zero
4-
// TODO: enable the test when L0 driver is honoring too large sizes.
5-
//
63
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t1.out
74
// RUN: %HOST_RUN_PLACEHOLDER %t1.out
85
// RUN: %CPU_RUN_PLACEHOLDER %t1.out

‎config_sycl/esimd_api_simd_binop_integer_promotion.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/ESIMD/api/simd_binop_integer_promotion.cpp

‎config_sycl/esimd_api_simd_subscript_operator.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/ESIMD/api/simd_subscript_operator.cpp

‎config_sycl/plugin_level_zero_event_leak.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/Plugin/level-zero-event-leak.cpp

‎llvm_test_suite_sycl.xml

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
<description>Port of llvm_test_suite_sycl.
44
Suite is autogenerated by suite_generator_sycl.pl that you can find in the root dir of suite
55
Sources repo git-amr-2.devtools.intel.com/gerrit/icl_tst-llvm-project-llvm-test-suite
6-
Last Changed Revision: 9b15e6c6b34315f7ff2e132d2712c92a98eeeaa8 Tue Jun 22 23:23:28 2021 -0700
6+
Last Changed Revision: de2cb00703b39cc8cd390c95a1deef6e02fdbb69 Tue Jun 29 07:53:25 2021 -0700
77
</description>
88
<files>
99
<file path="cmake" />
@@ -192,7 +192,9 @@ Last Changed Revision: 9b15e6c6b34315f7ff2e132d2712c92a98eeeaa8 Tue Jun 22 23
192192
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_accessor_gather_scatter" />
193193
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_accessor_load_store" />
194194
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_api_esimd_bit_ops" />
195+
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_api_simd_binop_integer_promotion" />
195196
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_api_simd_memory_access" />
197+
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_api_simd_subscript_operator" />
196198
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_bitonicsortk" />
197199
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_bitonicsortkv2" />
198200
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="esimd_dp4a" />
@@ -376,6 +378,7 @@ Last Changed Revision: 9b15e6c6b34315f7ff2e132d2712c92a98eeeaa8 Tue Jun 22 23
376378
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_interop_level_zero_interop_task_mem" />
377379
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_interop_level_zero_keep_ownership" />
378380
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_interop_opencl" />
381+
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_level_zero_event_leak" />
379382
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_max_malloc" />
380383
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_retain_events" />
381384
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl.xml" testName="plugin_sycl_ls" />

‎llvm_test_suite_sycl_valgrind.xml

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
<description>Port of llvm_test_suite_sycl.
44
Suite is autogenerated by suite_generator_sycl.pl that you can find in the root dir of suite
55
Sources repo git-amr-2.devtools.intel.com/gerrit/icl_tst-llvm-project-llvm-test-suite
6-
Last Changed Revision: 9b15e6c6b34315f7ff2e132d2712c92a98eeeaa8 Tue Jun 22 23:23:28 2021 -0700
6+
Last Changed Revision: de2cb00703b39cc8cd390c95a1deef6e02fdbb69 Tue Jun 29 07:53:25 2021 -0700
77
</description>
88
<files>
99
<file path="cmake" />
@@ -192,7 +192,9 @@ Last Changed Revision: 9b15e6c6b34315f7ff2e132d2712c92a98eeeaa8 Tue Jun 22 23
192192
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_accessor_gather_scatter" />
193193
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_accessor_load_store" />
194194
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_api_esimd_bit_ops" />
195+
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_api_simd_binop_integer_promotion" />
195196
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_api_simd_memory_access" />
197+
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_api_simd_subscript_operator" />
196198
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_bitonicsortk" />
197199
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_bitonicsortkv2" />
198200
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="esimd_dp4a" />
@@ -376,6 +378,7 @@ Last Changed Revision: 9b15e6c6b34315f7ff2e132d2712c92a98eeeaa8 Tue Jun 22 23
376378
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_interop_level_zero_interop_task_mem" />
377379
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_interop_level_zero_keep_ownership" />
378380
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_interop_opencl" />
381+
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_level_zero_event_leak" />
379382
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_max_malloc" />
380383
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_retain_events" />
381384
<test configFile="config_sycl/TEMPLATE_llvm_test_suite_sycl_valgrind.xml" testName="plugin_sycl_ls" />

0 commit comments

Comments
 (0)