[SYCL] Inline cl namespace to simplify SYCL API usage (#974)

The inlining is enabled by default and make using SYCL API a little bit
easier.

In order to disable cl namespace inlining user should define
__SYCL_DISABLE_NAMESPACE_INLINE__ macro.

Signed-off-by: Alexey Bader <[email protected]>

Author: bader

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)