[SYCL] Inline cl namespace to simplify SYCL API usage

clang/lib/AST/QualTypeNames.cpp

@@ -197,7 +197,7 @@ static NestedNameSpecifier *createOuterNNS(const ASTContext &Ctx, const Decl *D,
       // Ignore inline namespace;
       NS = dyn_cast<NamespaceDecl>(NS->getDeclContext());
     }
-    if (NS->getDeclName()) {
+    if (NS && NS->getDeclName()) {
       return createNestedNameSpecifier(Ctx, NS, WithGlobalNsPrefix);
     }
     return nullptr;  // no starting '::', no anonymous

clang/lib/Sema/SemaSYCL.cpp

@@ -1638,6 +1638,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
 
+  O << "#include <CL/sycl/detail/defines.hpp>\n";
   O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
 
   O << "\n";
@@ -1651,7 +1652,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   }
   O << "\n";
 
-  O << "namespace cl {\n";
+  O << "__SYCL_INLINE namespace cl {\n";
   O << "namespace sycl {\n";
   O << "namespace detail {\n";
 

clang/test/CodeGenSYCL/kernel-name.cpp

@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm -x c++ %s -o - | FileCheck %s
+
+inline namespace cl {
+  namespace sycl {
+    class kernel {};
+  }
+}
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+// CHECK: define spir_kernel {{.*}}2cl4sycl6kernel
+int main() {
+  kernel_single_task<class kernel>([]() {});
+  return 0;
+}

sycl/include/CL/sycl.hpp

@@ -40,3 +40,4 @@
 #include <CL/sycl/types.hpp>
 #include <CL/sycl/usm.hpp>
 #include <CL/sycl/version.hpp>
+

sycl/include/CL/sycl/access/access.hpp

@@ -7,7 +7,9 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
-namespace cl {
+#include <CL/sycl/detail/defines.hpp>
+
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace access {
 

sycl/include/CL/sycl/accessor.hpp

@@ -143,7 +143,7 @@
 // accessor_common contains several helpers common for both accessor(1) and
 // accessor(3)
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 template <typename DataT, int Dimensions, access::mode AccessMode,

sycl/include/CL/sycl/aliases.hpp

@@ -13,7 +13,7 @@
 #include <cstddef>
 #include <cstdint>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 template <typename T, int N> class vec;
 namespace detail {
@@ -69,7 +69,7 @@ using half = cl::sycl::detail::half_impl::half;
   MAKE_VECTOR_ALIASES_FOR_OPENCL_TYPES(N)                                      \
   MAKE_VECTOR_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES(N)
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 using byte = std::uint8_t;
 using schar = signed char;

sycl/include/CL/sycl/atomic.hpp

@@ -23,7 +23,7 @@
   static_assert(!std::is_same<T, float>::value,                                \
                 "SYCL atomic function not available for float type")
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 enum class memory_order : int { relaxed };
@@ -66,7 +66,7 @@ template <> struct GetSpirvMemoryScope<access::address_space::local_space> {
 
 #ifndef __SYCL_DEVICE_ONLY__
 // host implementation of SYCL atomics
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 // Translate cl::sycl::memory_order or __spv::MemorySemanticsMask
@@ -161,7 +161,7 @@ extern T __spirv_AtomicMax(std::atomic<T> *Ptr, __spv::Scope S,
 
 #endif // !defined(__SYCL_DEVICE_ONLY__)
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 template <typename T, access::address_space addressSpace =

sycl/include/CL/sycl/buffer.hpp

@@ -14,7 +14,7 @@
 
 // TODO: 4.3.4 Properties
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 class handler;
 class queue;

sycl/include/CL/sycl/builtins.hpp

@@ -17,7 +17,7 @@
 // TODO Decide whether to mark functions with this attribute.
 #define __NOEXC /*noexcept*/
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 #ifdef __SYCL_DEVICE_ONLY__
 #define __sycl_std
@@ -27,7 +27,7 @@ namespace __sycl_std = __host_std;
 } // namespace sycl
 } // namespace cl
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
 // genfloat acos (genfloat x)

sycl/include/CL/sycl/context.hpp

@@ -14,7 +14,7 @@
 #include <type_traits>
 // 4.6.2 Context class
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 // Forward declarations
 class device;

sycl/include/CL/sycl/detail/accessor_impl.hpp

@@ -16,7 +16,7 @@
 
 #include <memory>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/aligned_allocator.hpp

@@ -19,7 +19,7 @@
 #include <memory>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 template <typename T> class aligned_allocator {

sycl/include/CL/sycl/detail/array.hpp

@@ -12,7 +12,7 @@
 #include <functional>
 #include <stdexcept>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 template <int dimensions> class id;
 template <int dimensions> class range;

sycl/include/CL/sycl/detail/boolean.hpp

@@ -14,7 +14,7 @@
 #include <initializer_list>
 #include <type_traits>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/buffer_impl.hpp

@@ -29,7 +29,7 @@
 #include <memory>
 #include <type_traits>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 // Forward declarations
 template <typename DataT, int Dimensions, access::mode AccessMode,

sycl/include/CL/sycl/detail/builtins.hpp

@@ -67,7 +67,7 @@
   }
 
 #ifndef __SYCL_DEVICE_ONLY__
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace __host_std {
 #endif // __SYCL_DEVICE_ONLY__
 /* ----------------- 4.13.3 Math functions. ---------------------------------*/

sycl/include/CL/sycl/detail/cg.hpp

@@ -24,7 +24,7 @@
 #include <type_traits>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/clusm.hpp

@@ -17,7 +17,7 @@
 #include <mutex>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 namespace usm {

sycl/include/CL/sycl/detail/common.hpp

@@ -8,6 +8,8 @@
 
 #pragma once
 
+#include <CL/sycl/detail/defines.hpp>
+
 // Suppress a compiler warning about undefined CL_TARGET_OPENCL_VERSION
 // Khronos ICD supports only latest OpenCL version
 #define CL_TARGET_OPENCL_VERSION 220
@@ -20,7 +22,7 @@
 #define STRINGIFY_LINE_HELP(s) #s
 #define STRINGIFY_LINE(s) STRINGIFY_LINE_HELP(s)
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
@@ -85,21 +87,7 @@ static inline std::string codeToString(cl_int code){
 #define CHECK_OCL_CODE_NO_EXC(X) REPORT_OCL_ERR_TO_STREAM(X)
 #endif
 
-#ifndef __has_attribute
-#define __has_attribute(x) 0
-#endif
-
-#if __has_attribute(always_inline)
-#define ALWAYS_INLINE __attribute__((always_inline))
-#else
-#define ALWAYS_INLINE
-#endif
-
-#ifndef SYCL_EXTERNAL
-#define SYCL_EXTERNAL
-#endif
-
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
@@ -175,7 +163,7 @@ struct NDLoopIterateImpl {
   }
 };
 
-// spcialization for DIM=0 to terminate recursion
+// Specialization for DIM=0 to terminate recursion
 template <int NDIMS, template <int> class LoopBoundTy, typename FuncTy,
           template <int> class LoopIndexTy>
 struct NDLoopIterateImpl<NDIMS, 0, LoopBoundTy, FuncTy, LoopIndexTy> {

sycl/include/CL/sycl/detail/common_info.hpp

@@ -9,7 +9,7 @@
 #pragma once
 #include <CL/sycl/stl.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/context_impl.hpp

@@ -22,7 +22,7 @@
 #include <map>
 #include <memory>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 // Forward declaration
 class device;

sycl/include/CL/sycl/detail/context_info.hpp

@@ -11,7 +11,7 @@
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/defines.hpp

@@ -0,0 +1,29 @@
+//==---------- defines.hpp ----- Preprocessor directives -------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
+#define __SYCL_INLINE inline
+#else
+#define __SYCL_INLINE
+#endif // __SYCL_DISABLE_NAMESPACE_INLINE__
+
+#ifndef __has_attribute
+#define __has_attribute(x) 0
+#endif
+
+#if __has_attribute(always_inline)
+#define ALWAYS_INLINE __attribute__((always_inline))
+#else
+#define ALWAYS_INLINE
+#endif
+
+#ifndef SYCL_EXTERNAL
+#define SYCL_EXTERNAL
+#endif

sycl/include/CL/sycl/detail/device_impl.hpp

@@ -13,7 +13,7 @@
 #include <CL/sycl/stl.hpp>
 #include <memory>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 // Forward declaration

sycl/include/CL/sycl/detail/device_info.hpp

@@ -8,12 +8,13 @@
 
 #pragma once
 #include <CL/sycl/detail/common_info.hpp>
+#include <CL/sycl/detail/defines.hpp>
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/detail/platform_impl.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 #include <CL/sycl/platform.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/event_impl.hpp

@@ -15,7 +15,7 @@
 
 #include <cassert>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 class context;
 namespace detail {

sycl/include/CL/sycl/detail/event_info.hpp

@@ -11,7 +11,7 @@
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 

sycl/include/CL/sycl/detail/force_device.hpp

@@ -8,9 +8,10 @@
 
 #pragma once
 
+#include <CL/sycl/detail/defines.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {