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

Commit b327f9d

Browse files
bso-intelbso-intel
committed
Merge remote-tracking branch 'upstream/intel' into filter
2 parents b4d842d + 5d8eee5 commit b327f9d

File tree

79 files changed

+560
-238
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

79 files changed

+560
-238
lines changed

‎SYCL/Basic/alloc_pinned_host_memory.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,7 @@ int main() {
3737
}
3838
}
3939

40+
// CHECK:---> piMemBufferCreate
4041
// CHECK:---> piMemBufferCreate
4142
// CHECK-NEXT: {{.*}} : {{.*}}
42-
// CHECK-NEXT: {{.*}} : 9
43+
// CHECK-NEXT: {{.*}} : 17

‎SYCL/Basic/event_profiling_info.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,5 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
22
//
3-
// Looks like there is a bug in the test. There are rare sporadic failures of
4-
// this test on different devices.
5-
// REQUIRES: TEMPORARY_DISABLED
6-
//
73
// RUN: %HOST_RUN_PLACEHOLDER %t.out
84
// RUN: %CPU_RUN_PLACEHOLDER %t.out
95
// RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -67,6 +63,8 @@ int main() {
6763
event kernelEvent = kernelQueue.submit([&](sycl::handler &CGH) {
6864
CGH.single_task<class EmptyKernel>([=]() {});
6965
});
66+
copyEvent.wait();
67+
kernelEvent.wait();
7068

7169
assert(verifyProfiling(copyEvent) && verifyProfiling(kernelEvent));
7270
}

‎SYCL/Basic/parallel_for_indexers.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6,10 +6,6 @@
66
// RUN: %GPU_RUN_PLACEHOLDER %t2.out
77
// RUN: %ACC_RUN_PLACEHOLDER %t2.out
88

9-
// TODO: Unexpected result
10-
// TODO: _indexers.cpp:37: int main(): Assertion `id == -1' failed.
11-
// XFAIL: level_zero&&gpu
12-
139
#include <CL/sycl.hpp>
1410

1511
#include <cassert>

‎SYCL/ESIMD/BitonicSortK.cpp

Lines changed: 12 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818
#include <iostream>
1919

2020
using namespace cl::sycl;
21-
using namespace sycl::INTEL::gpu;
21+
using namespace sycl::ext::intel::experimental::esimd;
2222
using namespace std;
2323

2424
#define LOG2_ELEMENTS 16 // 24
@@ -598,11 +598,11 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
598598
auto e = pQueue_->submit([&](handler &cgh) {
599599
auto acci = bufi.get_access<access::mode::read>(cgh);
600600
auto acco = bufo.get_access<access::mode::write>(cgh);
601-
cgh.parallel_for<class Sort256>(SortGlobalRange * SortLocalRange,
602-
[=](id<1> i) SYCL_ESIMD_KERNEL {
603-
using namespace sycl::INTEL::gpu;
604-
cmk_bitonic_sort_256(acci, acco, i);
605-
});
601+
cgh.parallel_for<class Sort256>(
602+
SortGlobalRange * SortLocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
603+
using namespace sycl::ext::intel::experimental::esimd;
604+
cmk_bitonic_sort_256(acci, acco, i);
605+
});
606606
});
607607
e.wait();
608608
total_time += esimd_test::report_time("kernel time", e, e);
@@ -638,11 +638,12 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
638638
buffer<uint32_t, 1> buf(pOutputs, range<1>(size));
639639
mergeEvent[k] = pQueue_->submit([&](handler &cgh) {
640640
auto acc = buf.get_access<access::mode::read_write>(cgh);
641-
cgh.parallel_for<class Merge>(MergeGlobalRange * MergeLocalRange,
642-
[=](id<1> tid) SYCL_ESIMD_KERNEL {
643-
using namespace sycl::INTEL::gpu;
644-
cmk_bitonic_merge(acc, j, i, tid);
645-
});
641+
cgh.parallel_for<class Merge>(
642+
MergeGlobalRange * MergeLocalRange,
643+
[=](id<1> tid) SYCL_ESIMD_KERNEL {
644+
using namespace sycl::ext::intel::experimental::esimd;
645+
cmk_bitonic_merge(acc, j, i, tid);
646+
});
646647
});
647648
k++;
648649
}

‎SYCL/ESIMD/BitonicSortKv2.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
#include <iostream>
2020

2121
using namespace cl::sycl;
22-
using namespace sycl::INTEL::gpu;
22+
using namespace sycl::ext::intel::experimental::esimd;
2323
using namespace std;
2424

2525
#define LOG2_ELEMENTS 16 // 24
@@ -517,7 +517,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
517517
auto e = pQueue_->submit([&](handler &cgh) {
518518
cgh.parallel_for<class Sort256>(
519519
SortGlobalRange * SortLocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
520-
using namespace sycl::INTEL::gpu;
520+
using namespace sycl::ext::intel::experimental::esimd;
521521
cmk_bitonic_sort_256(pInputs, pOutputs, i);
522522
});
523523
});
@@ -553,12 +553,12 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
553553
// locally.
554554
for (int j = i; j >= 8; j--) {
555555
mergeEvent[k] = pQueue_->submit([&](handler &cgh) {
556-
cgh.parallel_for<class Merge>(MergeGlobalRange * MergeLocalRange,
557-
[=](id<1> tid) SYCL_ESIMD_KERNEL {
558-
using namespace sycl::INTEL::gpu;
559-
cmk_bitonic_merge(pOutputs, j, i,
560-
tid);
561-
});
556+
cgh.parallel_for<class Merge>(
557+
MergeGlobalRange * MergeLocalRange,
558+
[=](id<1> tid) SYCL_ESIMD_KERNEL {
559+
using namespace sycl::ext::intel::experimental::esimd;
560+
cmk_bitonic_merge(pOutputs, j, i, tid);
561+
});
562562
});
563563
// mergeEvent[k].wait();
564564
k++;

‎SYCL/ESIMD/PrefixSum.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@
3838
#define REMAINING_ENTRIES 64
3939

4040
using namespace cl::sycl;
41-
using namespace sycl::INTEL::gpu;
41+
using namespace sycl::ext::intel::experimental::esimd;
4242

4343
void compute_local_prefixsum(unsigned int prefixSum[], unsigned int size,
4444
unsigned elem_stride, unsigned thread_stride) {

‎SYCL/ESIMD/Prefix_Local_sum1.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@
3636
#define MIN_NUM_THREADS 1
3737

3838
using namespace cl::sycl;
39-
using namespace sycl::INTEL::gpu;
39+
using namespace sycl::ext::intel::experimental::esimd;
4040

4141
void compute_local_prefixsum(unsigned int input[], unsigned int prefixSum[],
4242
unsigned int size) {

‎SYCL/ESIMD/Prefix_Local_sum2.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@
3636
#define MIN_NUM_THREADS 1
3737

3838
using namespace cl::sycl;
39-
using namespace sycl::INTEL::gpu;
39+
using namespace sycl::ext::intel::experimental::esimd;
4040

4141
void compute_local_prefixsum(unsigned int input[], unsigned int prefixSum[],
4242
unsigned int size) {

‎SYCL/ESIMD/Prefix_Local_sum3.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@
3838
#define REMAINING_ENTRIES 64
3939

4040
using namespace cl::sycl;
41-
using namespace sycl::INTEL::gpu;
41+
using namespace sycl::ext::intel::experimental::esimd;
4242

4343
void compute_local_prefixsum(unsigned int prefixSum[], unsigned int size,
4444
unsigned elem_stride, unsigned thread_stride,

‎SYCL/ESIMD/Stencil.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -107,7 +107,7 @@ int main(int argc, char *argv[]) {
107107
auto e = q.submit([&](handler &cgh) {
108108
cgh.parallel_for<class Stencil_kernel>(
109109
GlobalRange * LocalRange, [=](item<2> it) SYCL_ESIMD_KERNEL {
110-
using namespace sycl::INTEL::gpu;
110+
using namespace sycl::ext::intel::experimental::esimd;
111111
uint h_pos = it.get_id(0);
112112
uint v_pos = it.get_id(1);
113113

‎SYCL/ESIMD/accessor.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ int main() {
5050

5151
cgh.parallel_for<class Test>(
5252
range<1>(1), [=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
53-
using namespace sycl::INTEL::gpu;
53+
using namespace sycl::ext::intel::experimental::esimd;
5454
unsigned int offset = 0;
5555
for (int k = 0; k < VL / 16; k++) {
5656
simd<Ty, 16> var = block_load<Ty, 16>(acc0, offset);

‎SYCL/ESIMD/accessor_gather_scatter.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ template <typename T, unsigned VL, unsigned STRIDE> struct Kernel {
3232
Kernel(Acc<T> acc) : acc(acc) {}
3333

3434
void operator()(id<1> i) const SYCL_ESIMD_KERNEL {
35-
using namespace sycl::INTEL::gpu;
35+
using namespace sycl::ext::intel::experimental::esimd;
3636
uint32_t ii = static_cast<uint32_t>(i.get(0));
3737
// every STRIDE threads (subgroups with sg_size=1) access contiguous block
3838
// of STRIDE*VL elements

‎SYCL/ESIMD/accessor_load_store.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ template <typename T> struct Kernel {
3030
Kernel(Acc<T> acc) : acc(acc) {}
3131

3232
void operator()(id<1> i) const SYCL_ESIMD_KERNEL {
33-
using namespace sycl::INTEL::gpu;
33+
using namespace sycl::ext::intel::experimental::esimd;
3434
uint32_t ii = static_cast<uint32_t>(i.get(0));
3535
T v = scalar_load<T>(acc, ii);
3636
v += ii;

‎SYCL/ESIMD/dp4a.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -57,7 +57,7 @@ int main(void) {
5757
auto e = q.submit([&](handler &cgh) {
5858
cgh.parallel_for<class Test>(
5959
Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
60-
using namespace sycl::INTEL::gpu;
60+
using namespace sycl::ext::intel::experimental::esimd;
6161

6262
simd<DTYPE, SIZE> src0(0);
6363
src0 = block_load<DTYPE, SIZE>(S0);

‎SYCL/ESIMD/ext_math.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
#include <iostream>
2020

2121
using namespace cl::sycl;
22-
using namespace sycl::INTEL::gpu;
22+
using namespace sycl::ext::intel::experimental::esimd;
2323

2424
// --- Data initialization functions
2525

‎SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -60,7 +60,7 @@ int main(void) {
6060
auto qq = q.submit([&](handler &cgh) {
6161
cgh.parallel_for<KernelID>(
6262
sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL {
63-
using namespace sycl::INTEL::gpu;
63+
using namespace sycl::ext::intel::experimental::esimd;
6464

6565
simd<a_data_t, SIZE> va(0);
6666
simd<b_data_t, SIZE> vb(0);

‎SYCL/ESIMD/fp_call_from_func.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ class KernelID;
2727
ESIMD_NOINLINE int add(int A, int B) { return A + B; }
2828

2929
template <typename AccTy> ESIMD_NOINLINE void test(AccTy acc, int A, int B) {
30-
using namespace sycl::INTEL::gpu;
30+
using namespace sycl::ext::intel::experimental::esimd;
3131

3232
auto foo = &add;
3333
auto res = foo(A, B);

‎SYCL/ESIMD/fp_call_recursive.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -51,15 +51,15 @@ int main(int argc, char **argv) {
5151
q.submit([&](handler &cgh) {
5252
auto acc = buf.get_access<access::mode::write>(cgh);
5353

54-
cgh.parallel_for<KernelID>(sycl::range<1>{1},
55-
[=](id<1> i) SYCL_ESIMD_KERNEL {
56-
using namespace sycl::INTEL::gpu;
54+
cgh.parallel_for<KernelID>(
55+
sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL {
56+
using namespace sycl::ext::intel::experimental::esimd;
5757

58-
auto foo = &add;
59-
auto res = foo(in1, in2, in3);
58+
auto foo = &add;
59+
auto res = foo(in1, in2, in3);
6060

61-
scalar_store(acc, 0, res);
62-
});
61+
scalar_store(acc, 0, res);
62+
});
6363
});
6464
} catch (cl::sycl::exception const &e) {
6565
std::cout << "SYCL exception caught: " << e.what() << std::endl;

‎SYCL/ESIMD/fp_in_phi.cpp

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -50,21 +50,21 @@ bool test(queue q, bool flag) {
5050
auto o_acc = o_buf.get_access<access::mode::write>(cgh);
5151
auto y_acc = y_buf.get_access<access::mode::write>(cgh);
5252

53-
cgh.parallel_for<KernelID>(sycl::range<1>{1},
54-
[=](id<1> i) SYCL_ESIMD_KERNEL {
55-
using namespace sycl::INTEL::gpu;
56-
using f = int (*)(int);
53+
cgh.parallel_for<KernelID>(
54+
sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL {
55+
using namespace sycl::ext::intel::experimental::esimd;
56+
using f = int (*)(int);
5757

58-
f a[] = {f1, f2};
59-
if (flag) {
60-
a[0] = f3;
61-
scalar_store(y_acc, 0, 2);
62-
}
58+
f a[] = {f1, f2};
59+
if (flag) {
60+
a[0] = f3;
61+
scalar_store(y_acc, 0, 2);
62+
}
6363

64-
auto res = a[0](in1) + a[1](in2);
64+
auto res = a[0](in1) + a[1](in2);
6565

66-
scalar_store(o_acc, 0, res);
67-
});
66+
scalar_store(o_acc, 0, res);
67+
});
6868
});
6969
} catch (cl::sycl::exception const &e) {
7070
std::cout << "SYCL exception caught: " << e.what() << std::endl;

‎SYCL/ESIMD/fp_in_select.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -45,15 +45,15 @@ bool test(queue q, bool flag) {
4545
q.submit([&](handler &cgh) {
4646
auto acc = buf.get_access<access::mode::write>(cgh);
4747

48-
cgh.parallel_for<KernelID>(sycl::range<1>{1},
49-
[=](id<1> i) SYCL_ESIMD_KERNEL {
50-
using namespace sycl::INTEL::gpu;
48+
cgh.parallel_for<KernelID>(
49+
sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL {
50+
using namespace sycl::ext::intel::experimental::esimd;
5151

52-
auto foo = flag ? &add : &sub;
53-
auto res = foo(in1, in2);
52+
auto foo = flag ? &add : &sub;
53+
auto res = foo(in1, in2);
5454

55-
scalar_store(acc, 0, res);
56-
});
55+
scalar_store(acc, 0, res);
56+
});
5757
});
5858
} catch (cl::sycl::exception const &e) {
5959
std::cout << "SYCL exception caught: " << e.what() << std::endl;

‎SYCL/ESIMD/histogram.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -153,7 +153,7 @@ int main(int argc, char *argv[]) {
153153

154154
cgh.parallel_for<class Hist>(
155155
Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
156-
using namespace sycl::INTEL::gpu;
156+
using namespace sycl::ext::intel::experimental::esimd;
157157

158158
// Get thread origin offsets
159159
uint tid = ndi.get_group(0);

‎SYCL/ESIMD/histogram_256_slm.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ static constexpr int BLOCK_WIDTH = 32;
2222
static constexpr int NUM_BLOCKS = 32;
2323

2424
using namespace cl::sycl;
25-
using namespace sycl::INTEL::gpu;
25+
using namespace sycl::ext::intel::experimental::esimd;
2626

2727
// Histogram kernel: computes the distribution of pixel intensities
2828
ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,

‎SYCL/ESIMD/histogram_256_slm_spec.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ static constexpr int BLOCK_WIDTH = 32;
2323
static constexpr int NUM_BLOCKS = 32;
2424

2525
using namespace cl::sycl;
26-
using namespace sycl::INTEL::gpu;
26+
using namespace sycl::ext::intel::experimental::esimd;
2727

2828
// Histogram kernel: computes the distribution of pixel intensities
2929
ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,

‎SYCL/ESIMD/histogram_2d.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,7 @@ int main(int argc, char *argv[]) {
148148

149149
cgh.parallel_for<class Hist>(
150150
Range, [=](nd_item<2> ndi) SYCL_ESIMD_KERNEL {
151-
using namespace sycl::INTEL::gpu;
151+
using namespace sycl::ext::intel::experimental::esimd;
152152

153153
// Get thread origin offsets
154154
uint h_pos = ndi.get_group(0) * BLOCK_WIDTH;

‎SYCL/ESIMD/histogram_raw_send.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ int checkHistogram(unsigned int *refHistogram, unsigned int *hist) {
6262
return 1;
6363
}
6464

65-
using namespace sycl::INTEL::gpu;
65+
using namespace sycl::ext::intel::experimental::esimd;
6666
template <EsimdAtomicOpType Op, typename T, int n>
6767
ESIMD_INLINE void atomic_write(T *bins, simd<unsigned, n> offset,
6868
simd<T, n> src0, simd<ushort, n> pred) {

‎SYCL/ESIMD/kmeans/kmeans.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@
2323
#include <vector>
2424

2525
using namespace cl::sycl;
26-
using namespace sycl::INTEL::gpu;
26+
using namespace sycl::ext::intel::experimental::esimd;
2727
using namespace std;
2828

2929
inline float dist(Point p, Centroid c) {

‎SYCL/ESIMD/linear/linear.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -87,7 +87,7 @@ int main(int argc, char *argv[]) {
8787

8888
cgh.parallel_for<class Test>(
8989
GlobalRange * LocalRange, [=](item<2> it) SYCL_ESIMD_KERNEL {
90-
using namespace sycl::INTEL::gpu;
90+
using namespace sycl::ext::intel::experimental::esimd;
9191

9292
simd<unsigned char, 8 * 32> vin;
9393
auto in = vin.format<unsigned char, 8, 32>();

‎SYCL/ESIMD/mandelbrot/mandelbrot.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818
#include <memory>
1919

2020
using namespace cl::sycl;
21-
using namespace sycl::INTEL::gpu;
21+
using namespace sycl::ext::intel::experimental::esimd;
2222

2323
#ifdef _SIM_MODE_
2424
#define CRUNCH 32

‎SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
#include <memory>
2020

2121
using namespace cl::sycl;
22-
using namespace sycl::INTEL::gpu;
22+
using namespace sycl::ext::intel::experimental::esimd;
2323

2424
#ifdef _SIM_MODE_
2525
#define CRUNCH 32

‎SYCL/ESIMD/matrix_transpose.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818

1919
using namespace cl::sycl;
2020
using namespace std;
21-
using namespace sycl::INTEL::gpu;
21+
using namespace sycl::ext::intel::experimental::esimd;
2222

2323
void initMatrix(int *M, unsigned N) {
2424
assert(N >= 8 && (((N - 1) & N) == 0) &&

0 commit comments

Comments
 (0)