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

Commit b2897f9

Browse files
sndmitrievsndmitriev
authored
[SYCL][ESIMD] Move some ESIMD APIs outside of experimental namespace (#892)
* [SYCL][ESIMD] Move some ESIMD APIs outside of experimental namespace Signed-off-by: Sergey Dmitriev <[email protected]> * Fix failures on precommit testing
1 parent f763b7f commit b2897f9

File tree

110 files changed

+333
-319
lines changed

Some content is hidden

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

110 files changed

+333
-319
lines changed

‎SYCL/DeprecatedFeatures/ESIMD/Inputs/spec_const_common.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@
1212
#include "esimd_test_utils.hpp"
1313

1414
#include <CL/sycl.hpp>
15-
#include <sycl/ext/intel/experimental/esimd.hpp>
15+
#include <sycl/ext/intel/esimd.hpp>
1616

1717
#include <iostream>
1818
#include <vector>
@@ -21,7 +21,7 @@ using namespace cl::sycl;
2121

2222
template <typename AccessorTy>
2323
ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) {
24-
using namespace sycl::ext::intel::experimental::esimd;
24+
using namespace sycl::ext::intel::esimd;
2525
// scatter function, that is used in scalar_store, can only process types
2626
// whose size is no more than 4 bytes.
2727
#if (STORE == 0)

‎SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -14,16 +14,16 @@
1414

1515
#include <CL/sycl.hpp>
1616
#include <iostream>
17-
#include <sycl/ext/intel/experimental/esimd.hpp>
17+
#include <sycl/ext/intel/esimd.hpp>
1818

1919
static constexpr int NUM_BINS = 256;
2020
static constexpr int SLM_SIZE = (NUM_BINS * 4);
2121
static constexpr int BLOCK_WIDTH = 32;
2222
static constexpr int NUM_BLOCKS = 32;
2323

2424
using namespace cl::sycl;
25-
using namespace sycl::ext::intel::experimental;
26-
using namespace sycl::ext::intel::experimental::esimd;
25+
using namespace sycl::ext::intel;
26+
using namespace sycl::ext::intel::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/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@
2222
#include "esimd_test_utils.hpp"
2323

2424
#include <CL/sycl.hpp>
25-
#include <sycl/ext/intel/experimental/esimd.hpp>
25+
#include <sycl/ext/intel/esimd.hpp>
2626

2727
#include <iostream>
2828
#include <vector>
@@ -76,8 +76,8 @@ int main(int argc, char **argv) {
7676
auto acc = buf.get_access<sycl::access::mode::write>(cgh);
7777
cgh.single_task<KernelAAA>(
7878
program.get_kernel<KernelAAA>(), [=]() SYCL_ESIMD_KERNEL {
79-
sycl::ext::intel::experimental::esimd::scalar_store(
80-
acc, i * sizeof(int), sc0.get() + sc1.get());
79+
sycl::ext::intel::esimd::scalar_store(acc, i * sizeof(int),
80+
sc0.get() + sc1.get());
8181
});
8282
});
8383
} catch (sycl::exception &e) {

‎SYCL/DeprecatedFeatures/program-merge-options.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#define SYCL2020_DISABLE_DEPRECATION_WARNINGS
22

33
#include <CL/sycl.hpp>
4-
#include <sycl/ext/intel/experimental/esimd.hpp>
4+
#include <sycl/ext/intel/esimd.hpp>
55

66
class KernelName;
77
void submitKernel() {

‎SYCL/ESIMD/BitonicSortK.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -15,10 +15,10 @@
1515
#include <CL/sycl.hpp>
1616
#include <algorithm>
1717
#include <iostream>
18-
#include <sycl/ext/intel/experimental/esimd.hpp>
18+
#include <sycl/ext/intel/esimd.hpp>
1919

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

2424
#define LOG2_ELEMENTS 16 // 24
@@ -617,7 +617,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
617617
auto acco = bufo.get_access<access::mode::write>(cgh);
618618
cgh.parallel_for<class Sort256>(
619619
SortGlobalRange * SortLocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
620-
using namespace sycl::ext::intel::experimental::esimd;
620+
using namespace sycl::ext::intel::esimd;
621621
cmk_bitonic_sort_256(acci, acco, i);
622622
});
623623
});
@@ -660,7 +660,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
660660
cgh.parallel_for<class Merge>(
661661
MergeGlobalRange * MergeLocalRange,
662662
[=](id<1> tid) SYCL_ESIMD_KERNEL {
663-
using namespace sycl::ext::intel::experimental::esimd;
663+
using namespace sycl::ext::intel::esimd;
664664
cmk_bitonic_merge(acc, j, i, tid);
665665
});
666666
});

‎SYCL/ESIMD/BitonicSortKv2.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -15,11 +15,11 @@
1515
#include <CL/sycl.hpp>
1616
#include <algorithm>
1717
#include <iostream>
18-
#include <sycl/ext/intel/experimental/esimd.hpp>
18+
#include <sycl/ext/intel/esimd.hpp>
1919

2020
using namespace cl::sycl;
21-
using namespace sycl::ext::intel::experimental;
22-
using namespace sycl::ext::intel::experimental::esimd;
21+
using namespace sycl::ext::intel;
22+
using namespace sycl::ext::intel::esimd;
2323
using namespace std;
2424

2525
#define LOG2_ELEMENTS 16 // 24
@@ -534,7 +534,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
534534
auto e = pQueue_->submit([&](handler &cgh) {
535535
cgh.parallel_for<class Sort256>(
536536
SortGlobalRange * SortLocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
537-
using namespace sycl::ext::intel::experimental::esimd;
537+
using namespace sycl::ext::intel::esimd;
538538
cmk_bitonic_sort_256(pInputs, pOutputs, i);
539539
});
540540
});
@@ -575,7 +575,7 @@ int BitonicSort::Solve(uint32_t *pInputs, uint32_t *pOutputs, uint32_t size) {
575575
cgh.parallel_for<class Merge>(
576576
MergeGlobalRange * MergeLocalRange,
577577
[=](id<1> tid) SYCL_ESIMD_KERNEL {
578-
using namespace sycl::ext::intel::experimental::esimd;
578+
using namespace sycl::ext::intel::esimd;
579579
cmk_bitonic_merge(pOutputs, j, i, tid);
580580
});
581581
});

‎SYCL/ESIMD/PrefixSum.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
#include <CL/sycl.hpp>
1616
#include <iostream>
17-
#include <sycl/ext/intel/experimental/esimd.hpp>
17+
#include <sycl/ext/intel/esimd.hpp>
1818

1919
#define MAX_TS_WIDTH 1024
2020
// kernel can handle TUPLE_SZ 1, 2, or 4
@@ -37,7 +37,7 @@
3737
#define REMAINING_ENTRIES 64
3838

3939
using namespace cl::sycl;
40-
using namespace sycl::ext::intel::experimental::esimd;
40+
using namespace sycl::ext::intel::esimd;
4141

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

‎SYCL/ESIMD/Prefix_Local_sum1.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
#include <CL/sycl.hpp>
1616
#include <iostream>
17-
#include <sycl/ext/intel/experimental/esimd.hpp>
17+
#include <sycl/ext/intel/esimd.hpp>
1818

1919
#define MAX_TS_WIDTH 1024
2020
// kernel can handle TUPLE_SZ 1, 2, or 4
@@ -35,7 +35,7 @@
3535
#define MIN_NUM_THREADS 1
3636

3737
using namespace cl::sycl;
38-
using namespace sycl::ext::intel::experimental::esimd;
38+
using namespace sycl::ext::intel::esimd;
3939

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

‎SYCL/ESIMD/Prefix_Local_sum2.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
#include <CL/sycl.hpp>
1616
#include <iostream>
17-
#include <sycl/ext/intel/experimental/esimd.hpp>
17+
#include <sycl/ext/intel/esimd.hpp>
1818

1919
#define MAX_TS_WIDTH 1024
2020
// kernel can handle TUPLE_SZ 1, 2, or 4
@@ -35,7 +35,7 @@
3535
#define MIN_NUM_THREADS 1
3636

3737
using namespace cl::sycl;
38-
using namespace sycl::ext::intel::experimental::esimd;
38+
using namespace sycl::ext::intel::esimd;
3939

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

‎SYCL/ESIMD/Prefix_Local_sum3.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
#include <CL/sycl.hpp>
1616
#include <iostream>
17-
#include <sycl/ext/intel/experimental/esimd.hpp>
17+
#include <sycl/ext/intel/esimd.hpp>
1818

1919
#define MAX_TS_WIDTH 1024
2020
// kernel can handle TUPLE_SZ 1, 2, or 4
@@ -37,7 +37,7 @@
3737
#define REMAINING_ENTRIES 64
3838

3939
using namespace cl::sycl;
40-
using namespace sycl::ext::intel::experimental::esimd;
40+
using namespace sycl::ext::intel::esimd;
4141

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

‎SYCL/ESIMD/Stencil.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
#include <CL/sycl.hpp>
1616
#include <iostream>
17-
#include <sycl/ext/intel/experimental/esimd.hpp>
17+
#include <sycl/ext/intel/esimd.hpp>
1818

1919
#define WIDTH 16
2020
#define HEIGHT 16
@@ -115,7 +115,7 @@ int main(int argc, char *argv[]) {
115115
auto e = q.submit([&](handler &cgh) {
116116
cgh.parallel_for<class Stencil_kernel>(
117117
GlobalRange * LocalRange, [=](item<2> it) SYCL_ESIMD_KERNEL {
118-
using namespace sycl::ext::intel::experimental::esimd;
118+
using namespace sycl::ext::intel::esimd;
119119
uint h_pos = it.get_id(0);
120120
uint v_pos = it.get_id(1);
121121

‎SYCL/ESIMD/accessor.cpp

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@
1515
#include "esimd_test_utils.hpp"
1616

1717
#include <CL/sycl.hpp>
18-
#include <sycl/ext/intel/experimental/esimd.hpp>
18+
#include <sycl/ext/intel/esimd.hpp>
1919

2020
#include <iostream>
2121

@@ -48,20 +48,20 @@ int main() {
4848
auto acc0 = buf0.get_access<access::mode::read_write>(cgh);
4949
auto acc1 = buf1.get_access<access::mode::write>(cgh);
5050

51-
cgh.parallel_for<class Test>(
52-
range<1>(1), [=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
53-
using namespace sycl::ext::intel::experimental::esimd;
54-
unsigned int offset = 0;
55-
for (int k = 0; k < VL / 16; k++) {
56-
simd<Ty, 16> var;
57-
var.copy_from(acc0, offset);
58-
var += VAL;
59-
var.copy_to(acc0, offset);
60-
var += 1;
61-
var.copy_to(acc1, offset);
62-
offset += 64;
63-
}
64-
});
51+
cgh.parallel_for<class Test>(range<1>(1),
52+
[=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
53+
using namespace sycl::ext::intel::esimd;
54+
unsigned int offset = 0;
55+
for (int k = 0; k < VL / 16; k++) {
56+
simd<Ty, 16> var;
57+
var.copy_from(acc0, offset);
58+
var += VAL;
59+
var.copy_to(acc0, offset);
60+
var += 1;
61+
var.copy_to(acc1, offset);
62+
offset += 64;
63+
}
64+
});
6565
});
6666

6767
q.wait();

‎SYCL/ESIMD/accessor_gather_scatter.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919

2020
#include <CL/sycl.hpp>
2121
#include <iostream>
22-
#include <sycl/ext/intel/experimental/esimd.hpp>
22+
#include <sycl/ext/intel/esimd.hpp>
2323

2424
using namespace cl::sycl;
2525

@@ -33,7 +33,7 @@ template <typename T, unsigned VL, unsigned STRIDE> struct Kernel {
3333
Kernel(Acc<T> acc) : acc(acc) {}
3434

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

‎SYCL/ESIMD/accessor_load_store.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919

2020
#include <CL/sycl.hpp>
2121
#include <iostream>
22-
#include <sycl/ext/intel/experimental/esimd.hpp>
22+
#include <sycl/ext/intel/esimd.hpp>
2323

2424
using namespace cl::sycl;
2525

@@ -31,7 +31,7 @@ template <typename T> struct Kernel {
3131
Kernel(Acc<T> acc) : acc(acc) {}
3232

3333
void operator()(id<1> i) const SYCL_ESIMD_KERNEL {
34-
using namespace sycl::ext::intel::experimental::esimd;
34+
using namespace sycl::ext::intel::esimd;
3535
uint32_t ii = static_cast<uint32_t>(i.get(0));
3636
T v = scalar_load<T>(acc, ii * sizeof(T));
3737
v += ii;

‎SYCL/ESIMD/aot_mixed.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@
2222

2323
#include <CL/sycl.hpp>
2424
#include <iostream>
25-
#include <sycl/ext/intel/experimental/esimd.hpp>
25+
#include <sycl/ext/intel/esimd.hpp>
2626

2727
using namespace cl::sycl;
2828

@@ -79,7 +79,7 @@ bool test_esimd(queue q) {
7979
auto PC = bufc.get_access<access::mode::write>(cgh);
8080
cgh.parallel_for<class TestESIMD>(
8181
Size / VL, [=](id<1> i) SYCL_ESIMD_KERNEL {
82-
using namespace sycl::ext::intel::experimental::esimd;
82+
using namespace sycl::ext::intel::esimd;
8383
unsigned int offset = i * VL * sizeof(float);
8484
simd<float, VL> va;
8585
va.copy_from(PA, offset);

‎SYCL/ESIMD/api/ballot.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -21,11 +21,11 @@
2121
#include <iostream>
2222
#include <random>
2323

24-
#include <sycl/ext/intel/experimental/esimd.hpp>
24+
#include <sycl/ext/intel/esimd.hpp>
2525

2626
using namespace cl::sycl;
27-
using namespace sycl::ext::intel::experimental;
28-
using namespace sycl::ext::intel::experimental::esimd;
27+
using namespace sycl::ext::intel;
28+
using namespace sycl::ext::intel::esimd;
2929

3030
template <class T, int N> bool test(queue &Q) {
3131
std::cout << " Running " << typeid(T).name() << " test, N=" << N << "...\n";

‎SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -27,10 +27,10 @@
2727

2828
#include <CL/sycl.hpp>
2929
#include <iostream>
30-
#include <sycl/ext/intel/experimental/esimd.hpp>
30+
#include <sycl/ext/intel/esimd.hpp>
3131

3232
using namespace cl::sycl;
33-
using namespace sycl::ext::intel::experimental::esimd;
33+
using namespace sycl::ext::intel::esimd;
3434

3535
template <class T1, class T2, int VL, class OpClass, class Ops> class TestID;
3636

@@ -39,14 +39,14 @@ template <class T1, class T2, class OpClass>
3939
using scalar_comp_t =
4040
std::conditional_t<std::is_same_v<OpClass, esimd_test::CmpOp>,
4141
typename simd_mask<8>::element_type,
42-
__SEIEED::computation_type_t<T1, T2>>;
42+
__ESIMD_DNS::computation_type_t<T1, T2>>;
4343

4444
// Result type of a vector binary Op
4545
template <class T1, class T2, class OpClass, int N = 0>
4646
using comp_t = std::conditional_t<
4747
N == 0, scalar_comp_t<T1, T2, OpClass>,
4848
std::conditional_t<std::is_same_v<OpClass, esimd_test::CmpOp>, simd_mask<N>,
49-
simd<__SEIEED::computation_type_t<T1, T2>, N>>>;
49+
simd<__ESIMD_DNS::computation_type_t<T1, T2>, N>>>;
5050

5151
// Helpers for printing
5252
template <class T> auto cast(T val) { return val; }

‎SYCL/ESIMD/api/esimd_bit_ops.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -17,13 +17,13 @@
1717
#include "../esimd_test_utils.hpp"
1818

1919
#include <CL/sycl.hpp>
20-
#include <sycl/ext/intel/experimental/esimd.hpp>
20+
#include <sycl/ext/intel/esimd.hpp>
2121

2222
#include <iostream>
2323

2424
using namespace cl::sycl;
25-
using namespace sycl::ext::intel::experimental;
26-
using namespace sycl::ext::intel::experimental::esimd;
25+
using namespace sycl::ext::intel;
26+
using namespace sycl::ext::intel::esimd;
2727

2828
struct bit_op {
2929
enum { cbit, fbl, fbh, num_ops };

‎SYCL/ESIMD/api/esimd_merge.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -17,10 +17,10 @@
1717

1818
#include <CL/sycl.hpp>
1919
#include <iostream>
20-
#include <sycl/ext/intel/experimental/esimd.hpp>
20+
#include <sycl/ext/intel/esimd.hpp>
2121

22-
using namespace sycl::ext::intel::experimental::esimd;
23-
using namespace sycl::ext::intel::experimental;
22+
using namespace sycl::ext::intel::esimd;
23+
using namespace sycl::ext::intel;
2424
using namespace cl::sycl;
2525

2626
template <class T> void prn(T *arr, int size, const char *title) {

0 commit comments

Comments
 (0)