[SYCL] Replace using of ONEAPI/INTEL namespaces with ext::oneapi/intel

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<int32_t, MyInt32Const> i32 =
-      prog.set_spec_constant<MyInt32Const>(10);
+  cl::sycl::ext::oneapi::experimental::spec_constant<int32_t, MyInt32Const>
+      i32 = prog.set_spec_constant<MyInt32Const>(10);
 
   prog.build_with_kernel_type<Kernel>();
 

SYCL/AtomicRef/accessor.cpp

@@ -10,7 +10,7 @@
 #include <numeric>
 #include <vector>
 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

SYCL/AtomicRef/assignment.h

@@ -21,8 +21,8 @@ template <typename T> void assignment_test(queue q, size_t N) {
           assignment_buf.template get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<assignment_kernel<T>>(range<1>(N), [=](item<1> it) {
         size_t gid = it.get_id(0);
-        auto atm = atomic_ref<T, ONEAPI::memory_order::relaxed,
-                              ONEAPI::memory_scope::device,
+        auto atm = atomic_ref<T, ext::oneapi::memory_order::relaxed,
+                              ext::oneapi::memory_scope::device,
                               access::address_space::global_space>(st[0]);
         atm = T(gid);
       });

SYCL/Basic/enqueue_barrier.cpp

@@ -8,7 +8,7 @@
 // UNSUPPORTED: cuda || windows
 
 #include <CL/sycl.hpp>
-#include <CL/sycl/INTEL/fpga_device_selector.hpp>
+#include <sycl/ext/intel/fpga_device_selector.hpp>
 
 int main() {
   sycl::context Context;

SYCL/Basic/fpga_tests/buffer_location.cpp

@@ -8,7 +8,8 @@ int main() {
   sycl::buffer<int, 1> 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<class Test>([=]() { Acc[0] = 42; });
   });

SYCL/Basic/fpga_tests/fpga_io_pipes.cpp

@@ -8,9 +8,9 @@
 //
 //===----------------------------------------------------------------------===//
 #include <CL/sycl.hpp>
-#include <CL/sycl/INTEL/fpga_extensions.hpp>
 #include <fstream>
 #include <iostream>
+#include <sycl/ext/intel/fpga_extensions.hpp>
 
 #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<cl::sycl::info::device::kernel_kernel_pipe_support>()) {
-    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;
   }

SYCL/Basic/fpga_tests/fpga_lsu.cpp

@@ -8,7 +8,7 @@
 //
 //===----------------------------------------------------------------------===//
 #include <CL/sycl.hpp>
-#include <CL/sycl/INTEL/fpga_extensions.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
 
 // 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::prefetch<true>,
-                                 cl::sycl::INTEL::statically_coalesce<false>>;
+        using PrefetchingLSU = cl::sycl::ext::intel::lsu<
+            cl::sycl::ext::intel::prefetch<true>,
+            cl::sycl::ext::intel::statically_coalesce<false>>;
 
-        using BurstCoalescedLSU =
-            cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>,
-                                 cl::sycl::INTEL::statically_coalesce<false>>;
+        using BurstCoalescedLSU = cl::sycl::ext::intel::lsu<
+            cl::sycl::ext::intel::burst_coalesce<true>,
+            cl::sycl::ext::intel::statically_coalesce<false>>;
 
-        using CachingLSU =
-            cl::sycl::INTEL::lsu<cl::sycl::INTEL::burst_coalesce<true>,
-                                 cl::sycl::INTEL::cache<1024>,
-                                 cl::sycl::INTEL::statically_coalesce<false>>;
+        using CachingLSU = cl::sycl::ext::intel::lsu<
+            cl::sycl::ext::intel::burst_coalesce<true>,
+            cl::sycl::ext::intel::cache<1024>,
+            cl::sycl::ext::intel::statically_coalesce<false>>;
 
-        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);
 }

SYCL/Basic/fpga_tests/fpga_pipes.cpp

@@ -12,8 +12,8 @@
 //
 //===----------------------------------------------------------------------===//
 #include <CL/sycl.hpp>
-#include <CL/sycl/INTEL/fpga_extensions.hpp>
 #include <iostream>
+#include <sycl/ext/intel/fpga_extensions.hpp>
 
 // Size of an array passing through a pipe
 constexpr size_t N = 10;
@@ -31,7 +31,7 @@ template <int N> class templ_nb_pipe;
 
 // For non-blocking multiple pipes
 template <int N>
-using PipeMulNb = cl::sycl::INTEL::pipe<class templ_nb_pipe<N>, int>;
+using PipeMulNb = cl::sycl::ext::intel::pipe<class templ_nb_pipe<N>, int>;
 
 // For simple blocking pipes with explicit type
 class some_bl_pipe;
@@ -46,7 +46,7 @@ template <int N> class templ_bl_pipe;
 
 // For blocking multiple pipes
 template <int N>
-using PipeMulBl = cl::sycl::INTEL::pipe<class templ_bl_pipe<N>, int>;
+using PipeMulBl = cl::sycl::ext::intel::pipe<class templ_bl_pipe<N>, int>;
 
 // Kernel names
 template <int TestNumber, int KernelNumber = 0> class writer;
@@ -57,7 +57,7 @@ template <typename PipeName, int TestNumber>
 int test_simple_nb_pipe(cl::sycl::queue Queue) {
   int data[] = {0};
 
-  using Pipe = cl::sycl::INTEL::pipe<PipeName, int>;
+  using Pipe = cl::sycl::ext::intel::pipe<PipeName, int>;
 
   cl::sycl::buffer<int, 1> readBuf(data, 1);
   Queue.submit([&](cl::sycl::handler &cgh) {
@@ -146,7 +146,7 @@ template <int TestNumber> int test_multiple_nb_pipe(cl::sycl::queue Queue) {
 // Test for array passing through a non-blocking pipe
 template <int TestNumber> int test_array_th_nb_pipe(cl::sycl::queue Queue) {
   int data[N] = {0};
-  using AnotherNbPipe = cl::sycl::INTEL::pipe<class another_nb_pipe, int>;
+  using AnotherNbPipe = cl::sycl::ext::intel::pipe<class another_nb_pipe, int>;
 
   Queue.submit([&](cl::sycl::handler &cgh) {
     cgh.single_task<class writer<TestNumber>>([=]() {
@@ -188,7 +188,7 @@ template <typename PipeName, int TestNumber>
 int test_simple_bl_pipe(cl::sycl::queue Queue) {
   int data[] = {0};
 
-  using Pipe = cl::sycl::INTEL::pipe<PipeName, int>;
+  using Pipe = cl::sycl::ext::intel::pipe<PipeName, int>;
 
   cl::sycl::buffer<int, 1> readBuf(data, 1);
   Queue.submit([&](cl::sycl::handler &cgh) {
@@ -256,7 +256,7 @@ template <int TestNumber> int test_multiple_bl_pipe(cl::sycl::queue Queue) {
 // Test for array passing through a blocking pipe
 template <int TestNumber> int test_array_th_bl_pipe(cl::sycl::queue Queue) {
   int data[N] = {0};
-  using AnotherBlPipe = cl::sycl::INTEL::pipe<class another_bl_pipe, int>;
+  using AnotherBlPipe = cl::sycl::ext::intel::pipe<class another_bl_pipe, int>;
 
   Queue.submit([&](cl::sycl::handler &cgh) {
     cgh.single_task<class writer<TestNumber>>([=]() {
@@ -290,7 +290,7 @@ int main() {
 
   if (!Queue.get_device()
            .get_info<cl::sycl::info::device::kernel_kernel_pipe_support>()) {
-    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;
   }

SYCL/Basic/fpga_tests/global_fpga_device_selector.cpp

@@ -4,12 +4,12 @@
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 #include <CL/sycl.hpp>
-#include <CL/sycl/INTEL/fpga_extensions.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
 
 // 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; }

SYCL/Basic/fpga_tests/io_pipe_def.h

@@ -1,12 +1,12 @@
-#include <CL/sycl/INTEL/fpga_extensions.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
 
 namespace intelfpga {
 template <unsigned ID> struct ethernet_pipe_id {
   static constexpr unsigned id = ID;
 };
 
 using ethernet_read_pipe =
-    sycl::INTEL::kernel_readable_io_pipe<ethernet_pipe_id<0>, int, 0>;
+    sycl::ext::intel::kernel_readable_io_pipe<ethernet_pipe_id<0>, int, 0>;
 using ethernet_write_pipe =
-    sycl::INTEL::kernel_writeable_io_pipe<ethernet_pipe_id<1>, int, 0>;
+    sycl::ext::intel::kernel_writeable_io_pipe<ethernet_pipe_id<1>, int, 0>;
 } // namespace intelfpga

SYCL/Basic/fpga_tests/pipes_info.cpp

@@ -23,7 +23,7 @@ int main() {
       Device.get_info<cl::sycl::info::device::kernel_kernel_pipe_support>();
 
   // 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<cl::sycl::info::platform::name>();
   bool SupposedToBeSupported =

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];
           });

SYCL/Basic/stream/blocking_pipes_and_stream.cpp

@@ -28,7 +28,7 @@
 // CHECK-NEXT: 9
 
 #include <CL/sycl.hpp>
-#include <CL/sycl/INTEL/fpga_extensions.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
 
 using namespace cl::sycl;
 

SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp

@@ -33,7 +33,8 @@ int main() {
   P.build_with_kernel_type<class K>();
   cl::sycl::kernel KE = P.get_kernel<class K>();
 
-  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<cl::sycl::access::mode::read>(CGH);
     CGH.parallel_for<class K>(
         KE, cl::sycl::range<1>(Size), [=](cl::sycl::id<1> Index) {
-          auto Fptr =
-              cl::sycl::ONEAPI::to_device_func_ptr<decltype(add)>(FptrStorage);
+          auto Fptr = cl::sycl::ext::oneapi::to_device_func_ptr<decltype(add)>(
+              FptrStorage);
           AccA[Index] = Fptr(AccA[Index], AccB[Index]);
         });
   });

SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp

@@ -35,12 +35,13 @@ int main() {
   P.build_with_kernel_type<class K>();
   cl::sycl::kernel KE = P.get_kernel<class K>();
 
-  cl::sycl::buffer<cl::sycl::ONEAPI::device_func_ptr_holder_t> DispatchTable(2);
+  cl::sycl::buffer<cl::sycl::ext::oneapi::device_func_ptr_holder_t>
+      DispatchTable(2);
   {
     auto DTAcc =
         DispatchTable.get_access<cl::sycl::access::mode::discard_write>();
-    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<cl::sycl::access::mode::read>(CGH);
       CGH.parallel_for<class K>(
           KE, cl::sycl::range<1>(Size), [=](cl::sycl::id<1> Index) {
-            auto FP = cl::sycl::ONEAPI::to_device_func_ptr<int(int, int)>(
+            auto FP = cl::sycl::ext::oneapi::to_device_func_ptr<int(int, int)>(
                 AccDT[Mode]);
 
             AccA[Index] = FP(AccA[Index], AccB[Index]);

SYCL/DeviceLib/built-ins/printf.cpp

@@ -41,7 +41,7 @@ int main() {
     Queue.submit([&](handler &CGH) {
       CGH.single_task<class integral>([=]() {
         // 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<int, 4>::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<class stream_string>(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();