diff --git a/SYCL/AOT/spec_const_aot.cpp b/SYCL/AOT/spec_const_aot.cpp index 706b618438..2c88408509 100644 --- a/SYCL/AOT/spec_const_aot.cpp +++ b/SYCL/AOT/spec_const_aot.cpp @@ -36,8 +36,8 @@ int main(int argc, char **argv) { << "\n"; cl::sycl::program prog(q.get_context()); - cl::sycl::ONEAPI::experimental::spec_constant i32 = - prog.set_spec_constant(10); + cl::sycl::ext::oneapi::experimental::spec_constant + i32 = prog.set_spec_constant(10); prog.build_with_kernel_type(); diff --git a/SYCL/AtomicRef/accessor.cpp b/SYCL/AtomicRef/accessor.cpp index 60380d8e53..4726e2f888 100644 --- a/SYCL/AtomicRef/accessor.cpp +++ b/SYCL/AtomicRef/accessor.cpp @@ -10,7 +10,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; // Equivalent to add_test from add.cpp // Uses atomic_accessor instead of atomic_ref diff --git a/SYCL/AtomicRef/assignment.h b/SYCL/AtomicRef/assignment.h index 8c0b8ae0e3..e07b7803ad 100644 --- a/SYCL/AtomicRef/assignment.h +++ b/SYCL/AtomicRef/assignment.h @@ -21,8 +21,8 @@ template void assignment_test(queue q, size_t N) { assignment_buf.template get_access(cgh); cgh.parallel_for>(range<1>(N), [=](item<1> it) { size_t gid = it.get_id(0); - auto atm = atomic_ref(st[0]); atm = T(gid); }); diff --git a/SYCL/Basic/enqueue_barrier.cpp b/SYCL/Basic/enqueue_barrier.cpp index d65428c248..658a4a04e4 100644 --- a/SYCL/Basic/enqueue_barrier.cpp +++ b/SYCL/Basic/enqueue_barrier.cpp @@ -8,7 +8,7 @@ // UNSUPPORTED: cuda || windows #include -#include +#include int main() { sycl::context Context; diff --git a/SYCL/Basic/fpga_tests/buffer_location.cpp b/SYCL/Basic/fpga_tests/buffer_location.cpp index aad641c462..d6e8560bcf 100644 --- a/SYCL/Basic/fpga_tests/buffer_location.cpp +++ b/SYCL/Basic/fpga_tests/buffer_location.cpp @@ -8,7 +8,8 @@ int main() { sycl::buffer Buf{sycl::range{1}}; Queue.submit([&](sycl::handler &CGH) { - sycl::ONEAPI::accessor_property_list PL{sycl::INTEL::buffer_location<1>}; + sycl::ext::oneapi::accessor_property_list PL{ + sycl::ext::intel::buffer_location<1>}; sycl::accessor Acc(Buf, CGH, sycl::write_only, PL); CGH.single_task([=]() { Acc[0] = 42; }); }); diff --git a/SYCL/Basic/fpga_tests/fpga_io_pipes.cpp b/SYCL/Basic/fpga_tests/fpga_io_pipes.cpp index d1d62cdbc7..e30489b6df 100644 --- a/SYCL/Basic/fpga_tests/fpga_io_pipes.cpp +++ b/SYCL/Basic/fpga_tests/fpga_io_pipes.cpp @@ -8,9 +8,9 @@ // //===----------------------------------------------------------------------===// #include -#include #include #include +#include #include "io_pipe_def.h" @@ -113,11 +113,11 @@ int test_io_bl_pipe(cl::sycl::queue Queue) { } int main() { - cl::sycl::queue Queue{cl::sycl::INTEL::fpga_emulator_selector{}}; + cl::sycl::queue Queue{cl::sycl::ext::intel::fpga_emulator_selector{}}; if (!Queue.get_device() .get_info()) { - std::cout << "SYCL_INTEL_data_flow_pipes not supported, skipping" + std::cout << "SYCL_ext_intel_data_flow_pipes not supported, skipping" << std::endl; return 0; } diff --git a/SYCL/Basic/fpga_tests/fpga_lsu.cpp b/SYCL/Basic/fpga_tests/fpga_lsu.cpp index 8acb73199f..74beb0b873 100644 --- a/SYCL/Basic/fpga_tests/fpga_lsu.cpp +++ b/SYCL/Basic/fpga_tests/fpga_lsu.cpp @@ -8,7 +8,7 @@ // //===----------------------------------------------------------------------===// #include -#include +#include // TODO: run is disabled, since no support added in FPGA backend yet. Check // implementation correctness from CXX and SYCL languages perspective. @@ -38,20 +38,20 @@ int test_lsu(cl::sycl::queue Queue) { auto input_ptr = input_accessor.get_pointer(); auto output_ptr = output_accessor.get_pointer(); - using PrefetchingLSU = - cl::sycl::INTEL::lsu, - cl::sycl::INTEL::statically_coalesce>; + using PrefetchingLSU = cl::sycl::ext::intel::lsu< + cl::sycl::ext::intel::prefetch, + cl::sycl::ext::intel::statically_coalesce>; - using BurstCoalescedLSU = - cl::sycl::INTEL::lsu, - cl::sycl::INTEL::statically_coalesce>; + using BurstCoalescedLSU = cl::sycl::ext::intel::lsu< + cl::sycl::ext::intel::burst_coalesce, + cl::sycl::ext::intel::statically_coalesce>; - using CachingLSU = - cl::sycl::INTEL::lsu, - cl::sycl::INTEL::cache<1024>, - cl::sycl::INTEL::statically_coalesce>; + using CachingLSU = cl::sycl::ext::intel::lsu< + cl::sycl::ext::intel::burst_coalesce, + cl::sycl::ext::intel::cache<1024>, + cl::sycl::ext::intel::statically_coalesce>; - using PipelinedLSU = cl::sycl::INTEL::lsu<>; + using PipelinedLSU = cl::sycl::ext::intel::lsu<>; int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0] int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1] @@ -74,7 +74,7 @@ int test_lsu(cl::sycl::queue Queue) { } int main() { - cl::sycl::queue Queue{cl::sycl::INTEL::fpga_emulator_selector{}}; + cl::sycl::queue Queue{cl::sycl::ext::intel::fpga_emulator_selector{}}; return test_lsu(Queue); } diff --git a/SYCL/Basic/fpga_tests/fpga_pipes.cpp b/SYCL/Basic/fpga_tests/fpga_pipes.cpp index d5e61b9a6c..a3b10e65bd 100644 --- a/SYCL/Basic/fpga_tests/fpga_pipes.cpp +++ b/SYCL/Basic/fpga_tests/fpga_pipes.cpp @@ -12,8 +12,8 @@ // //===----------------------------------------------------------------------===// #include -#include #include +#include // Size of an array passing through a pipe constexpr size_t N = 10; @@ -31,7 +31,7 @@ template class templ_nb_pipe; // For non-blocking multiple pipes template -using PipeMulNb = cl::sycl::INTEL::pipe, int>; +using PipeMulNb = cl::sycl::ext::intel::pipe, int>; // For simple blocking pipes with explicit type class some_bl_pipe; @@ -46,7 +46,7 @@ template class templ_bl_pipe; // For blocking multiple pipes template -using PipeMulBl = cl::sycl::INTEL::pipe, int>; +using PipeMulBl = cl::sycl::ext::intel::pipe, int>; // Kernel names template class writer; @@ -57,7 +57,7 @@ template int test_simple_nb_pipe(cl::sycl::queue Queue) { int data[] = {0}; - using Pipe = cl::sycl::INTEL::pipe; + using Pipe = cl::sycl::ext::intel::pipe; cl::sycl::buffer readBuf(data, 1); Queue.submit([&](cl::sycl::handler &cgh) { @@ -146,7 +146,7 @@ template int test_multiple_nb_pipe(cl::sycl::queue Queue) { // Test for array passing through a non-blocking pipe template int test_array_th_nb_pipe(cl::sycl::queue Queue) { int data[N] = {0}; - using AnotherNbPipe = cl::sycl::INTEL::pipe; + using AnotherNbPipe = cl::sycl::ext::intel::pipe; Queue.submit([&](cl::sycl::handler &cgh) { cgh.single_task>([=]() { @@ -188,7 +188,7 @@ template int test_simple_bl_pipe(cl::sycl::queue Queue) { int data[] = {0}; - using Pipe = cl::sycl::INTEL::pipe; + using Pipe = cl::sycl::ext::intel::pipe; cl::sycl::buffer readBuf(data, 1); Queue.submit([&](cl::sycl::handler &cgh) { @@ -256,7 +256,7 @@ template int test_multiple_bl_pipe(cl::sycl::queue Queue) { // Test for array passing through a blocking pipe template int test_array_th_bl_pipe(cl::sycl::queue Queue) { int data[N] = {0}; - using AnotherBlPipe = cl::sycl::INTEL::pipe; + using AnotherBlPipe = cl::sycl::ext::intel::pipe; Queue.submit([&](cl::sycl::handler &cgh) { cgh.single_task>([=]() { @@ -290,7 +290,7 @@ int main() { if (!Queue.get_device() .get_info()) { - std::cout << "SYCL_INTEL_data_flow_pipes not supported, skipping" + std::cout << "SYCL_ext_intel_data_flow_pipes not supported, skipping" << std::endl; return 0; } diff --git a/SYCL/Basic/fpga_tests/global_fpga_device_selector.cpp b/SYCL/Basic/fpga_tests/global_fpga_device_selector.cpp index f63d16c264..88a5ab87ff 100644 --- a/SYCL/Basic/fpga_tests/global_fpga_device_selector.cpp +++ b/SYCL/Basic/fpga_tests/global_fpga_device_selector.cpp @@ -4,12 +4,12 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out #include -#include +#include // Check that FPGA emulator device is found if we try to initialize inline // global variable using fpga_emulator_selector parameter. inline cl::sycl::queue fpga_emu_queue_inlined{ - cl::sycl::INTEL::fpga_emulator_selector{}}; + cl::sycl::ext::intel::fpga_emulator_selector{}}; int main() { return 0; } diff --git a/SYCL/Basic/fpga_tests/io_pipe_def.h b/SYCL/Basic/fpga_tests/io_pipe_def.h index 935170c907..54a5b48e34 100644 --- a/SYCL/Basic/fpga_tests/io_pipe_def.h +++ b/SYCL/Basic/fpga_tests/io_pipe_def.h @@ -1,4 +1,4 @@ -#include +#include namespace intelfpga { template struct ethernet_pipe_id { @@ -6,7 +6,7 @@ template struct ethernet_pipe_id { }; using ethernet_read_pipe = - sycl::INTEL::kernel_readable_io_pipe, int, 0>; + sycl::ext::intel::kernel_readable_io_pipe, int, 0>; using ethernet_write_pipe = - sycl::INTEL::kernel_writeable_io_pipe, int, 0>; + sycl::ext::intel::kernel_writeable_io_pipe, int, 0>; } // namespace intelfpga diff --git a/SYCL/Basic/fpga_tests/pipes_info.cpp b/SYCL/Basic/fpga_tests/pipes_info.cpp index 08a5db1c86..12a19f157a 100644 --- a/SYCL/Basic/fpga_tests/pipes_info.cpp +++ b/SYCL/Basic/fpga_tests/pipes_info.cpp @@ -23,7 +23,7 @@ int main() { Device.get_info(); // Query for platform string. We expect only Intel FPGA platforms to support - // SYCL_INTEL_data_flow_pipes extension. + // SYCL_ext_intel_data_flow_pipes extension. std::string platform_name = Platform.get_info(); bool SupposedToBeSupported = diff --git a/SYCL/Basic/linear-sub_group.cpp b/SYCL/Basic/linear-sub_group.cpp index 2954544359..3feab42d6e 100644 --- a/SYCL/Basic/linear-sub_group.cpp +++ b/SYCL/Basic/linear-sub_group.cpp @@ -38,7 +38,7 @@ int main(int argc, char *argv[]) { nd_range<2>(range<2>(outer, inner), range<2>(outer, inner)), [=](nd_item<2> it) { id<2> idx = it.get_global_id(); - ONEAPI::sub_group sg = it.get_sub_group(); + ext::oneapi::sub_group sg = it.get_sub_group(); output[idx] = sg.get_group_id()[0] * sg.get_local_range()[0] + sg.get_local_id()[0]; }); diff --git a/SYCL/Basic/stream/blocking_pipes_and_stream.cpp b/SYCL/Basic/stream/blocking_pipes_and_stream.cpp index f58c590081..cca007b637 100644 --- a/SYCL/Basic/stream/blocking_pipes_and_stream.cpp +++ b/SYCL/Basic/stream/blocking_pipes_and_stream.cpp @@ -28,7 +28,7 @@ // CHECK-NEXT: 9 #include -#include +#include using namespace cl::sycl; diff --git a/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp b/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp index 13168032e2..bf00a343a5 100644 --- a/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp +++ b/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp @@ -33,7 +33,8 @@ int main() { P.build_with_kernel_type(); cl::sycl::kernel KE = P.get_kernel(); - auto FptrStorage = cl::sycl::ONEAPI::get_device_func_ptr(&add, "add", P, D); + auto FptrStorage = + cl::sycl::ext::oneapi::get_device_func_ptr(&add, "add", P, D); if (!D.is_host()) { // FIXME: update this check with query to supported extension // For now, we don't have runtimes that report required OpenCL extension and @@ -56,8 +57,8 @@ int main() { auto AccB = BufB.template get_access(CGH); CGH.parallel_for( KE, cl::sycl::range<1>(Size), [=](cl::sycl::id<1> Index) { - auto Fptr = - cl::sycl::ONEAPI::to_device_func_ptr(FptrStorage); + auto Fptr = cl::sycl::ext::oneapi::to_device_func_ptr( + FptrStorage); AccA[Index] = Fptr(AccA[Index], AccB[Index]); }); }); diff --git a/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp b/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp index ee26e47089..922c802257 100644 --- a/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp +++ b/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp @@ -35,12 +35,13 @@ int main() { P.build_with_kernel_type(); cl::sycl::kernel KE = P.get_kernel(); - cl::sycl::buffer DispatchTable(2); + cl::sycl::buffer + DispatchTable(2); { auto DTAcc = DispatchTable.get_access(); - DTAcc[0] = cl::sycl::ONEAPI::get_device_func_ptr(&add, "add", P, D); - DTAcc[1] = cl::sycl::ONEAPI::get_device_func_ptr(&sub, "sub", P, D); + DTAcc[0] = cl::sycl::ext::oneapi::get_device_func_ptr(&add, "add", P, D); + DTAcc[1] = cl::sycl::ext::oneapi::get_device_func_ptr(&sub, "sub", P, D); if (!D.is_host()) { // FIXME: update this check with query to supported extension // For now, we don't have runtimes that report required OpenCL extension @@ -70,7 +71,7 @@ int main() { DispatchTable.template get_access(CGH); CGH.parallel_for( KE, cl::sycl::range<1>(Size), [=](cl::sycl::id<1> Index) { - auto FP = cl::sycl::ONEAPI::to_device_func_ptr( + auto FP = cl::sycl::ext::oneapi::to_device_func_ptr( AccDT[Mode]); AccA[Index] = FP(AccA[Index], AccB[Index]); diff --git a/SYCL/DeviceLib/built-ins/printf.cpp b/SYCL/DeviceLib/built-ins/printf.cpp index db5e59cbd4..e84f4664b5 100644 --- a/SYCL/DeviceLib/built-ins/printf.cpp +++ b/SYCL/DeviceLib/built-ins/printf.cpp @@ -41,7 +41,7 @@ int main() { Queue.submit([&](handler &CGH) { CGH.single_task([=]() { // String - ONEAPI::experimental::printf(format_hello_world); + ext::oneapi::experimental::printf(format_hello_world); // Due to a bug in Intel CPU Runtime for OpenCL on Windows, information // printed using such format strings (without %-specifiers) might // appear in different order if output is redirected to a file or @@ -50,8 +50,8 @@ int main() { // CHECK: {{(Hello, World!)?}} // Integral types - ONEAPI::experimental::printf(format_int, (int32_t)123); - ONEAPI::experimental::printf(format_int, (int32_t)-123); + ext::oneapi::experimental::printf(format_int, (int32_t)123); + ext::oneapi::experimental::printf(format_int, (int32_t)-123); // CHECK: 123 // CHECK-NEXT: -123 @@ -60,8 +60,8 @@ int main() { // You can declare format string in non-global scope, but in this case // static keyword is required static const CONSTANT char format[] = "%f\n"; - ONEAPI::experimental::printf(format, 33.4f); - ONEAPI::experimental::printf(format, -33.4f); + ext::oneapi::experimental::printf(format, 33.4f); + ext::oneapi::experimental::printf(format, -33.4f); } // CHECK-NEXT: 33.4 // CHECK-NEXT: -33.4 @@ -73,23 +73,23 @@ int main() { using ocl_int4 = cl::sycl::vec::vector_t; { static const CONSTANT char format[] = "%v4d\n"; - ONEAPI::experimental::printf(format, (ocl_int4)v4); + ext::oneapi::experimental::printf(format, (ocl_int4)v4); } // However, you are still able to print them by-element: { - ONEAPI::experimental::printf(format_vec, (int32_t)v4.w(), - (int32_t)v4.z(), (int32_t)v4.y(), - (int32_t)v4.x()); + ext::oneapi::experimental::printf(format_vec, (int32_t)v4.w(), + (int32_t)v4.z(), (int32_t)v4.y(), + (int32_t)v4.x()); } #else // On host side you always have to print them by-element: - ONEAPI::experimental::printf(format_vec, (int32_t)v4.x(), - (int32_t)v4.y(), (int32_t)v4.z(), - (int32_t)v4.w()); - ONEAPI::experimental::printf(format_vec, (int32_t)v4.w(), - (int32_t)v4.z(), (int32_t)v4.y(), - (int32_t)v4.x()); + ext::oneapi::experimental::printf(format_vec, (int32_t)v4.x(), + (int32_t)v4.y(), (int32_t)v4.z(), + (int32_t)v4.w()); + ext::oneapi::experimental::printf(format_vec, (int32_t)v4.w(), + (int32_t)v4.z(), (int32_t)v4.y(), + (int32_t)v4.x()); #endif // __SYCL_DEVICE_ONLY__ // CHECK-NEXT: 5,6,7,8 // CHECK-NEXT: 8,7,6,5 @@ -100,7 +100,7 @@ int main() { // According to OpenCL spec, argument should be a void pointer { static const CONSTANT char format[] = "%p\n"; - ONEAPI::experimental::printf(format, (void *)Ptr); + ext::oneapi::experimental::printf(format, (void *)Ptr); } // CHECK-NEXT: {{(0x)?[0-9a-fA-F]+$}} }); @@ -114,7 +114,8 @@ int main() { Queue.submit([&](handler &CGH) { CGH.parallel_for(range<1>(10), [=](id<1> i) { // cast to uint64_t to be sure that we pass 64-bit unsigned value - ONEAPI::experimental::printf(format_hello_world_2, (uint64_t)i.get(0)); + ext::oneapi::experimental::printf(format_hello_world_2, + (uint64_t)i.get(0)); }); }); Queue.wait(); diff --git a/SYCL/DotProduct/dot_product_int_test.cpp b/SYCL/DotProduct/dot_product_int_test.cpp index c6eaec1002..6f51200c2d 100644 --- a/SYCL/DotProduct/dot_product_int_test.cpp +++ b/SYCL/DotProduct/dot_product_int_test.cpp @@ -8,10 +8,10 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out #include -#include #include #include #include +#include // Change if tests are added/removed static int testCount = 4; @@ -19,7 +19,7 @@ static int passCount; using namespace cl::sycl; using namespace cl::sycl::detail::gtl; -using namespace cl::sycl::ONEAPI; +using namespace cl::sycl::ext::oneapi; constexpr int RangeLength = 100; diff --git a/SYCL/DotProduct/dot_product_vec_test.cpp b/SYCL/DotProduct/dot_product_vec_test.cpp index 37570df178..ce9c87bab8 100644 --- a/SYCL/DotProduct/dot_product_vec_test.cpp +++ b/SYCL/DotProduct/dot_product_vec_test.cpp @@ -8,10 +8,10 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out #include -#include #include #include #include +#include // Change if tests are added/removed static int testCount = 4; @@ -19,7 +19,7 @@ static int passCount; using namespace cl::sycl; using namespace cl::sycl::detail::gtl; -using namespace cl::sycl::ONEAPI; +using namespace cl::sycl::ext::oneapi; constexpr int RangeLength = 100; diff --git a/SYCL/ESIMD/BitonicSortK.cpp b/SYCL/ESIMD/BitonicSortK.cpp index 1942477c40..f6331d36ee 100644 --- a/SYCL/ESIMD/BitonicSortK.cpp +++ b/SYCL/ESIMD/BitonicSortK.cpp @@ -13,9 +13,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; diff --git a/SYCL/ESIMD/BitonicSortKv2.cpp b/SYCL/ESIMD/BitonicSortKv2.cpp index 42466539b0..d96d2214f6 100644 --- a/SYCL/ESIMD/BitonicSortKv2.cpp +++ b/SYCL/ESIMD/BitonicSortKv2.cpp @@ -14,9 +14,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; diff --git a/SYCL/ESIMD/PrefixSum.cpp b/SYCL/ESIMD/PrefixSum.cpp index 46452cb091..72a38cb75e 100644 --- a/SYCL/ESIMD/PrefixSum.cpp +++ b/SYCL/ESIMD/PrefixSum.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include #define MAX_TS_WIDTH 1024 // kernel can handle TUPLE_SZ 1, 2, or 4 diff --git a/SYCL/ESIMD/Prefix_Local_sum1.cpp b/SYCL/ESIMD/Prefix_Local_sum1.cpp index 9993800a94..037c5d4d2f 100644 --- a/SYCL/ESIMD/Prefix_Local_sum1.cpp +++ b/SYCL/ESIMD/Prefix_Local_sum1.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include #define MAX_TS_WIDTH 1024 // kernel can handle TUPLE_SZ 1, 2, or 4 diff --git a/SYCL/ESIMD/Prefix_Local_sum2.cpp b/SYCL/ESIMD/Prefix_Local_sum2.cpp index 550c97e9e9..01404da74c 100644 --- a/SYCL/ESIMD/Prefix_Local_sum2.cpp +++ b/SYCL/ESIMD/Prefix_Local_sum2.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include #define MAX_TS_WIDTH 1024 // kernel can handle TUPLE_SZ 1, 2, or 4 diff --git a/SYCL/ESIMD/Prefix_Local_sum3.cpp b/SYCL/ESIMD/Prefix_Local_sum3.cpp index 887d350d01..438ac6e1b0 100644 --- a/SYCL/ESIMD/Prefix_Local_sum3.cpp +++ b/SYCL/ESIMD/Prefix_Local_sum3.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include #define MAX_TS_WIDTH 1024 // kernel can handle TUPLE_SZ 1, 2, or 4 diff --git a/SYCL/ESIMD/Stencil.cpp b/SYCL/ESIMD/Stencil.cpp index 8bc0d4f90e..e28e596f28 100644 --- a/SYCL/ESIMD/Stencil.cpp +++ b/SYCL/ESIMD/Stencil.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include #define WIDTH 16 #define HEIGHT 16 diff --git a/SYCL/ESIMD/accessor.cpp b/SYCL/ESIMD/accessor.cpp index 976e14b607..33a5d94d1c 100644 --- a/SYCL/ESIMD/accessor.cpp +++ b/SYCL/ESIMD/accessor.cpp @@ -15,7 +15,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/accessor_gather_scatter.cpp b/SYCL/ESIMD/accessor_gather_scatter.cpp index a1e93f9dde..89907d6b29 100644 --- a/SYCL/ESIMD/accessor_gather_scatter.cpp +++ b/SYCL/ESIMD/accessor_gather_scatter.cpp @@ -16,8 +16,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/accessor_load_store.cpp b/SYCL/ESIMD/accessor_load_store.cpp index 573b3f1c10..7ef3f209bd 100644 --- a/SYCL/ESIMD/accessor_load_store.cpp +++ b/SYCL/ESIMD/accessor_load_store.cpp @@ -16,8 +16,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/dp4a.cpp b/SYCL/ESIMD/dp4a.cpp index e85951d3d7..e635da422e 100644 --- a/SYCL/ESIMD/dp4a.cpp +++ b/SYCL/ESIMD/dp4a.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/ext_math.cpp b/SYCL/ESIMD/ext_math.cpp index 5fc87d850e..2985af16d6 100644 --- a/SYCL/ESIMD/ext_math.cpp +++ b/SYCL/ESIMD/ext_math.cpp @@ -15,9 +15,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp index 48e1dba1c5..0dfbc281f7 100644 --- a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp +++ b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp @@ -19,8 +19,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include static_assert(SIZE >= VL, "Size must greater than or equal to VL"); static_assert(SIZE % VL == 0, "Size must be multiple of VL"); diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp index ebd6251aae..292d03f370 100644 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -21,7 +21,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/fp_call_recursive.cpp b/SYCL/ESIMD/fp_call_recursive.cpp index f44ce0e0aa..5dd43e1230 100644 --- a/SYCL/ESIMD/fp_call_recursive.cpp +++ b/SYCL/ESIMD/fp_call_recursive.cpp @@ -18,7 +18,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index c34b729459..ae63a56f25 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -20,7 +20,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include #include diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index 8f96afa585..08c70655ba 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -22,7 +22,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/histogram.cpp b/SYCL/ESIMD/histogram.cpp index c13d7299da..97fb4b97d5 100644 --- a/SYCL/ESIMD/histogram.cpp +++ b/SYCL/ESIMD/histogram.cpp @@ -14,9 +14,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/histogram_256_slm.cpp b/SYCL/ESIMD/histogram_256_slm.cpp index 4dc963fe97..55d8910ed1 100644 --- a/SYCL/ESIMD/histogram_256_slm.cpp +++ b/SYCL/ESIMD/histogram_256_slm.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include static constexpr int NUM_BINS = 256; static constexpr int SLM_SIZE = (NUM_BINS * 4); diff --git a/SYCL/ESIMD/histogram_256_slm_spec.cpp b/SYCL/ESIMD/histogram_256_slm_spec.cpp index 07f615146a..8713b7c1ba 100644 --- a/SYCL/ESIMD/histogram_256_slm_spec.cpp +++ b/SYCL/ESIMD/histogram_256_slm_spec.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include static constexpr int NUM_BINS = 256; static constexpr int SLM_SIZE = (NUM_BINS * 4); @@ -160,7 +160,7 @@ int main(int argc, char **argv) { } cl::sycl::program prg(q.get_context()); - sycl::ONEAPI::experimental::spec_constant + sycl::ext::oneapi::experimental::spec_constant num_blocks_const = prg.set_spec_constant(num_blocks); prg.build_with_kernel_type(); diff --git a/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp b/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp index af444ce2df..689265abc7 100644 --- a/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp +++ b/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp @@ -7,8 +7,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include static constexpr int NUM_BINS = 256; static constexpr int SLM_SIZE = (NUM_BINS * 4); diff --git a/SYCL/ESIMD/histogram_2d.cpp b/SYCL/ESIMD/histogram_2d.cpp index 016e03aba3..ce7c32badd 100644 --- a/SYCL/ESIMD/histogram_2d.cpp +++ b/SYCL/ESIMD/histogram_2d.cpp @@ -14,9 +14,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/histogram_raw_send.cpp b/SYCL/ESIMD/histogram_raw_send.cpp index c62ce2b670..3fcae5afa2 100644 --- a/SYCL/ESIMD/histogram_raw_send.cpp +++ b/SYCL/ESIMD/histogram_raw_send.cpp @@ -15,9 +15,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/kmeans/kmeans.cpp b/SYCL/ESIMD/kmeans/kmeans.cpp index 9a74b0083b..e2fa4e4823 100644 --- a/SYCL/ESIMD/kmeans/kmeans.cpp +++ b/SYCL/ESIMD/kmeans/kmeans.cpp @@ -15,10 +15,10 @@ #include "esimd_test_utils.hpp" #include -#include #include #include #include +#include #include using namespace cl::sycl; diff --git a/SYCL/ESIMD/linear/linear.cpp b/SYCL/ESIMD/linear/linear.cpp index 1d2ccd83c0..938d5043b8 100644 --- a/SYCL/ESIMD/linear/linear.cpp +++ b/SYCL/ESIMD/linear/linear.cpp @@ -15,9 +15,9 @@ #include "esimd_test_utils.hpp" #include -#include #include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/mandelbrot/mandelbrot.cpp b/SYCL/ESIMD/mandelbrot/mandelbrot.cpp index 8ad6acacb6..c2bff56b15 100644 --- a/SYCL/ESIMD/mandelbrot/mandelbrot.cpp +++ b/SYCL/ESIMD/mandelbrot/mandelbrot.cpp @@ -12,10 +12,10 @@ #include "esimd_test_utils.hpp" #include -#include #include #include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; diff --git a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp index 5207f10563..44b7bf36b0 100644 --- a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp +++ b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp @@ -13,10 +13,10 @@ #include "esimd_test_utils.hpp" #include -#include #include #include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; @@ -126,16 +126,16 @@ int main(int argc, char *argv[]) { << ", thrs = " << thrs << "\n"; } cl::sycl::program prg(q.get_context()); - sycl::ONEAPI::experimental::spec_constant crunch_const = - prg.set_spec_constant(crunch); - sycl::ONEAPI::experimental::spec_constant xoff_const = - prg.set_spec_constant(xoff); - sycl::ONEAPI::experimental::spec_constant yoff_const = - prg.set_spec_constant(yoff); - sycl::ONEAPI::experimental::spec_constant scale_const = - prg.set_spec_constant(scale); - sycl::ONEAPI::experimental::spec_constant thrs_const = - prg.set_spec_constant(thrs); + sycl::ext::oneapi::experimental::spec_constant + crunch_const = prg.set_spec_constant(crunch); + sycl::ext::oneapi::experimental::spec_constant + xoff_const = prg.set_spec_constant(xoff); + sycl::ext::oneapi::experimental::spec_constant + yoff_const = prg.set_spec_constant(yoff); + sycl::ext::oneapi::experimental::spec_constant + scale_const = prg.set_spec_constant(scale); + sycl::ext::oneapi::experimental::spec_constant + thrs_const = prg.set_spec_constant(thrs); prg.build_with_kernel_type(); auto e = q.submit([&](cl::sycl::handler &cgh) { diff --git a/SYCL/ESIMD/matrix_transpose.cpp b/SYCL/ESIMD/matrix_transpose.cpp index 97d88608f1..02d63ba4b4 100644 --- a/SYCL/ESIMD/matrix_transpose.cpp +++ b/SYCL/ESIMD/matrix_transpose.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; using namespace std; diff --git a/SYCL/ESIMD/matrix_transpose2.cpp b/SYCL/ESIMD/matrix_transpose2.cpp index 610999e0a3..32ff8f0ff5 100644 --- a/SYCL/ESIMD/matrix_transpose2.cpp +++ b/SYCL/ESIMD/matrix_transpose2.cpp @@ -15,8 +15,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; using namespace std; diff --git a/SYCL/ESIMD/matrix_transpose_glb.cpp b/SYCL/ESIMD/matrix_transpose_glb.cpp index 0ded82f1bb..e56438446e 100644 --- a/SYCL/ESIMD/matrix_transpose_glb.cpp +++ b/SYCL/ESIMD/matrix_transpose_glb.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; using namespace std; diff --git a/SYCL/ESIMD/matrix_transpose_usm.cpp b/SYCL/ESIMD/matrix_transpose_usm.cpp index 78bf00dbec..d192945d5d 100644 --- a/SYCL/ESIMD/matrix_transpose_usm.cpp +++ b/SYCL/ESIMD/matrix_transpose_usm.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; using namespace std; diff --git a/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp b/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp index 6aca9821af..a43c3a9ffe 100644 --- a/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp +++ b/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp @@ -18,8 +18,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include static_assert(SIZE >= VL, "Size must greater than or equal to VL"); static_assert(SIZE % VL == 0, "Size must be multiple of VL"); diff --git a/SYCL/ESIMD/noinline_call_from_func.cpp b/SYCL/ESIMD/noinline_call_from_func.cpp index fe45037c67..4633bb84a4 100644 --- a/SYCL/ESIMD/noinline_call_from_func.cpp +++ b/SYCL/ESIMD/noinline_call_from_func.cpp @@ -20,7 +20,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/noinline_call_recursive.cpp b/SYCL/ESIMD/noinline_call_recursive.cpp index 9649b38961..c6083d01f6 100644 --- a/SYCL/ESIMD/noinline_call_recursive.cpp +++ b/SYCL/ESIMD/noinline_call_recursive.cpp @@ -18,7 +18,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/printf.cpp b/SYCL/ESIMD/printf.cpp index 9b0459333a..0d2820fded 100644 --- a/SYCL/ESIMD/printf.cpp +++ b/SYCL/ESIMD/printf.cpp @@ -21,7 +21,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include #include diff --git a/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp b/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp index 099471a28f..a0a1fe27be 100644 --- a/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp +++ b/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/reduction.cpp b/SYCL/ESIMD/reduction.cpp index 4238cbf8d9..9e7fe9869e 100644 --- a/SYCL/ESIMD/reduction.cpp +++ b/SYCL/ESIMD/reduction.cpp @@ -15,8 +15,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/regression/big_const_initializer.cpp b/SYCL/ESIMD/regression/big_const_initializer.cpp index f9411814dc..9cbafaf5d7 100644 --- a/SYCL/ESIMD/regression/big_const_initializer.cpp +++ b/SYCL/ESIMD/regression/big_const_initializer.cpp @@ -16,7 +16,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include diff --git a/SYCL/ESIMD/regression/dgetrf.cpp b/SYCL/ESIMD/regression/dgetrf.cpp index dd6951edc8..d8ac134e4d 100644 --- a/SYCL/ESIMD/regression/dgetrf.cpp +++ b/SYCL/ESIMD/regression/dgetrf.cpp @@ -17,11 +17,11 @@ // corresponding to LU input sizes; all internal functions are inlined. // #include -#include #include #include #include #include +#include #define ABS(x) ((x) >= 0 ? (x) : -(x)) #define MIN(x, y) ((x) <= (y) ? (x) : (y)) diff --git a/SYCL/ESIMD/regression/operator_decrement.cpp b/SYCL/ESIMD/regression/operator_decrement.cpp index a3f5a8d381..9745fc7fe9 100644 --- a/SYCL/ESIMD/regression/operator_decrement.cpp +++ b/SYCL/ESIMD/regression/operator_decrement.cpp @@ -15,8 +15,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/regression/replicate_bug.cpp b/SYCL/ESIMD/regression/replicate_bug.cpp index a5773ad6b2..ae8aedfdcf 100644 --- a/SYCL/ESIMD/regression/replicate_bug.cpp +++ b/SYCL/ESIMD/regression/replicate_bug.cpp @@ -16,8 +16,8 @@ #include "../esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; constexpr int VL = 8; diff --git a/SYCL/ESIMD/regression/store_zero_const.cpp b/SYCL/ESIMD/regression/store_zero_const.cpp index 9043a7dbd2..698b31b333 100644 --- a/SYCL/ESIMD/regression/store_zero_const.cpp +++ b/SYCL/ESIMD/regression/store_zero_const.cpp @@ -19,7 +19,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include #include diff --git a/SYCL/ESIMD/regression/unused_load.cpp b/SYCL/ESIMD/regression/unused_load.cpp index a979554117..7b6e1d6ed7 100644 --- a/SYCL/ESIMD/regression/unused_load.cpp +++ b/SYCL/ESIMD/regression/unused_load.cpp @@ -14,7 +14,7 @@ // copy_from invocation. #include -#include +#include #include diff --git a/SYCL/ESIMD/slm_barrier.cpp b/SYCL/ESIMD/slm_barrier.cpp index c28f64ac8f..92b3b7af23 100644 --- a/SYCL/ESIMD/slm_barrier.cpp +++ b/SYCL/ESIMD/slm_barrier.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; diff --git a/SYCL/ESIMD/slm_split_barrier.cpp b/SYCL/ESIMD/slm_split_barrier.cpp index db248a2200..5970a56423 100644 --- a/SYCL/ESIMD/slm_split_barrier.cpp +++ b/SYCL/ESIMD/slm_split_barrier.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; using namespace sycl::ext::intel::experimental::esimd; diff --git a/SYCL/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp index 29a29b93c1..d3739855c3 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp @@ -5,7 +5,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include #include diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index 2dca350838..af09a78dcd 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -12,7 +12,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include #include diff --git a/SYCL/ESIMD/spec_const_redefine_esimd.cpp b/SYCL/ESIMD/spec_const_redefine_esimd.cpp index 1e1ccabf2f..7ceda7a795 100755 --- a/SYCL/ESIMD/spec_const_redefine_esimd.cpp +++ b/SYCL/ESIMD/spec_const_redefine_esimd.cpp @@ -23,7 +23,7 @@ #include "esimd_test_utils.hpp" #include -#include +#include #include #include @@ -63,9 +63,9 @@ int main(int argc, char **argv) { for (int i = 0; i < n_sc_sets; i++) { sycl::program program(q.get_context()); const int *sc_set = &sc_vals[i][0]; - sycl::ONEAPI::experimental::spec_constant sc0 = + sycl::ext::oneapi::experimental::spec_constant sc0 = program.set_spec_constant(sc_set[0]); - sycl::ONEAPI::experimental::spec_constant sc1 = + sycl::ext::oneapi::experimental::spec_constant sc1 = program.set_spec_constant(sc_set[1]); program.build_with_kernel_type(); diff --git a/SYCL/ESIMD/stencil2.cpp b/SYCL/ESIMD/stencil2.cpp index 791a893c61..1f6cd4c066 100644 --- a/SYCL/ESIMD/stencil2.cpp +++ b/SYCL/ESIMD/stencil2.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include #define WIDTH 16 #define HEIGHT 16 diff --git a/SYCL/ESIMD/sycl_esimd_mix.cpp b/SYCL/ESIMD/sycl_esimd_mix.cpp index ac8f38876b..f741972b57 100644 --- a/SYCL/ESIMD/sycl_esimd_mix.cpp +++ b/SYCL/ESIMD/sycl_esimd_mix.cpp @@ -17,8 +17,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/test_id_3d.cpp b/SYCL/ESIMD/test_id_3d.cpp index 8d8fa665d7..1863251376 100644 --- a/SYCL/ESIMD/test_id_3d.cpp +++ b/SYCL/ESIMD/test_id_3d.cpp @@ -17,8 +17,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/usm_gather_scatter_rgba.cpp b/SYCL/ESIMD/usm_gather_scatter_rgba.cpp index 50f73e38c6..c72babedd5 100644 --- a/SYCL/ESIMD/usm_gather_scatter_rgba.cpp +++ b/SYCL/ESIMD/usm_gather_scatter_rgba.cpp @@ -16,8 +16,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/vadd_1d.cpp b/SYCL/ESIMD/vadd_1d.cpp index 5ac586d4b9..de66524362 100644 --- a/SYCL/ESIMD/vadd_1d.cpp +++ b/SYCL/ESIMD/vadd_1d.cpp @@ -13,8 +13,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/vadd_2d.cpp b/SYCL/ESIMD/vadd_2d.cpp index 1317db6cd4..28e762777e 100644 --- a/SYCL/ESIMD/vadd_2d.cpp +++ b/SYCL/ESIMD/vadd_2d.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/vadd_2d_acc.cpp b/SYCL/ESIMD/vadd_2d_acc.cpp index 260e8a67da..22d42e5300 100644 --- a/SYCL/ESIMD/vadd_2d_acc.cpp +++ b/SYCL/ESIMD/vadd_2d_acc.cpp @@ -17,8 +17,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/vadd_raw_send.cpp b/SYCL/ESIMD/vadd_raw_send.cpp index d16343bdcd..eb81dcc68d 100644 --- a/SYCL/ESIMD/vadd_raw_send.cpp +++ b/SYCL/ESIMD/vadd_raw_send.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/ESIMD/vadd_usm.cpp b/SYCL/ESIMD/vadd_usm.cpp index bbd1529dd0..ce59f40096 100644 --- a/SYCL/ESIMD/vadd_usm.cpp +++ b/SYCL/ESIMD/vadd_usm.cpp @@ -14,8 +14,8 @@ #include "esimd_test_utils.hpp" #include -#include #include +#include using namespace cl::sycl; diff --git a/SYCL/FilterSelector/reuse.cpp b/SYCL/FilterSelector/reuse.cpp index f8197a908b..bfb5596aac 100644 --- a/SYCL/FilterSelector/reuse.cpp +++ b/SYCL/FilterSelector/reuse.cpp @@ -17,7 +17,7 @@ using namespace cl::sycl; // TODO: change to 'using namespace cl::sycl::oneapi' after PR intel/llvm:4014 // is merged -using namespace cl::sycl::ONEAPI; +using namespace cl::sycl::ext::oneapi; int main() { std::vector Devs; @@ -29,15 +29,15 @@ int main() { if (Devs.size() > 1) { // TODO: change all occurrences of filter_selector to 'filter_selector' or // 'oneapi::filter_selector' after PR intel/llvm:4014 is merged - ONEAPI::filter_selector filter("1"); + ext::oneapi::filter_selector filter("1"); device d1(filter); device d2(filter); assert(d1 == d2); - ONEAPI::filter_selector f1("0"); - ONEAPI::filter_selector f2("1"); + ext::oneapi::filter_selector f1("0"); + ext::oneapi::filter_selector f2("1"); device d3(f1); device d4(f2); diff --git a/SYCL/FilterSelector/select.cpp b/SYCL/FilterSelector/select.cpp index 693ff949a4..a0485457b0 100644 --- a/SYCL/FilterSelector/select.cpp +++ b/SYCL/FilterSelector/select.cpp @@ -17,7 +17,7 @@ using namespace cl::sycl; // TODO: change to 'using namespace cl::sycl::oneapi' after PR intel/llvm:4014 // is merged -using namespace cl::sycl::ONEAPI; +using namespace cl::sycl::ext::oneapi; int main() { std::vector CPUs; @@ -75,52 +75,52 @@ int main() { std::cout << "Test 'cpu'"; // TODO: change all occurrences of filter_selector to 'filter_selector' or // 'oneapi::filter_selector' after PR intel/llvm:4014 is merged - device d1(ONEAPI::filter_selector("cpu")); + device d1(ext::oneapi::filter_selector("cpu")); assert(d1.is_cpu()); std::cout << "...PASS" << std::endl; } if (!GPUs.empty()) { std::cout << "Test 'gpu'"; - device d2(ONEAPI::filter_selector("gpu")); + device d2(ext::oneapi::filter_selector("gpu")); assert(d2.is_gpu()); std::cout << "...PASS" << std::endl; } if (!CPUs.empty() || !GPUs.empty()) { std::cout << "Test 'cpu,gpu'"; - device d3(ONEAPI::filter_selector("cpu,gpu")); + device d3(ext::oneapi::filter_selector("cpu,gpu")); assert((d3.is_gpu() || d3.is_cpu())); std::cout << "...PASS" << std::endl; } if (HasOpenCLDevices) { std::cout << "Test 'opencl'"; - device d4(ONEAPI::filter_selector("opencl")); + device d4(ext::oneapi::filter_selector("opencl")); assert(d4.get_platform().get_backend() == backend::opencl); std::cout << "...PASS" << std::endl; if (!CPUs.empty()) { std::cout << "Test 'opencl:cpu'"; - device d5(ONEAPI::filter_selector("opencl:cpu")); + device d5(ext::oneapi::filter_selector("opencl:cpu")); assert(d5.is_cpu() && d5.get_platform().get_backend() == backend::opencl); std::cout << "...PASS" << std::endl; std::cout << "Test 'opencl:cpu:0'"; - device d6(ONEAPI::filter_selector("opencl:cpu:0")); + device d6(ext::oneapi::filter_selector("opencl:cpu:0")); assert(d6.is_cpu() && d6.get_platform().get_backend() == backend::opencl); std::cout << "...PASS" << std::endl; } if (HasOpenCLGPU) { std::cout << "Test 'opencl:gpu'" << std::endl; - device d7(ONEAPI::filter_selector("opencl:gpu")); + device d7(ext::oneapi::filter_selector("opencl:gpu")); assert(d7.is_gpu() && d7.get_platform().get_backend() == backend::opencl); } } std::cout << "Test '0'"; - device d8(ONEAPI::filter_selector("0")); + device d8(ext::oneapi::filter_selector("0")); std::cout << "...PASS" << std::endl; std::string ErrorMesg( @@ -128,7 +128,7 @@ int main() { try { // pick something crazy - device d9(ONEAPI::filter_selector("gpu:999")); + device d9(ext::oneapi::filter_selector("gpu:999")); std::cout << "d9 = " << d9.get_info() << std::endl; } catch (const sycl::runtime_error &e) { assert(ErrorMesg.find_first_of(e.what()) == 0); @@ -137,7 +137,7 @@ int main() { try { // pick something crazy - device d10(ONEAPI::filter_selector("bob:gpu")); + device d10(ext::oneapi::filter_selector("bob:gpu")); std::cout << "d10 = " << d10.get_info() << std::endl; } catch (const sycl::runtime_error &e) { assert(ErrorMesg.find_first_of(e.what()) == 0); @@ -146,7 +146,7 @@ int main() { try { // pick something crazy - device d11(ONEAPI::filter_selector("opencl:bob")); + device d11(ext::oneapi::filter_selector("opencl:bob")); std::cout << "d11 = " << d11.get_info() << std::endl; } catch (const sycl::runtime_error &e) { assert(ErrorMesg.find_first_of(e.what()) == 0); @@ -155,19 +155,19 @@ int main() { if (HasLevelZeroDevices && HasLevelZeroGPU) { std::cout << "Test 'level_zero'"; - device d12(ONEAPI::filter_selector("level_zero")); + device d12(ext::oneapi::filter_selector("level_zero")); assert(d12.get_platform().get_backend() == backend::level_zero); std::cout << "...PASS" << std::endl; std::cout << "Test 'level_zero:gpu'"; - device d13(ONEAPI::filter_selector("level_zero:gpu")); + device d13(ext::oneapi::filter_selector("level_zero:gpu")); assert(d13.is_gpu() && d13.get_platform().get_backend() == backend::level_zero); std::cout << "...PASS" << std::endl; if (HasOpenCLDevices && !CPUs.empty()) { std::cout << "Test 'level_zero:gpu,cpu'"; - device d14(ONEAPI::filter_selector("level_zero:gpu,cpu")); + device d14(ext::oneapi::filter_selector("level_zero:gpu,cpu")); assert((d14.is_gpu() || d14.is_cpu())); std::cout << "...PASS 1/2" << std::endl; if (d14.is_gpu()) { @@ -179,37 +179,37 @@ int main() { if (Devs.size() > 1) { std::cout << "Test '1'"; - device d15(ONEAPI::filter_selector("1")); + device d15(ext::oneapi::filter_selector("1")); std::cout << "...PASS" << std::endl; } if (HasCUDADevices) { std::cout << "Test 'cuda'"; - device d16(ONEAPI::filter_selector("cuda")); + device d16(ext::oneapi::filter_selector("cuda")); assert(d16.get_platform().get_backend() == backend::cuda); std::cout << "...PASS" << std::endl; std::cout << "Test 'cuda:gpu'"; - device d17(ONEAPI::filter_selector("cuda:gpu")); + device d17(ext::oneapi::filter_selector("cuda:gpu")); assert(d17.is_gpu() && d17.get_platform().get_backend() == backend::cuda); std::cout << "...PASS" << std::endl; } if (!Accels.empty()) { std::cout << "Test 'accelerator'"; - device d18(ONEAPI::filter_selector("accelerator")); + device d18(ext::oneapi::filter_selector("accelerator")); assert(d18.is_accelerator()); std::cout << "...PASS" << std::endl; } if (HasROCmDevices) { std::cout << "Test 'rocm'"; - device d19(ONEAPI::filter_selector("rocm")); + device d19(ext::oneapi::filter_selector("rocm")); assert(d19.get_platform().get_backend() == backend::rocm); std::cout << "...PASS" << std::endl; std::cout << "test 'rocm:gpu'"; - device d20(ONEAPI::filter_selector("rocm:gpu")); + device d20(ext::oneapi::filter_selector("rocm:gpu")); assert(d20.is_gpu() && d19.get_platform().get_backend() == backend::rocm); std::cout << "...PASS" << std::endl; } diff --git a/SYCL/GroupAlgorithm/all_of.cpp b/SYCL/GroupAlgorithm/all_of.cpp index 844a48c740..f1abf43ea0 100644 --- a/SYCL/GroupAlgorithm/all_of.cpp +++ b/SYCL/GroupAlgorithm/all_of.cpp @@ -10,7 +10,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template class all_of_kernel; diff --git a/SYCL/GroupAlgorithm/any_of.cpp b/SYCL/GroupAlgorithm/any_of.cpp index dcdd392c78..d01a0a7814 100644 --- a/SYCL/GroupAlgorithm/any_of.cpp +++ b/SYCL/GroupAlgorithm/any_of.cpp @@ -10,7 +10,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template class any_of_kernel; diff --git a/SYCL/GroupAlgorithm/broadcast.cpp b/SYCL/GroupAlgorithm/broadcast.cpp index 97f7161bde..4ccd629aba 100644 --- a/SYCL/GroupAlgorithm/broadcast.cpp +++ b/SYCL/GroupAlgorithm/broadcast.cpp @@ -11,7 +11,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template diff --git a/SYCL/GroupAlgorithm/exclusive_scan.cpp b/SYCL/GroupAlgorithm/exclusive_scan.cpp index 02d4f87830..186b6da851 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan.cpp @@ -21,7 +21,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template class exclusive_scan_kernel; @@ -140,26 +140,27 @@ int main() { std::iota(input.begin(), input.end(), 0); std::fill(output.begin(), output.end(), 0); - test(q, input, output, ONEAPI::plus<>(), 0); - test(q, input, output, ONEAPI::minimum<>(), + test(q, input, output, ext::oneapi::plus<>(), 0); + test(q, input, output, ext::oneapi::minimum<>(), std::numeric_limits::max()); - test(q, input, output, ONEAPI::maximum<>(), + test(q, input, output, ext::oneapi::maximum<>(), std::numeric_limits::lowest()); - test(q, input, output, ONEAPI::plus(), 0); - test(q, input, output, ONEAPI::minimum(), + test(q, input, output, ext::oneapi::plus(), 0); + test(q, input, output, ext::oneapi::minimum(), std::numeric_limits::max()); - test(q, input, output, ONEAPI::maximum(), + test(q, input, output, ext::oneapi::maximum(), std::numeric_limits::lowest()); #ifdef SPIRV_1_3 test(q, input, output, - ONEAPI::multiplies(), 1); - test(q, input, output, ONEAPI::bit_or(), 0); + ext::oneapi::multiplies(), 1); + test(q, input, output, ext::oneapi::bit_or(), + 0); test(q, input, output, - ONEAPI::bit_xor(), 0); + ext::oneapi::bit_xor(), 0); test(q, input, output, - ONEAPI::bit_and(), ~0); + ext::oneapi::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/inclusive_scan.cpp b/SYCL/GroupAlgorithm/inclusive_scan.cpp index 9bc4991cb7..690ce0b21c 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan.cpp @@ -21,7 +21,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template class inclusive_scan_kernel; @@ -140,27 +140,27 @@ int main() { std::iota(input.begin(), input.end(), 0); std::fill(output.begin(), output.end(), 0); - test(q, input, output, ONEAPI::plus<>(), 0); - test(q, input, output, ONEAPI::minimum<>(), + test(q, input, output, ext::oneapi::plus<>(), 0); + test(q, input, output, ext::oneapi::minimum<>(), std::numeric_limits::max()); - test(q, input, output, ONEAPI::maximum<>(), + test(q, input, output, ext::oneapi::maximum<>(), std::numeric_limits::lowest()); - test(q, input, output, ONEAPI::plus(), 0); - test(q, input, output, ONEAPI::minimum(), + test(q, input, output, ext::oneapi::plus(), 0); + test(q, input, output, ext::oneapi::minimum(), std::numeric_limits::max()); - test(q, input, output, ONEAPI::maximum(), + test(q, input, output, ext::oneapi::maximum(), std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - ONEAPI::multiplies(), 1); + test( + q, input, output, ext::oneapi::multiplies(), 1); test(q, input, output, - ONEAPI::bit_or(), 0); + ext::oneapi::bit_or(), 0); test(q, input, output, - ONEAPI::bit_xor(), 0); + ext::oneapi::bit_xor(), 0); test(q, input, output, - ONEAPI::bit_and(), ~0); + ext::oneapi::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/leader.cpp b/SYCL/GroupAlgorithm/leader.cpp index e24fc4fbde..1a353db33a 100644 --- a/SYCL/GroupAlgorithm/leader.cpp +++ b/SYCL/GroupAlgorithm/leader.cpp @@ -7,7 +7,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; class leader_kernel; diff --git a/SYCL/GroupAlgorithm/none_of.cpp b/SYCL/GroupAlgorithm/none_of.cpp index 3df59a805e..d7f3d9d5b1 100644 --- a/SYCL/GroupAlgorithm/none_of.cpp +++ b/SYCL/GroupAlgorithm/none_of.cpp @@ -10,7 +10,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template class none_of_kernel; diff --git a/SYCL/GroupAlgorithm/ns_compat.h b/SYCL/GroupAlgorithm/ns_compat.h index 39cf703fee..0c61f26535 100644 --- a/SYCL/GroupAlgorithm/ns_compat.h +++ b/SYCL/GroupAlgorithm/ns_compat.h @@ -1,6 +1,6 @@ /* https://github.com/intel/llvm/pull/2231: namespace compatibility mode*/ __SYCL_INLINE_NAMESPACE(cl){namespace sycl{namespace intel{}; -namespace ONEAPI { +namespace ext::oneapi { using namespace cl::sycl::intel; } } diff --git a/SYCL/GroupAlgorithm/reduce.cpp b/SYCL/GroupAlgorithm/reduce.cpp index e9729b0441..740dabb7a1 100644 --- a/SYCL/GroupAlgorithm/reduce.cpp +++ b/SYCL/GroupAlgorithm/reduce.cpp @@ -20,7 +20,7 @@ #include #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; template @@ -76,26 +76,27 @@ int main() { std::iota(input.begin(), input.end(), 0); std::fill(output.begin(), output.end(), 0); - test(q, input, output, ONEAPI::plus<>(), 0); - test(q, input, output, ONEAPI::minimum<>(), + test(q, input, output, ext::oneapi::plus<>(), 0); + test(q, input, output, ext::oneapi::minimum<>(), std::numeric_limits::max()); - test(q, input, output, ONEAPI::maximum<>(), + test(q, input, output, ext::oneapi::maximum<>(), std::numeric_limits::lowest()); - test(q, input, output, ONEAPI::plus(), 0); - test(q, input, output, ONEAPI::minimum(), + test(q, input, output, ext::oneapi::plus(), 0); + test(q, input, output, ext::oneapi::minimum(), std::numeric_limits::max()); - test(q, input, output, ONEAPI::maximum(), + test(q, input, output, ext::oneapi::maximum(), std::numeric_limits::lowest()); #ifdef SPIRV_1_3 test(q, input, output, - ONEAPI::multiplies(), 1); + ext::oneapi::multiplies(), 1); test(q, input, output, - ONEAPI::bit_or(), 0); - test(q, input, output, ONEAPI::bit_xor(), 0); + ext::oneapi::bit_or(), 0); + test(q, input, output, ext::oneapi::bit_xor(), + 0); test(q, input, output, - ONEAPI::bit_and(), ~0); + ext::oneapi::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/support.h b/SYCL/GroupAlgorithm/support.h index c880958b2c..c2b33af04c 100644 --- a/SYCL/GroupAlgorithm/support.h +++ b/SYCL/GroupAlgorithm/support.h @@ -1,6 +1,6 @@ #include using namespace sycl; -using namespace sycl::ONEAPI; +using namespace sycl::ext::oneapi; bool isSupportedDevice(device D) { std::string PlatformName = D.get_platform().get_info(); diff --git a/SYCL/KernelAndProgram/spec_consts.hpp b/SYCL/KernelAndProgram/spec_consts.hpp index da4507e6a1..9ce468d749 100644 --- a/SYCL/KernelAndProgram/spec_consts.hpp +++ b/SYCL/KernelAndProgram/spec_consts.hpp @@ -17,9 +17,8 @@ int global_val = 10; // Fetch a value at runtime. int get_value() { return global_val; } -float foo( - const cl::sycl::ONEAPI::experimental::spec_constant - &f32) { +float foo(const cl::sycl::ext::oneapi::experimental::spec_constant< + float, MyFloatConst> &f32) { return f32; } @@ -28,13 +27,13 @@ struct SCWrapper { : SC1(p.set_spec_constant(4)), SC2(p.set_spec_constant(2)) {} - cl::sycl::ONEAPI::experimental::spec_constant SC1; - cl::sycl::ONEAPI::experimental::spec_constant SC2; + cl::sycl::ext::oneapi::experimental::spec_constant SC1; + cl::sycl::ext::oneapi::experimental::spec_constant SC2; }; // MyKernel is used to test default constructor using AccT = sycl::accessor; -using ScT = sycl::ONEAPI::experimental::spec_constant; +using ScT = sycl::ext::oneapi::experimental::spec_constant; struct MyKernel { MyKernel(AccT &Acc) : Acc(Acc) {} @@ -75,13 +74,13 @@ int main(int argc, char **argv) { // TODO make this floating point once supported by the compiler const float goldf = get_value(); - cl::sycl::ONEAPI::experimental::spec_constant i32 = - program1.set_spec_constant(goldi); + cl::sycl::ext::oneapi::experimental::spec_constant + i32 = program1.set_spec_constant(goldi); - cl::sycl::ONEAPI::experimental::spec_constant f32 = + cl::sycl::ext::oneapi::experimental::spec_constant f32 = program2.set_spec_constant(goldf); - cl::sycl::ONEAPI::experimental::spec_constant sc = + cl::sycl::ext::oneapi::experimental::spec_constant sc = program4.set_spec_constant(goldi); program1.build_with_kernel_type(); diff --git a/SYCL/Matrix/joint_matrix_bf16.cpp b/SYCL/Matrix/joint_matrix_bf16.cpp index 631a744323..20218f294f 100644 --- a/SYCL/Matrix/joint_matrix_bf16.cpp +++ b/SYCL/Matrix/joint_matrix_bf16.cpp @@ -72,16 +72,17 @@ void matrix_multiply(big_matrix &C, const auto sg_startx = global_idx - spmd_item.get_local_id(0); const auto sg_starty = global_idy - spmd_item.get_local_id(1); - ONEAPI::sub_group sg = spmd_item.get_sub_group(); - joint_matrix sub_a(sg); + ext::oneapi::sub_group sg = spmd_item.get_sub_group(); + joint_matrix sub_a( + sg); // For B, since current implementation does not support non-packed // layout, users need to specify the updated VNNI sizes along with // the packed_b layout. By default, the layout is row_major and size // is (TK, TN). - joint_matrix sub_b(sg); - joint_matrix sub_c(sg); + joint_matrix sub_c(sg); joint_matrix_load(sg, sub_c, accC.get_pointer() + (sg_startx * TM) * N + diff --git a/SYCL/OnlineCompiler/online_compiler_L0.cpp b/SYCL/OnlineCompiler/online_compiler_L0.cpp index 4624a93fb3..7a152423a2 100644 --- a/SYCL/OnlineCompiler/online_compiler_L0.cpp +++ b/SYCL/OnlineCompiler/online_compiler_L0.cpp @@ -6,12 +6,12 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %th.out // RUN: %HOST_RUN_PLACEHOLDER %th.out -// This test checks INTEL feature class online_compiler for Level-Zero. +// This test checks ext::intel feature class online_compiler for Level-Zero. // All Level-Zero specific code is kept here and the common part that can be // re-used by other backends is kept in online_compiler_common.hpp file. #include -#include +#include #include diff --git a/SYCL/OnlineCompiler/online_compiler_OpenCL.cpp b/SYCL/OnlineCompiler/online_compiler_OpenCL.cpp index c131f72733..87cf30016e 100644 --- a/SYCL/OnlineCompiler/online_compiler_OpenCL.cpp +++ b/SYCL/OnlineCompiler/online_compiler_OpenCL.cpp @@ -6,12 +6,12 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s %opencl_lib -o %th.out // RUN: %HOST_RUN_PLACEHOLDER %th.out -// This test checks INTEL feature class online_compiler for OpenCL. +// This test checks ext::intel feature class online_compiler for OpenCL. // All OpenCL specific code is kept here and the common part that can be // re-used by other backends is kept in online_compiler_common.hpp file. #include -#include +#include #include diff --git a/SYCL/OnlineCompiler/online_compiler_common.hpp b/SYCL/OnlineCompiler/online_compiler_common.hpp index 63458548e6..5cb5ed1755 100644 --- a/SYCL/OnlineCompiler/online_compiler_common.hpp +++ b/SYCL/OnlineCompiler/online_compiler_common.hpp @@ -1,5 +1,5 @@ #include -#include +#include #include #include @@ -25,7 +25,7 @@ void cm_kernel() { } )==="; -using namespace sycl::INTEL; +using namespace sycl::ext::intel; #ifdef RUN_KERNELS void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { @@ -56,7 +56,8 @@ int main(int argc, char **argv) { { // Compile and run a trivial OpenCL kernel. std::cout << "Test case1\n"; - sycl::INTEL::online_compiler + sycl::ext::intel::online_compiler< + sycl::ext::intel::source_language::opencl_c> Compiler; std::vector IL; try { @@ -79,7 +80,8 @@ int main(int argc, char **argv) { { // Compile and run a trivial OpenCL kernel using online_compiler() // constructor accepting SYCL device. std::cout << "Test case2\n"; - sycl::INTEL::online_compiler + sycl::ext::intel::online_compiler< + sycl::ext::intel::source_language::opencl_c> Compiler(Device); std::vector IL; try { @@ -99,7 +101,8 @@ int main(int argc, char **argv) { // PATHs to clangFEWrapper library properly. { // Compile a trivial CM kernel. std::cout << "Test case3\n"; - sycl::INTEL::online_compiler Compiler; + sycl::ext::intel::online_compiler + Compiler; try { std::vector IL = Compiler.compile(CMSource); @@ -113,7 +116,8 @@ int main(int argc, char **argv) { { // Compile a source with syntax errors. std::cout << "Test case4\n"; - sycl::INTEL::online_compiler + sycl::ext::intel::online_compiler< + sycl::ext::intel::source_language::opencl_c> Compiler; std::vector IL; bool TestPassed = false; @@ -133,7 +137,8 @@ int main(int argc, char **argv) { { // Compile a good CL source using unrecognized compilation options. std::cout << "Test case5\n"; - sycl::INTEL::online_compiler + sycl::ext::intel::online_compiler< + sycl::ext::intel::source_language::opencl_c> Compiler; std::vector IL; bool TestPassed = false; @@ -156,7 +161,8 @@ int main(int argc, char **argv) { { // Try compiling CM source with OpenCL compiler. std::cout << "Test case6\n"; - sycl::INTEL::online_compiler + sycl::ext::intel::online_compiler< + sycl::ext::intel::source_language::opencl_c> Compiler; std::vector IL; bool TestPassed = false; diff --git a/SYCL/Reduction/reduction_nd_N_vars.cpp b/SYCL/Reduction/reduction_nd_N_vars.cpp index 0f4b4ead79..71e8e25636 100644 --- a/SYCL/Reduction/reduction_nd_N_vars.cpp +++ b/SYCL/Reduction/reduction_nd_N_vars.cpp @@ -23,7 +23,7 @@ template void printNVarsTestLabel(bool IsSYCL2020, const RangeT &Range, bool ToCERR = false) { std::ostream &OS = ToCERR ? std::cerr : std::cout; - std::string Mode = IsSYCL2020 ? "SYCL2020" : "ONEAPI "; + std::string Mode = IsSYCL2020 ? "SYCL2020" : "ext::oneapi "; OS << (ToCERR ? "Error" : "Start") << ": Mode=" << Mode << ", Range=" << Range; if (!ToCERR) diff --git a/SYCL/Reduction/reduction_usm.cpp b/SYCL/Reduction/reduction_usm.cpp index e3c22b17d2..a75c3950ad 100644 --- a/SYCL/Reduction/reduction_usm.cpp +++ b/SYCL/Reduction/reduction_usm.cpp @@ -89,7 +89,7 @@ void testUSM(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) { NumErrors += test, true, T, BinaryOperation>( Q, Identity, Init, WGSize, NWItems, usm::alloc::device); - // Test ext::oneapi reductions + // Test ext::oneapi:: reductions NumErrors += test, false, T, BinaryOperation>( Q, Identity, Init, WGSize, NWItems, usm::alloc::shared); NumErrors += test, false, T, BinaryOperation>( diff --git a/SYCL/Reduction/reduction_utils.hpp b/SYCL/Reduction/reduction_utils.hpp index bf84148ee9..293570365a 100644 --- a/SYCL/Reduction/reduction_utils.hpp +++ b/SYCL/Reduction/reduction_utils.hpp @@ -188,7 +188,7 @@ std::ostream &operator<<(std::ostream &OS, const nd_range &Range) { template void printTestLabel(bool IsSYCL2020, const RangeT &Range, bool ToCERR = false) { std::ostream &OS = ToCERR ? std::cerr : std::cout; - std::string Mode = IsSYCL2020 ? "SYCL2020" : "ONEAPI "; + std::string Mode = IsSYCL2020 ? "SYCL2020" : "ext::oneapi "; OS << (ToCERR ? "Error" : "Start") << ": Mode=" << Mode << ", T=" << typeid(T).name() << ", BOp=" << typeid(BinaryOperation).name() << ", Range=" << Range; diff --git a/SYCL/SpecConstants/1.2.1/composite-in-functor.cpp b/SYCL/SpecConstants/1.2.1/composite-in-functor.cpp index 8b13e8b420..4f62d4211b 100644 --- a/SYCL/SpecConstants/1.2.1/composite-in-functor.cpp +++ b/SYCL/SpecConstants/1.2.1/composite-in-functor.cpp @@ -30,7 +30,7 @@ struct pod_t { class my_kernel_t { public: using sc_t = - sycl::ONEAPI::experimental::spec_constant; + sycl::ext::oneapi::experimental::spec_constant; my_kernel_t(const sc_t &sc, const cl::sycl::stream &strm) : sc_(sc), strm_(strm) {} diff --git a/SYCL/SpecConstants/1.2.1/composite-type.cpp b/SYCL/SpecConstants/1.2.1/composite-type.cpp index 873583a229..54a1a5fc55 100644 --- a/SYCL/SpecConstants/1.2.1/composite-type.cpp +++ b/SYCL/SpecConstants/1.2.1/composite-type.cpp @@ -59,7 +59,7 @@ int main(int argc, char **argv) { POD gold = {{{goldi, goldf}, {goldi, goldf}}, goldi}; - cl::sycl::ONEAPI::experimental::spec_constant pod = + cl::sycl::ext::oneapi::experimental::spec_constant pod = program.set_spec_constant(gold); program.build_with_kernel_type(); diff --git a/SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp b/SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp index 6bba43cd3d..3f2f678acb 100644 --- a/SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp +++ b/SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp @@ -26,7 +26,8 @@ struct pod_t { template class kernel_t { public: - using sc_t = sycl::ONEAPI::experimental::spec_constant; + using sc_t = + sycl::ext::oneapi::experimental::spec_constant; kernel_t(const sc_t &sc, cl::sycl::stream &strm) : sc_(sc), strm_(strm) {} diff --git a/SYCL/SpecConstants/1.2.1/spec_const_hw.cpp b/SYCL/SpecConstants/1.2.1/spec_const_hw.cpp index 2e1f271700..cee34d2952 100644 --- a/SYCL/SpecConstants/1.2.1/spec_const_hw.cpp +++ b/SYCL/SpecConstants/1.2.1/spec_const_hw.cpp @@ -35,9 +35,8 @@ int global_val = 10; // Fetch a value at runtime. int get_value() { return global_val; } -float foo( - const cl::sycl::ONEAPI::experimental::spec_constant - &f32) { +float foo(const cl::sycl::ext::oneapi::experimental::spec_constant< + float, MyFloatConst> &f32) { return f32; } @@ -46,13 +45,13 @@ struct SCWrapper { : SC1(p.set_spec_constant(4)), SC2(p.set_spec_constant(2)) {} - cl::sycl::ONEAPI::experimental::spec_constant SC1; - cl::sycl::ONEAPI::experimental::spec_constant SC2; + cl::sycl::ext::oneapi::experimental::spec_constant SC1; + cl::sycl::ext::oneapi::experimental::spec_constant SC2; }; // MyKernel is used to test default constructor using AccT = sycl::accessor; -using ScT = sycl::ONEAPI::experimental::spec_constant; +using ScT = sycl::ext::oneapi::experimental::spec_constant; struct MyKernel { MyKernel(AccT &Acc) : Acc(Acc) {} @@ -93,13 +92,13 @@ int main(int argc, char **argv) { // TODO make this floating point once supported by the compiler float goldf = (float)get_value(); - cl::sycl::ONEAPI::experimental::spec_constant i32 = - program1.set_spec_constant(goldi); + cl::sycl::ext::oneapi::experimental::spec_constant + i32 = program1.set_spec_constant(goldi); - cl::sycl::ONEAPI::experimental::spec_constant f32 = + cl::sycl::ext::oneapi::experimental::spec_constant f32 = program2.set_spec_constant(goldf); - cl::sycl::ONEAPI::experimental::spec_constant sc = + cl::sycl::ext::oneapi::experimental::spec_constant sc = program4.set_spec_constant(goldi); program1.build_with_kernel_type(); diff --git a/SYCL/SpecConstants/1.2.1/spec_const_neg.cpp b/SYCL/SpecConstants/1.2.1/spec_const_neg.cpp index fb0f9d1aa1..6d270bbb43 100644 --- a/SYCL/SpecConstants/1.2.1/spec_const_neg.cpp +++ b/SYCL/SpecConstants/1.2.1/spec_const_neg.cpp @@ -45,8 +45,8 @@ int main(int argc, char **argv) { << "\n"; cl::sycl::program program1(q.get_context()); - cl::sycl::ONEAPI::experimental::spec_constant i32 = - program1.set_spec_constant(10); + cl::sycl::ext::oneapi::experimental::spec_constant + i32 = program1.set_spec_constant(10); std::vector veci(1); bool passed = false; @@ -56,8 +56,8 @@ int main(int argc, char **argv) { try { // This is an attempt to set a spec constant after the program has been // built - spec_const_error should be thrown - cl::sycl::ONEAPI::experimental::spec_constant i32 = - program1.set_spec_constant(10); + cl::sycl::ext::oneapi::experimental::spec_constant + i32 = program1.set_spec_constant(10); cl::sycl::buffer bufi(veci.data(), veci.size()); @@ -66,7 +66,7 @@ int main(int argc, char **argv) { cgh.single_task(program1.get_kernel(), [=]() { acci[0] = i32.get(); }); }); - } catch (cl::sycl::ONEAPI::experimental::spec_const_error &sc_err) { + } catch (cl::sycl::ext::oneapi::experimental::spec_const_error &sc_err) { passed = true; } catch (cl::sycl::exception &e) { std::cout << "*** Exception caught: " << e.what() << "\n"; diff --git a/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp b/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp index b983258fd0..f1990c2a42 100644 --- a/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp +++ b/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp @@ -70,9 +70,9 @@ int main(int argc, char **argv) { for (int i = 0; i < n_sc_sets; i++) { cl::sycl::program program(q.get_context()); const int *sc_set = &sc_vals[i][0]; - cl::sycl::ONEAPI::experimental::spec_constant sc0 = + cl::sycl::ext::oneapi::experimental::spec_constant sc0 = program.set_spec_constant(sc_set[0]); - cl::sycl::ONEAPI::experimental::spec_constant sc1 = + cl::sycl::ext::oneapi::experimental::spec_constant sc1 = program.set_spec_constant(sc_set[1]); program.build_with_kernel_type(); diff --git a/SYCL/SpecConstants/1.2.1/specialization_constants.cpp b/SYCL/SpecConstants/1.2.1/specialization_constants.cpp index c281baa64d..e7604df102 100644 --- a/SYCL/SpecConstants/1.2.1/specialization_constants.cpp +++ b/SYCL/SpecConstants/1.2.1/specialization_constants.cpp @@ -84,32 +84,32 @@ int main(int argc, char **argv) { program prog(q.get_context()); // Create specialization constants. - ONEAPI::experimental::spec_constant i1 = + ext::oneapi::experimental::spec_constant i1 = prog.set_spec_constant(bool_ref); - ONEAPI::experimental::spec_constant i8 = + ext::oneapi::experimental::spec_constant i8 = prog.set_spec_constant(int8_ref); - ONEAPI::experimental::spec_constant ui8 = + ext::oneapi::experimental::spec_constant ui8 = prog.set_spec_constant(uint8_ref); - ONEAPI::experimental::spec_constant i16 = + ext::oneapi::experimental::spec_constant i16 = prog.set_spec_constant(int16_ref); - ONEAPI::experimental::spec_constant ui16 = + ext::oneapi::experimental::spec_constant ui16 = prog.set_spec_constant(uint16_ref); - ONEAPI::experimental::spec_constant i32 = + ext::oneapi::experimental::spec_constant i32 = prog.set_spec_constant(int32_ref); - ONEAPI::experimental::spec_constant ui32 = + ext::oneapi::experimental::spec_constant ui32 = prog.set_spec_constant(uint32_ref); - ONEAPI::experimental::spec_constant i64 = + ext::oneapi::experimental::spec_constant i64 = prog.set_spec_constant(int64_ref); - ONEAPI::experimental::spec_constant ui64 = + ext::oneapi::experimental::spec_constant ui64 = prog.set_spec_constant(uint64_ref); #if HALF - ONEAPI::experimental::spec_constant f16 = + ext::oneapi::experimental::spec_constant f16 = prog.set_spec_constant(half_ref); #endif - ONEAPI::experimental::spec_constant f32 = + ext::oneapi::experimental::spec_constant f32 = prog.set_spec_constant(float_ref); - ONEAPI::experimental::spec_constant f64 = + ext::oneapi::experimental::spec_constant f64 = prog.set_spec_constant(double_ref); prog.build_with_kernel_type(); diff --git a/SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp b/SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp index 318d0db48e..28920badc2 100644 --- a/SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp +++ b/SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp @@ -64,8 +64,8 @@ int main(int argc, char **argv) { program prog(q.get_context()); // Create specialization constants. - ONEAPI::experimental::spec_constant ui32 = - prog.set_spec_constant(uint32_ref); + ext::oneapi::experimental::spec_constant + ui32 = prog.set_spec_constant(uint32_ref); prog.build_with_kernel_type(); @@ -78,7 +78,7 @@ int main(int argc, char **argv) { bool exception_was_thrown = false; try { ui32 = prog.set_spec_constant(uint32_ref + 1); - } catch (const ONEAPI::experimental::spec_const_error &e) { + } catch (const ext::oneapi::experimental::spec_const_error &e) { exception_was_thrown = true; } if (!exception_was_thrown) { diff --git a/SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp b/SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp index efd164fed1..d6ce5e40d3 100644 --- a/SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp +++ b/SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp @@ -71,12 +71,12 @@ int main(int argc, char **argv) { program prog(q.get_context()); // Create specialization constants. - ONEAPI::experimental::spec_constant i1 = + ext::oneapi::experimental::spec_constant i1 = prog.set_spec_constant(bool_ref); - ONEAPI::experimental::spec_constant ui32 = - prog.set_spec_constant(uint32_ref); - ONEAPI::experimental::spec_constant f64 = - prog.set_spec_constant(double_ref); + ext::oneapi::experimental::spec_constant + ui32 = prog.set_spec_constant(uint32_ref); + ext::oneapi::experimental::spec_constant + f64 = prog.set_spec_constant(double_ref); // Override specialization constants. i1 = prog.set_spec_constant(bool_ref_override); diff --git a/SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp b/SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp index 11029c6739..16180fc748 100644 --- a/SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp +++ b/SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp @@ -31,7 +31,8 @@ struct pod_t { template class kernel_t { public: - using sc_t = sycl::ONEAPI::experimental::spec_constant; + using sc_t = + sycl::ext::oneapi::experimental::spec_constant; kernel_t(const sc_t &sc, cl::sycl::stream &strm) : sc_(sc), strm_(strm) {} diff --git a/SYCL/SubGroup/barrier.cpp b/SYCL/SubGroup/barrier.cpp index 0f32cab3b4..d973e587dd 100644 --- a/SYCL/SubGroup/barrier.cpp +++ b/SYCL/SubGroup/barrier.cpp @@ -33,7 +33,7 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { cgh.parallel_for>( NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); size_t lid = SG.get_local_id().get(0); size_t gid = NdItem.get_global_id(0); size_t SGoff = gid - lid; diff --git a/SYCL/SubGroup/broadcast.hpp b/SYCL/SubGroup/broadcast.hpp index ed9803cef3..67b932cf53 100644 --- a/SYCL/SubGroup/broadcast.hpp +++ b/SYCL/SubGroup/broadcast.hpp @@ -20,11 +20,11 @@ template void check(queue &Queue) { auto syclacc = syclbuf.template get_access(cgh); auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); /*Broadcast GID of element with SGLID == SGID % SGMLR*/ - syclacc[NdItem.get_global_id()] = - ONEAPI::broadcast(SG, T(NdItem.get_global_id(0)), - SG.get_group_id() % SG.get_max_local_range()[0]); + syclacc[NdItem.get_global_id()] = ext::oneapi::broadcast( + SG, T(NdItem.get_global_id(0)), + SG.get_group_id() % SG.get_max_local_range()[0]); if (NdItem.get_global_id(0) == 0) sgsizeacc[0] = SG.get_max_local_range()[0]; }); diff --git a/SYCL/SubGroup/common.cpp b/SYCL/SubGroup/common.cpp index ec2ba7f58b..8ddba4a08d 100644 --- a/SYCL/SubGroup/common.cpp +++ b/SYCL/SubGroup/common.cpp @@ -33,7 +33,7 @@ void check(queue &Queue, unsigned int G, unsigned int L) { auto sgsizeacc = sgsizebuf.get_access(cgh); auto syclacc = syclbuf.get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); syclacc[NdItem.get_global_id()].local_id = SG.get_local_id().get(0); syclacc[NdItem.get_global_id()].local_range = SG.get_local_range().get(0); diff --git a/SYCL/SubGroup/generic-shuffle.cpp b/SYCL/SubGroup/generic-shuffle.cpp index e47908b42c..5db1b56511 100644 --- a/SYCL/SubGroup/generic-shuffle.cpp +++ b/SYCL/SubGroup/generic-shuffle.cpp @@ -40,7 +40,7 @@ void check_pointer(queue &Queue, size_t G = 256, size_t L = 64) { cgh.parallel_for( NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); uint32_t wggid = NdItem.get_global_id(0); uint32_t sgid = SG.get_group_id().get(0); if (wggid == 0) @@ -142,7 +142,7 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 256, size_t L = 64) { cgh.parallel_for( NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); uint32_t wggid = NdItem.get_global_id(0); uint32_t sgid = SG.get_group_id().get(0); if (wggid == 0) diff --git a/SYCL/SubGroup/generic_reduce.cpp b/SYCL/SubGroup/generic_reduce.cpp index c0da01e718..b0efa498ef 100644 --- a/SYCL/SubGroup/generic_reduce.cpp +++ b/SYCL/SubGroup/generic_reduce.cpp @@ -22,13 +22,13 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false, auto sgsizeacc = sgsizebuf.get_access(cgh); auto acc = buf.template get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group sg = NdItem.get_sub_group(); + ext::oneapi::sub_group sg = NdItem.get_sub_group(); if (skip_init) { acc[NdItem.get_global_id(0)] = - ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), op); + ext::oneapi::reduce(sg, T(NdItem.get_global_id(0)), op); } else { acc[NdItem.get_global_id(0)] = - ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), init, op); + ext::oneapi::reduce(sg, T(NdItem.get_global_id(0)), init, op); } if (NdItem.get_global_id(0) == 0) sgsizeacc[0] = sg.get_max_local_range()[0]; @@ -74,8 +74,8 @@ int main() { // Test user-defined type // Use complex as a proxy for this using UDT = std::complex; - check_op(Queue, UDT(L, L), ONEAPI::plus(), false, G, L); - check_op(Queue, UDT(0, 0), ONEAPI::plus(), true, G, L); + check_op(Queue, UDT(L, L), ext::oneapi::plus(), false, G, L); + check_op(Queue, UDT(0, 0), ext::oneapi::plus(), true, G, L); // Test user-defined operator auto UDOp = [=](const auto &lhs, const auto &rhs) { return lhs + rhs; }; diff --git a/SYCL/SubGroup/load_store.cpp b/SYCL/SubGroup/load_store.cpp index 1ad742fbef..bbf02b19d1 100644 --- a/SYCL/SubGroup/load_store.cpp +++ b/SYCL/SubGroup/load_store.cpp @@ -46,7 +46,7 @@ template void check(queue &Queue) { accessor LocalMem( {L + max_sg_size * N}, cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); auto SGid = SG.get_group_id().get(0); auto SGsize = SG.get_max_local_range().get(0); /* Avoid overlapping data ranges inside and between local groups */ @@ -134,7 +134,7 @@ template void check(queue &Queue) { accessor LocalMem( {L}, cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); if (NdItem.get_global_id(0) == 0) sgsizeacc[0] = SG.get_max_local_range()[0]; size_t SGOffset = diff --git a/SYCL/SubGroup/reduce.hpp b/SYCL/SubGroup/reduce.hpp index 7a3cbac095..68605785da 100644 --- a/SYCL/SubGroup/reduce.hpp +++ b/SYCL/SubGroup/reduce.hpp @@ -25,13 +25,13 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false, auto acc = buf.template get_access(cgh); cgh.parallel_for( NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group sg = NdItem.get_sub_group(); + ext::oneapi::sub_group sg = NdItem.get_sub_group(); if (skip_init) { acc[NdItem.get_global_id(0)] = - ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), op); + ext::oneapi::reduce(sg, T(NdItem.get_global_id(0)), op); } else { acc[NdItem.get_global_id(0)] = - ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), init, op); + ext::oneapi::reduce(sg, T(NdItem.get_global_id(0)), init, op); } if (NdItem.get_global_id(0) == 0) sgsizeacc[0] = sg.get_max_local_range()[0]; @@ -74,85 +74,85 @@ void check(queue &Queue, size_t G = 256, size_t L = 64) { check_op< sycl_subgr, T>( - Queue, T(L), ONEAPI::plus(), false, G, L); + Queue, T(L), ext::oneapi::plus(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::plus(), true, G, L); + Queue, T(0), ext::oneapi::plus(), true, G, L); check_op, - T>(Queue, T(0), ONEAPI::minimum(), false, G, L); + T>(Queue, T(0), ext::oneapi::minimum(), false, G, L); check_op, - T>(Queue, T(G), ONEAPI::minimum(), true, G, L); + T>(Queue, T(G), ext::oneapi::minimum(), true, G, L); check_op, - T>(Queue, T(G), ONEAPI::maximum(), false, G, L); + T>(Queue, T(G), ext::oneapi::maximum(), false, G, L); check_op, - T>(Queue, T(0), ONEAPI::maximum(), true, G, L); + T>(Queue, T(0), ext::oneapi::maximum(), true, G, L); // Transparent operator functors. check_op, - T>(Queue, T(L), ONEAPI::plus<>(), false, G, L); + T>(Queue, T(L), ext::oneapi::plus<>(), false, G, L); check_op, - T>(Queue, T(0), ONEAPI::plus<>(), true, G, L); + T>(Queue, T(0), ext::oneapi::plus<>(), true, G, L); check_op, - T>(Queue, T(0), ONEAPI::minimum<>(), false, G, L); + T>(Queue, T(0), ext::oneapi::minimum<>(), false, G, L); check_op, - T>(Queue, T(G), ONEAPI::minimum<>(), true, G, L); + T>(Queue, T(G), ext::oneapi::minimum<>(), true, G, L); check_op, - T>(Queue, T(G), ONEAPI::maximum<>(), false, G, L); + T>(Queue, T(G), ext::oneapi::maximum<>(), false, G, L); check_op< sycl_subgr, - T>(Queue, T(0), ONEAPI::maximum<>(), true, G, L); + T>(Queue, T(0), ext::oneapi::maximum<>(), true, G, L); } template void check_mul(queue &Queue, size_t G = 256, size_t L = 4) { check_op, T>( - Queue, T(G), ONEAPI::multiplies(), false, G, L); + Queue, T(G), ext::oneapi::multiplies(), false, G, L); check_op, T>( - Queue, T(1), ONEAPI::multiplies(), true, G, L); + Queue, T(1), ext::oneapi::multiplies(), true, G, L); // Transparent operator functors. check_op, T>( - Queue, T(G), ONEAPI::multiplies<>(), false, G, L); + Queue, T(G), ext::oneapi::multiplies<>(), false, G, L); check_op, T>( - Queue, T(1), ONEAPI::multiplies<>(), true, G, L); + Queue, T(1), ext::oneapi::multiplies<>(), true, G, L); } template void check_bit_ops(queue &Queue, size_t G = 256, size_t L = 4) { check_op, T>( - Queue, T(G), ONEAPI::bit_or(), false, G, L); + Queue, T(G), ext::oneapi::bit_or(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_or(), true, G, L); + Queue, T(0), ext::oneapi::bit_or(), true, G, L); check_op, T>( - Queue, T(G), ONEAPI::bit_xor(), false, G, L); + Queue, T(G), ext::oneapi::bit_xor(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_xor(), true, G, L); + Queue, T(0), ext::oneapi::bit_xor(), true, G, L); check_op, T>( - Queue, T(G), ONEAPI::bit_and(), false, G, L); + Queue, T(G), ext::oneapi::bit_and(), false, G, L); check_op, T>( - Queue, ~T(0), ONEAPI::bit_and(), true, G, L); + Queue, ~T(0), ext::oneapi::bit_and(), true, G, L); // Transparent operator functors check_op, T>( - Queue, T(G), ONEAPI::bit_or(), false, G, L); + Queue, T(G), ext::oneapi::bit_or(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_or(), true, G, L); + Queue, T(0), ext::oneapi::bit_or(), true, G, L); check_op, T>( - Queue, T(G), ONEAPI::bit_xor(), false, G, L); + Queue, T(G), ext::oneapi::bit_xor(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_xor(), true, G, L); + Queue, T(0), ext::oneapi::bit_xor(), true, G, L); check_op, T>( - Queue, T(G), ONEAPI::bit_and(), false, G, L); + Queue, T(G), ext::oneapi::bit_and(), false, G, L); check_op, T>( - Queue, ~T(0), ONEAPI::bit_and(), true, G, L); + Queue, ~T(0), ext::oneapi::bit_and(), true, G, L); } diff --git a/SYCL/SubGroup/scan.hpp b/SYCL/SubGroup/scan.hpp index 5d9c00dd20..0f3ae4d806 100644 --- a/SYCL/SubGroup/scan.hpp +++ b/SYCL/SubGroup/scan.hpp @@ -27,16 +27,16 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false, auto inacc = inbuf.template get_access(cgh); cgh.parallel_for( NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group sg = NdItem.get_sub_group(); + ext::oneapi::sub_group sg = NdItem.get_sub_group(); if (skip_init) { - exacc[NdItem.get_global_id(0)] = - ONEAPI::exclusive_scan(sg, T(NdItem.get_global_id(0)), op); - inacc[NdItem.get_global_id(0)] = - ONEAPI::inclusive_scan(sg, T(NdItem.get_global_id(0)), op); + exacc[NdItem.get_global_id(0)] = ext::oneapi::exclusive_scan( + sg, T(NdItem.get_global_id(0)), op); + inacc[NdItem.get_global_id(0)] = ext::oneapi::inclusive_scan( + sg, T(NdItem.get_global_id(0)), op); } else { - exacc[NdItem.get_global_id(0)] = ONEAPI::exclusive_scan( + exacc[NdItem.get_global_id(0)] = ext::oneapi::exclusive_scan( sg, T(NdItem.get_global_id(0)), init, op); - inacc[NdItem.get_global_id(0)] = ONEAPI::inclusive_scan( + inacc[NdItem.get_global_id(0)] = ext::oneapi::inclusive_scan( sg, T(NdItem.get_global_id(0)), op, init); } if (NdItem.get_global_id(0) == 0) @@ -81,120 +81,121 @@ void check(queue &Queue, size_t G = 256, size_t L = 64) { } check_op, - T>(Queue, T(L), ONEAPI::plus(), false, G, L); + T>(Queue, T(L), ext::oneapi::plus(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::plus(), true, G, L); + Queue, T(0), ext::oneapi::plus(), true, G, L); check_op< sycl_subgr, - T>(Queue, T(0), ONEAPI::minimum(), false, G, L); + T>(Queue, T(0), ext::oneapi::minimum(), false, G, L); if (std::is_floating_point::value || std::is_same::value) { check_op< sycl_subgr, - T>(Queue, std::numeric_limits::infinity(), ONEAPI::minimum(), + T>(Queue, std::numeric_limits::infinity(), ext::oneapi::minimum(), true, G, L); } else { check_op, - T>(Queue, std::numeric_limits::max(), ONEAPI::minimum(), + T>(Queue, std::numeric_limits::max(), ext::oneapi::minimum(), true, G, L); } check_op< sycl_subgr, - T>(Queue, T(G), ONEAPI::maximum(), false, G, L); + T>(Queue, T(G), ext::oneapi::maximum(), false, G, L); if (std::is_floating_point::value || std::is_same::value) { check_op< sycl_subgr, - T>(Queue, -std::numeric_limits::infinity(), ONEAPI::maximum(), - true, G, L); + T>(Queue, -std::numeric_limits::infinity(), + ext::oneapi::maximum(), true, G, L); } else { check_op, T>( - Queue, std::numeric_limits::min(), ONEAPI::maximum(), true, G, L); + Queue, std::numeric_limits::min(), ext::oneapi::maximum(), true, + G, L); } // Transparent operator functors. check_op, T>( - Queue, T(L), ONEAPI::plus<>(), false, G, L); + Queue, T(L), ext::oneapi::plus<>(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::plus<>(), true, G, L); + Queue, T(0), ext::oneapi::plus<>(), true, G, L); check_op< sycl_subgr, - T>(Queue, T(0), ONEAPI::minimum<>(), false, G, L); + T>(Queue, T(0), ext::oneapi::minimum<>(), false, G, L); if (std::is_floating_point::value || std::is_same::value) { check_op< sycl_subgr, - T>(Queue, std::numeric_limits::infinity(), ONEAPI::minimum<>(), true, - G, L); + T>(Queue, std::numeric_limits::infinity(), ext::oneapi::minimum<>(), + true, G, L); } else { check_op, - T>(Queue, std::numeric_limits::max(), ONEAPI::minimum<>(), true, - G, L); + T>(Queue, std::numeric_limits::max(), ext::oneapi::minimum<>(), + true, G, L); } check_op, T>( - Queue, T(G), ONEAPI::maximum<>(), false, G, L); + Queue, T(G), ext::oneapi::maximum<>(), false, G, L); if (std::is_floating_point::value || std::is_same::value) { check_op, T>( - Queue, -std::numeric_limits::infinity(), ONEAPI::maximum<>(), true, - G, L); + Queue, -std::numeric_limits::infinity(), ext::oneapi::maximum<>(), + true, G, L); } else { check_op< sycl_subgr, - T>(Queue, std::numeric_limits::min(), ONEAPI::maximum<>(), true, G, - L); + T>(Queue, std::numeric_limits::min(), ext::oneapi::maximum<>(), true, + G, L); } } template void check_mul(queue &Queue, size_t G = 256, size_t L = 4) { check_op, T>( - Queue, T(L), ONEAPI::multiplies(), false, G, L); + Queue, T(L), ext::oneapi::multiplies(), false, G, L); check_op, T>( - Queue, T(1), ONEAPI::multiplies<>(), true, G, L); + Queue, T(1), ext::oneapi::multiplies<>(), true, G, L); check_op, T>( - Queue, T(L), ONEAPI::multiplies(), false, G, L); + Queue, T(L), ext::oneapi::multiplies(), false, G, L); check_op, T>( - Queue, T(1), ONEAPI::multiplies<>(), true, G, L); + Queue, T(1), ext::oneapi::multiplies<>(), true, G, L); } template void check_bit_ops(queue &Queue, size_t G = 256, size_t L = 4) { check_op, T>( - Queue, T(L), ONEAPI::bit_or(), false, G, L); + Queue, T(L), ext::oneapi::bit_or(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_or(), true, G, L); + Queue, T(0), ext::oneapi::bit_or(), true, G, L); check_op, T>( - Queue, T(L), ONEAPI::bit_xor(), false, G, L); + Queue, T(L), ext::oneapi::bit_xor(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_xor(), true, G, L); + Queue, T(0), ext::oneapi::bit_xor(), true, G, L); check_op, T>( - Queue, T(L), ONEAPI::bit_and(), false, G, L); + Queue, T(L), ext::oneapi::bit_and(), false, G, L); check_op, T>( - Queue, ~T(0), ONEAPI::bit_and(), true, G, L); + Queue, ~T(0), ext::oneapi::bit_and(), true, G, L); // Transparent operator functors. check_op, T>( - Queue, T(L), ONEAPI::bit_or<>(), false, G, L); + Queue, T(L), ext::oneapi::bit_or<>(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_or<>(), true, G, L); + Queue, T(0), ext::oneapi::bit_or<>(), true, G, L); check_op, T>( - Queue, T(L), ONEAPI::bit_xor<>(), false, G, L); + Queue, T(L), ext::oneapi::bit_xor<>(), false, G, L); check_op, T>( - Queue, T(0), ONEAPI::bit_xor<>(), true, G, L); + Queue, T(0), ext::oneapi::bit_xor<>(), true, G, L); check_op, T>( - Queue, T(L), ONEAPI::bit_and<>(), false, G, L); + Queue, T(L), ext::oneapi::bit_and<>(), false, G, L); check_op, T>( - Queue, ~T(0), ONEAPI::bit_and<>(), true, G, L); + Queue, ~T(0), ext::oneapi::bit_and<>(), true, G, L); } diff --git a/SYCL/SubGroup/shuffle.hpp b/SYCL/SubGroup/shuffle.hpp index ff18284d3f..8c84e971e5 100644 --- a/SYCL/SubGroup/shuffle.hpp +++ b/SYCL/SubGroup/shuffle.hpp @@ -34,7 +34,7 @@ void check(queue &Queue, size_t G = 256, size_t L = 64) { auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); uint32_t wggid = NdItem.get_global_id(0); uint32_t sgid = SG.get_group_id().get(0); vec vwggid(wggid), vsgid(sgid); @@ -117,7 +117,7 @@ template void check(queue &Queue, size_t G = 256, size_t L = 64) { auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); uint32_t wggid = NdItem.get_global_id(0); uint32_t sgid = SG.get_group_id().get(0); if (wggid == 0) diff --git a/SYCL/SubGroup/sub_group_as.cpp b/SYCL/SubGroup/sub_group_as.cpp index e2eb92f0b5..72968f343b 100644 --- a/SYCL/SubGroup/sub_group_as.cpp +++ b/SYCL/SubGroup/sub_group_as.cpp @@ -35,7 +35,7 @@ int main(int argc, char *argv[]) { cgh.parallel_for( cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) { - cl::sycl::ONEAPI::sub_group sg = it.get_sub_group(); + cl::sycl::ext::oneapi::sub_group sg = it.get_sub_group(); if (!it.get_local_id(0)) { int end = it.get_global_id(0) + it.get_local_range()[0]; for (int i = it.get_global_id(0); i < end; i++) { diff --git a/SYCL/SubGroup/sub_group_as_vec.cpp b/SYCL/SubGroup/sub_group_as_vec.cpp index c4bd23a479..7a99d5d85d 100644 --- a/SYCL/SubGroup/sub_group_as_vec.cpp +++ b/SYCL/SubGroup/sub_group_as_vec.cpp @@ -36,7 +36,7 @@ int main(int argc, char *argv[]) { local(N, cgh); cgh.parallel_for( cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) { - cl::sycl::ONEAPI::sub_group sg = it.get_sub_group(); + cl::sycl::ext::oneapi::sub_group sg = it.get_sub_group(); if (!it.get_local_id(0)) { int end = it.get_global_id(0) + it.get_local_range()[0]; for (int i = it.get_global_id(0); i < end; i++) { diff --git a/SYCL/SubGroup/vote.cpp b/SYCL/SubGroup/vote.cpp index 07172e5a3c..4ecf7f3d0e 100644 --- a/SYCL/SubGroup/vote.cpp +++ b/SYCL/SubGroup/vote.cpp @@ -46,14 +46,14 @@ void check(queue Queue, const int G, const int L, const int D, const int R) { auto sganyacc = sganybuf.get_access(cgh); auto sgallacc = sgallbuf.get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { - ONEAPI::sub_group SG = NdItem.get_sub_group(); + ext::oneapi::sub_group SG = NdItem.get_sub_group(); /* Set to 1 if any local ID in subgroup devided by D has remainder R */ - if (ONEAPI::any_of(SG, SG.get_local_id().get(0) % D == R)) { + if (ext::oneapi::any_of(SG, SG.get_local_id().get(0) % D == R)) { sganyacc[NdItem.get_global_id()] = 1; } /* Set to 1 if remainder of division of subgroup local ID by D is less * than R for all work items in subgroup */ - if (ONEAPI::all_of(SG, SG.get_local_id().get(0) % D < R)) { + if (ext::oneapi::all_of(SG, SG.get_local_id().get(0) % D < R)) { sgallacc[NdItem.get_global_id()] = 1; } });