[SYCL] Move general language extensions to the ONEAPI namespace

clang/lib/Sema/SemaSYCL.cpp

@@ -2816,7 +2816,7 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
   std::array<DeclContextDesc, 4> Scopes = {
       Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
       Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "experimental"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"},
       Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
   return matchQualifiedTypeName(Ty, Scopes);
 }

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -242,7 +242,7 @@ struct get_kernel_name_t<auto_name, Type> {
   using name = Type;
 };
 
-namespace experimental {
+namespace ONEAPI {
 template <typename T, typename ID = T>
 class spec_constant {
 public:
@@ -256,7 +256,7 @@ class spec_constant {
     return get();
   }
 };
-} // namespace experimental
+} // namespace ONEAPI
 
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>

clang/test/CodeGenSYCL/int_header_spec_const.cpp

@@ -20,18 +20,18 @@ class MyDoubleConst;
 
 int main() {
   // Create specialization constants.
-  cl::sycl::experimental::spec_constant<bool, MyBoolConst> i1(false);
-  cl::sycl::experimental::spec_constant<char, MyInt8Const> i8(0);
-  cl::sycl::experimental::spec_constant<unsigned char, MyUInt8Const> ui8(0);
-  cl::sycl::experimental::spec_constant<short, MyInt16Const> i16(0);
-  cl::sycl::experimental::spec_constant<unsigned short, MyUInt16Const> ui16(0);
-  cl::sycl::experimental::spec_constant<int, MyInt32Const> i32(0);
+  cl::sycl::ONEAPI::spec_constant<bool, MyBoolConst> i1(false);
+  cl::sycl::ONEAPI::spec_constant<char, MyInt8Const> i8(0);
+  cl::sycl::ONEAPI::spec_constant<unsigned char, MyUInt8Const> ui8(0);
+  cl::sycl::ONEAPI::spec_constant<short, MyInt16Const> i16(0);
+  cl::sycl::ONEAPI::spec_constant<unsigned short, MyUInt16Const> ui16(0);
+  cl::sycl::ONEAPI::spec_constant<int, MyInt32Const> i32(0);
   // Constant used twice, but there must be single entry in the int header,
   // otherwise compilation error would be issued.
-  cl::sycl::experimental::spec_constant<int, MyInt32Const> i32_1(0);
-  cl::sycl::experimental::spec_constant<unsigned int, MyUInt32Const> ui32(0);
-  cl::sycl::experimental::spec_constant<float, MyFloatConst> f32(0);
-  cl::sycl::experimental::spec_constant<double, MyDoubleConst> f64(0);
+  cl::sycl::ONEAPI::spec_constant<int, MyInt32Const> i32_1(0);
+  cl::sycl::ONEAPI::spec_constant<unsigned int, MyUInt32Const> ui32(0);
+  cl::sycl::ONEAPI::spec_constant<float, MyFloatConst> f32(0);
+  cl::sycl::ONEAPI::spec_constant<double, MyDoubleConst> f64(0);
 
   double val;
   double *ptr = &val; // to avoid "unused" warnings

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -208,7 +208,7 @@ class handler {
   }
 };
 
-namespace experimental {
+namespace ONEAPI {
 
 template <typename T, typename ID = T>
 class spec_constant {};

clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp

@@ -10,9 +10,9 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
 }
 
 int main() {
-  cl::sycl::experimental::spec_constant<char, class MyInt32Const> spec_const;
+  cl::sycl::ONEAPI::spec_constant<char, class MyInt32Const> spec_const;
   cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessor;
-  // CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::experimental::spec_constant<char, class MyInt32Const>'
+  // CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::ONEAPI::spec_constant<char, class MyInt32Const>'
   // CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>'
   kernel<class MyKernel>([spec_const, accessor]() {});
   return 0;

sycl/include/CL/sycl.hpp

@@ -22,17 +22,17 @@
 #include <CL/sycl/handler.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/image.hpp>
-#include <CL/sycl/intel/atomic.hpp>
-#include <CL/sycl/intel/builtins.hpp>
-#include <CL/sycl/intel/function_pointer.hpp>
-#include <CL/sycl/intel/group_algorithm.hpp>
-#include <CL/sycl/intel/reduction.hpp>
-#include <CL/sycl/intel/sub_group.hpp>
 #include <CL/sycl/item.hpp>
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/multi_ptr.hpp>
 #include <CL/sycl/nd_item.hpp>
 #include <CL/sycl/nd_range.hpp>
+#include <CL/sycl/ONEAPI/atomic.hpp>
+#include <CL/sycl/ONEAPI/builtins.hpp>
+#include <CL/sycl/ONEAPI/function_pointer.hpp>
+#include <CL/sycl/ONEAPI/group_algorithm.hpp>
+#include <CL/sycl/ONEAPI/reduction.hpp>
+#include <CL/sycl/ONEAPI/sub_group.hpp>
 #include <CL/sycl/pipes.hpp>
 #include <CL/sycl/platform.hpp>
 #include <CL/sycl/pointers.hpp>

sycl/include/CL/sycl/intel/atomic.hpp → sycl/include/CL/sycl/ONEAPI/atomic.hpp

@@ -1,4 +1,4 @@
-//==---------------- atomic.hpp - SYCL_INTEL_extended_atomics --------------==//
+//==--------------- atomic.hpp - SYCL_ONEAPI_extended_atomics --------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -8,7 +8,7 @@
 
 #pragma once
 
-#include <CL/sycl/intel/atomic_accessor.hpp>
-#include <CL/sycl/intel/atomic_enums.hpp>
-#include <CL/sycl/intel/atomic_fence.hpp>
-#include <CL/sycl/intel/atomic_ref.hpp>
+#include <CL/sycl/ONEAPI/atomic_accessor.hpp>
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
+#include <CL/sycl/ONEAPI/atomic_fence.hpp>
+#include <CL/sycl/ONEAPI/atomic_ref.hpp>

sycl/include/CL/sycl/intel/atomic_accessor.hpp → sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp

@@ -9,12 +9,12 @@
 #pragma once
 
 #include <CL/sycl/access/access.hpp>
-#include <CL/sycl/intel/atomic_enums.hpp>
-#include <CL/sycl/intel/atomic_ref.hpp>
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
+#include <CL/sycl/ONEAPI/atomic_ref.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
+namespace ONEAPI {
 
 #if __cplusplus > 201402L
 
@@ -123,6 +123,6 @@ atomic_accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
 
 #endif
 
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/intel/atomic_enums.hpp → sycl/include/CL/sycl/ONEAPI/atomic_enums.hpp

@@ -1,4 +1,4 @@
-//==---------------- atomic_enums.hpp - SYCL_INTEL_extended_atomics enums --==//
+//==--------------- atomic_enums.hpp - SYCL_ONEAPI_extended_atomics enums --==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -20,7 +20,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
+namespace ONEAPI {
 
 enum class memory_order : int {
   relaxed,
@@ -63,7 +63,7 @@ namespace detail {
 // Nested ternary conditions in else branch required for C++11
 #if __cplusplus >= 201402L
 static inline constexpr std::memory_order
-getStdMemoryOrder(::cl::sycl::intel::memory_order order) {
+getStdMemoryOrder(::cl::sycl::ONEAPI::memory_order order) {
   switch (order) {
   case memory_order::relaxed:
     return std::memory_order_relaxed;
@@ -81,7 +81,7 @@ getStdMemoryOrder(::cl::sycl::intel::memory_order order) {
 }
 #else
 static inline constexpr std::memory_order
-getStdMemoryOrder(::cl::sycl::intel::memory_order order) {
+getStdMemoryOrder(::cl::sycl::ONEAPI::memory_order order) {
   return (order == memory_order::relaxed)
              ? std::memory_order_relaxed
              : (order == memory_order::__consume_unsupported)
@@ -98,6 +98,6 @@ getStdMemoryOrder(::cl::sycl::intel::memory_order order) {
 } // namespace detail
 #endif // __SYCL_DEVICE_ONLY__
 
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/intel/atomic_fence.hpp → sycl/include/CL/sycl/ONEAPI/atomic_fence.hpp

@@ -1,4 +1,4 @@
-//==----- atomic_fence.hpp - SYCL_INTEL_extended_atomics atomic_fence ------==//
+//==---- atomic_fence.hpp - SYCL_ONEAPI_extended_atomics atomic_fence ------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -10,15 +10,15 @@
 
 #include <CL/__spirv/spirv_ops.hpp>
 #include <CL/sycl/detail/spirv.hpp>
-#include <CL/sycl/intel/atomic_enums.hpp>
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
 
 #ifndef __SYCL_DEVICE_ONLY__
 #include <atomic>
 #endif
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
+namespace ONEAPI {
 namespace detail {
 using namespace cl::sycl::detail;
 }
@@ -35,6 +35,6 @@ static inline void atomic_fence(memory_order order, memory_scope scope) {
 #endif
 }
 
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/intel/atomic_ref.hpp → sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp

@@ -1,4 +1,4 @@
-//==----- atomic_ref.hpp - SYCL_INTEL_extended_atomics atomic_ref ----------==//
+//==----- atomic_ref.hpp - SYCL_ONEAPI_extended_atomics atomic_ref ----------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -13,7 +13,7 @@
 #include <CL/sycl/atomic.hpp>
 #include <CL/sycl/detail/defines.hpp>
 #include <CL/sycl/detail/helpers.hpp>
-#include <CL/sycl/intel/atomic_enums.hpp>
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
 
 #ifndef __SYCL_DEVICE_ONLY__
 #include <atomic>
@@ -27,14 +27,14 @@ namespace sycl {
 template <typename pointerT, access::address_space AddressSpace>
 class multi_ptr;
 
-namespace intel {
+namespace ONEAPI {
 namespace detail {
 
-// Import from detail:: into intel::detail:: to improve readability later
+// Import from detail:: into ONEAPI::detail:: to improve readability later
 using namespace ::cl::sycl::detail;
 
-using memory_order = cl::sycl::intel::memory_order;
-using memory_scope = cl::sycl::intel::memory_scope;
+using memory_order = cl::sycl::ONEAPI::memory_order;
+using memory_scope = cl::sycl::ONEAPI::memory_scope;
 
 template <typename T>
 using IsValidAtomicType =
@@ -127,14 +127,14 @@ class atomic_ref_base {
       detail::IsValidAtomicType<T>::value,
       "Invalid atomic type.  Valid types are arithmetic and pointer types");
   static_assert(!std::is_same<T, bool>::value,
-                "intel::atomic_ref does not support bool type");
+                "ONEAPI::atomic_ref does not support bool type");
   static_assert(!(std::is_same<T, char>::value ||
                   std::is_same<T, signed char>::value ||
                   std::is_same<T, unsigned char>::value),
-                "intel::atomic_ref does not support char type");
+                "ONEAPI::atomic_ref does not support char type");
   static_assert(!(std::is_same<T, short>::value ||
                   std::is_same<T, unsigned short>::value),
-                "intel::atomic_ref does not support short type");
+                "ONEAPI::atomic_ref does not support short type");
   static_assert(detail::IsValidAtomicAddressSpace<AddressSpace>::value,
                 "Invalid atomic address_space.  Valid address spaces are: "
                 "global_space, local_space, global_device_space");
@@ -651,6 +651,6 @@ class atomic_ref : public detail::atomic_ref_impl<T, DefaultOrder, DefaultScope,
                                 AddressSpace>::atomic_ref_impl;
 };
 
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/intel/builtins.hpp → sycl/include/CL/sycl/ONEAPI/builtins.hpp

@@ -18,8 +18,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
-namespace experimental {
+namespace ONEAPI {
 
 // Provides functionality to print data from kernels in a C way:
 // - On non-host devices this function is directly mapped to printf from
@@ -67,8 +66,7 @@ int printf(const CONSTANT_AS char *__format, Args... args) {
 #endif
 }
 
-} // namespace experimental
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 

sycl/include/CL/sycl/intel/function_pointer.hpp → sycl/include/CL/sycl/ONEAPI/function_pointer.hpp

@@ -21,7 +21,7 @@ namespace detail {
 __SYCL_EXPORT cl_ulong getDeviceFunctionPointerImpl(device &D, program &P,
                                                     const char *FuncName);
 }
-namespace intel {
+namespace ONEAPI {
 
 // This is a preview extension implementation, intended to provide early
 // access to a feature for review and community feedback.
@@ -83,6 +83,6 @@ device_func_ptr_holder_t get_device_func_ptr(FuncType F, const char *FuncName,
 
   return sycl::detail::getDeviceFunctionPointerImpl(D, P, FuncName);
 }
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/intel/functional.hpp → sycl/include/CL/sycl/ONEAPI/functional.hpp

@@ -11,7 +11,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
+namespace ONEAPI {
 
 template <typename T = void> struct minimum {
   T operator()(const T &lhs, const T &rhs) const {
@@ -57,7 +57,7 @@ template <typename T = void> using bit_or = std::bit_or<T>;
 template <typename T = void> using bit_xor = std::bit_xor<T>;
 template <typename T = void> using bit_and = std::bit_and<T>;
 
-} // namespace intel
+} // namespace ONEAPI
 
 #ifdef __SYCL_DEVICE_ONLY__
 namespace detail {
@@ -93,15 +93,15 @@ struct GroupOpTag<T, detail::enable_if_t<detail::is_sgenfloat<T>::value>> {
     return Ret;                                                                \
   }
 
-__SYCL_CALC_OVERLOAD(GroupOpISigned, SMin, intel::minimum<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMin, intel::minimum<T>)
-__SYCL_CALC_OVERLOAD(GroupOpFP, FMin, intel::minimum<T>)
-__SYCL_CALC_OVERLOAD(GroupOpISigned, SMax, intel::maximum<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMax, intel::maximum<T>)
-__SYCL_CALC_OVERLOAD(GroupOpFP, FMax, intel::maximum<T>)
-__SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, intel::plus<T>)
-__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, intel::plus<T>)
-__SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, intel::plus<T>)
+__SYCL_CALC_OVERLOAD(GroupOpISigned, SMin, ONEAPI::minimum<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMin, ONEAPI::minimum<T>)
+__SYCL_CALC_OVERLOAD(GroupOpFP, FMin, ONEAPI::minimum<T>)
+__SYCL_CALC_OVERLOAD(GroupOpISigned, SMax, ONEAPI::maximum<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMax, ONEAPI::maximum<T>)
+__SYCL_CALC_OVERLOAD(GroupOpFP, FMax, ONEAPI::maximum<T>)
+__SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, ONEAPI::plus<T>)
+__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, ONEAPI::plus<T>)
+__SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, ONEAPI::plus<T>)
 
 #undef __SYCL_CALC_OVERLOAD