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

Commit 7e75a62

Browse files
bso-intelbso-intel
committed
Merge remote-tracking branch 'upstream/intel' into filter
2 parents 40af33f + 30f2cc8 commit 7e75a62

File tree

9 files changed

+387
-39
lines changed

9 files changed

+387
-39
lines changed

‎SYCL/Basic/parallel_for_range.cpp

Lines changed: 0 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -78,42 +78,6 @@ int main() {
7878
<< std::endl;
7979
return 1;
8080
}
81-
82-
// parallel_for, (16, 16, 16) global, null local, reqd_wg_size(4, 4, 4) //
83-
// -> fail
84-
try {
85-
Q.submit([&](handler &CGH) {
86-
CGH.parallel_for<class ReqdWGSizeNegativeB>(
87-
range<3>(16, 16, 16), [=](item<3>) { reqd_wg_size_helper(); });
88-
});
89-
Q.wait_and_throw();
90-
std::cerr
91-
<< "Test case ReqdWGSizeNegativeB failed: no exception has been "
92-
"thrown\n";
93-
return 1; // We shouldn't be here, exception is expected
94-
} catch (nd_range_error &E) {
95-
if (string_class(E.what()).find("OpenCL 1.x and 2.0 requires to pass "
96-
"local size argument even if "
97-
"required work-group size was "
98-
"specified in the program source") ==
99-
string_class::npos) {
100-
std::cerr
101-
<< "Test case ReqdWGSizeNegativeB failed: unexpected exception: "
102-
<< E.what() << std::endl;
103-
return 1;
104-
}
105-
} catch (runtime_error &E) {
106-
std::cerr
107-
<< "Test case ReqdWGSizeNegativeB failed: unexpected exception: "
108-
<< E.what() << std::endl;
109-
return 1;
110-
} catch (...) {
111-
std::cerr
112-
<< "Test case ReqdWGSizeNegativeB failed: something unexpected "
113-
"has been caught"
114-
<< std::endl;
115-
return 1;
116-
}
11781
}
11882

11983
// Positive test-cases that should pass on any underlying OpenCL runtime

‎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/ESIMD/regression/noinline_bypointers_vadd.cpp

Lines changed: 112 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,112 @@
1+
//==------ noinline_bypointers_vadd.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+
// XFAIL: *
13+
14+
// Vector BE bug.
15+
// Memory access via a pointer type argument is broken if the function was not
16+
// inlined.
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+
using ptr = float *;
29+
static inline constexpr unsigned VL = 32;
30+
31+
SYCL_EXTERNAL ESIMD_NOINLINE void do_add(ptr A, float *B,
32+
ptr C) SYCL_ESIMD_FUNCTION {
33+
simd<float, VL> va;
34+
va.copy_from(A);
35+
simd<float, VL> vb;
36+
vb.copy_from(B);
37+
simd<float, VL> vc = va + vb;
38+
vc.copy_to(C);
39+
}
40+
41+
int main(void) {
42+
constexpr unsigned Size = 1024;
43+
constexpr unsigned GroupSize = 8;
44+
45+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
46+
47+
auto dev = q.get_device();
48+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
49+
auto ctxt = q.get_context();
50+
float *A =
51+
static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
52+
float *B =
53+
static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
54+
float *C =
55+
static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
56+
57+
for (unsigned i = 0; i < Size; ++i) {
58+
A[i] = B[i] = i;
59+
}
60+
61+
// We need that many workitems. Each processes VL elements of data.
62+
cl::sycl::range<1> GlobalRange{Size / VL};
63+
// Number of workitems in each workgroup.
64+
cl::sycl::range<1> LocalRange{GroupSize};
65+
66+
cl::sycl::nd_range<1> Range(GlobalRange, LocalRange);
67+
68+
try {
69+
auto e = q.submit([&](handler &cgh) {
70+
cgh.parallel_for<class Test>(
71+
Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
72+
int i = ndi.get_global_id(0);
73+
do_add(ptr{A + i * VL}, B + i * VL, ptr{C + i * VL});
74+
});
75+
});
76+
e.wait();
77+
} catch (cl::sycl::exception const &e) {
78+
std::cout << "SYCL exception caught: " << e.what() << '\n';
79+
80+
free(A, ctxt);
81+
free(B, ctxt);
82+
free(C, ctxt);
83+
84+
return e.get_cl_code();
85+
}
86+
87+
int err_cnt = 0;
88+
89+
for (unsigned i = 0; i < Size; ++i) {
90+
if (A[i] + B[i] != C[i]) {
91+
if (++err_cnt < 10) {
92+
std::cout << "failed at index " << i << ", " << C[i] << " != " << A[i]
93+
<< " + " << B[i] << "\n";
94+
}
95+
}
96+
}
97+
98+
free(A, ctxt);
99+
free(B, ctxt);
100+
free(C, ctxt);
101+
102+
if (err_cnt > 0) {
103+
std::cout << " pass rate: "
104+
<< ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
105+
<< (Size - err_cnt) << "/" << Size << ")\n";
106+
std::cout << "FAILED\n";
107+
return 1;
108+
}
109+
110+
std::cout << "Passed\n";
111+
return 0;
112+
}

‎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.

0 commit comments

Comments
 (0)