[SYCL] Move general language extensions to the ONEAPI namespace (#2231)

This concentrates general language extensions into the ONEAPI namespace.

It includes:
Atomics
Printf
Group Algorithms
Reductions
Sub-groups
Specialization Constants
Function Pointers

Signed-off-by: James Brodman [email protected]

Author: jbrodman

clang/lib/Sema/SemaSYCL.cpp

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

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -253,6 +253,7 @@ struct get_kernel_name_t<auto_name, Type> {
   using name = Type;
 };
 
+namespace ONEAPI {
 namespace experimental {
 template <typename T, typename ID = T>
 class spec_constant {
@@ -268,6 +269,7 @@ class spec_constant {
   }
 };
 } // 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::experimental::spec_constant<bool, MyBoolConst> i1(false);
+  cl::sycl::ONEAPI::experimental::spec_constant<char, MyInt8Const> i8(0);
+  cl::sycl::ONEAPI::experimental::spec_constant<unsigned char, MyUInt8Const> ui8(0);
+  cl::sycl::ONEAPI::experimental::spec_constant<short, MyInt16Const> i16(0);
+  cl::sycl::ONEAPI::experimental::spec_constant<unsigned short, MyUInt16Const> ui16(0);
+  cl::sycl::ONEAPI::experimental::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::experimental::spec_constant<int, MyInt32Const> i32_1(0);
+  cl::sycl::ONEAPI::experimental::spec_constant<unsigned int, MyUInt32Const> ui32(0);
+  cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst> f32(0);
+  cl::sycl::ONEAPI::experimental::spec_constant<double, MyDoubleConst> f64(0);
 
   double val;
   double *ptr = &val; // to avoid "unused" warnings

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -218,12 +218,12 @@ class handler {
   }
 };
 
+namespace ONEAPI {
 namespace experimental {
-
 template <typename T, typename ID = T>
 class spec_constant {};
 } // namespace experimental
-
+} // namespace ONEAPI
 } // namespace sycl
 } // namespace cl
 

clang/test/SemaSYCL/spec-const-kernel-arg.cpp

@@ -6,12 +6,12 @@
 #include <sycl.hpp>
 
 struct SpecConstantsWrapper {
-  cl::sycl::experimental::spec_constant<int, class sc_name1> SC1;
-  cl::sycl::experimental::spec_constant<int, class sc_name2> SC2;
+  cl::sycl::ONEAPI::experimental::spec_constant<int, class sc_name1> SC1;
+  cl::sycl::ONEAPI::experimental::spec_constant<int, class sc_name2> SC2;
 };
 
 int main() {
-  cl::sycl::experimental::spec_constant<char, class MyInt32Const> SC;
+  cl::sycl::ONEAPI::experimental::spec_constant<char, class MyInt32Const> SC;
   SpecConstantsWrapper W;
   cl::sycl::kernel_single_task<class kernel_sc>(
       [=]() {
@@ -23,7 +23,7 @@ int main() {
 // CHECK: FunctionDecl {{.*}}kernel_sc{{.*}} 'void ()'
 // CHECK: VarDecl {{.*}}'(lambda at {{.*}}'
 // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::experimental::spec_constant<char, class MyInt32Const>':'cl::sycl::experimental::spec_constant<char, MyInt32Const>'
+// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::ONEAPI::experimental::spec_constant<char, class MyInt32Const>':'cl::sycl::ONEAPI::experimental::spec_constant<char, MyInt32Const>'
 // CHECK-NEXT: InitListExpr {{.*}} 'SpecConstantsWrapper'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::experimental::spec_constant<int, class sc_name1>':'cl::sycl::experimental::spec_constant<int, sc_name1>'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::experimental::spec_constant<int, class sc_name2>':'cl::sycl::experimental::spec_constant<int, sc_name2>'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::ONEAPI::experimental::spec_constant<int, class sc_name1>':'cl::sycl::ONEAPI::experimental::spec_constant<int, sc_name1>'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::ONEAPI::experimental::spec_constant<int, class sc_name2>':'cl::sycl::ONEAPI::experimental::spec_constant<int, sc_name2>'

clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp

@@ -10,9 +10,9 @@ __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
 }
 
 int main() {
-  cl::sycl::experimental::spec_constant<char, class MyInt32Const> spec_const;
+  cl::sycl::ONEAPI::experimental::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::experimental::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;

llvm/test/tools/sycl-post-link/sc_sym_two_refs.ll

@@ -7,7 +7,7 @@
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir64-unknown-unknown-sycldevice"
 
-%"sycl::experimental::spec_constant" = type { i8 }
+%"spec_constant" = type { i8 }
 
 @SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
 ; CHECK-NOT: @SCSymID
@@ -21,7 +21,7 @@ define weak_odr dso_local spir_kernel void @Kernel() {
 }
 
 ; Function Attrs: norecurse
-define dso_local spir_func float @foo_float(%"sycl::experimental::spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
+define dso_local spir_func float @foo_float(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
   %2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
   ret float %2
 }

sycl/CMakeLists.txt

@@ -15,7 +15,7 @@ include(AddSYCLExecutable)
 set(SYCL_MAJOR_VERSION 3)
 set(SYCL_MINOR_VERSION 0)
 set(SYCL_PATCH_VERSION 0)
-set(SYCL_DEV_ABI_VERSION 0)
+set(SYCL_DEV_ABI_VERSION 1)
 if (SYCL_ADD_DEV_VERSION_POSTFIX)
   set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}")
 endif()

sycl/include/CL/sycl.hpp

@@ -8,6 +8,12 @@
 
 #pragma once
 
+#include <CL/sycl/ONEAPI/atomic.hpp>
+#include <CL/sycl/ONEAPI/experimental/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/accessor.hpp>
 #include <CL/sycl/aspects.hpp>
 #include <CL/sycl/atomic.hpp>
@@ -23,12 +29,6 @@
 #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>

sycl/include/CL/sycl/intel/atomic.hpp renamed to 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 renamed to sycl/include/CL/sycl/ONEAPI/atomic_accessor.hpp

@@ -1,4 +1,4 @@
-//==--- atomic_accessor.hpp - SYCL_INTEL_extended_atomics atomic_accessor --==//
+//==-- atomic_accessor.hpp - SYCL_ONEAPI_extended_atomics atomic_accessor --==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -8,13 +8,14 @@
 
 #pragma once
 
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
+#include <CL/sycl/ONEAPI/atomic_ref.hpp>
 #include <CL/sycl/access/access.hpp>
-#include <CL/sycl/intel/atomic_enums.hpp>
-#include <CL/sycl/intel/atomic_ref.hpp>
+#include <CL/sycl/accessor.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
+namespace ONEAPI {
 
 #if __cplusplus > 201402L
 
@@ -123,6 +124,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 renamed to 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 renamed to 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.
@@ -9,16 +9,16 @@
 #pragma once
 
 #include <CL/__spirv/spirv_ops.hpp>
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
 #include <CL/sycl/detail/spirv.hpp>
-#include <CL/sycl/intel/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 renamed to 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.
@@ -9,11 +9,11 @@
 #pragma once
 
 #include <CL/__spirv/spirv_ops.hpp>
+#include <CL/sycl/ONEAPI/atomic_enums.hpp>
 #include <CL/sycl/access/access.hpp>
 #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/detail/spirv.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 renamed to sycl/include/CL/sycl/ONEAPI/experimental/builtins.hpp

@@ -18,7 +18,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace intel {
+namespace ONEAPI {
 namespace experimental {
 
 // Provides functionality to print data from kernels in a C way:
@@ -68,7 +68,7 @@ int printf(const CONSTANT_AS char *__format, Args... args) {
 }
 
 } // namespace experimental
-} // namespace intel
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 

sycl/include/CL/sycl/experimental/spec_constant.hpp renamed to sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp

@@ -1,4 +1,4 @@
-//==----- spec_constant.hpp - SYCL public experimental API header file -----==//
+//==----------- spec_constant.hpp - SYCL public ONEAPI API header file -----==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -22,6 +22,9 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+class program;
+
+namespace ONEAPI {
 namespace experimental {
 
 class spec_const_error : public compile_program_error {
@@ -60,5 +63,6 @@ template <typename T, typename ID = T> class spec_constant {
 };
 
 } // namespace experimental
+} // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)