diff --git a/clang/lib/CodeGen/CGClass.cpp b/clang/lib/CodeGen/CGClass.cpp
index 3b5aa6645caa..4988274827a6 100644
--- a/clang/lib/CodeGen/CGClass.cpp
+++ b/clang/lib/CodeGen/CGClass.cpp
@@ -2491,7 +2491,8 @@ void CodeGenFunction::InitializeVTablePointer(const VPtr &Vptr) {
       llvm::FunctionType::get(CGM.Int32Ty, /*isVarArg=*/true)
           ->getPointerTo()
           ->getPointerTo();
-  VTableField = Builder.CreateBitCast(VTableField, VTablePtrTy->getPointerTo());
+  VTableField = Builder.CreateBitCast(VTableField, VTablePtrTy->getPointerTo(
+      VTableField.getType()->getAddressSpace()));
   VTableAddressPoint = Builder.CreateBitCast(VTableAddressPoint, VTablePtrTy);
 
   llvm::StoreInst *Store = Builder.CreateStore(VTableAddressPoint, VTableField);
diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index 1a3bbb768365..3be4da7b8670 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -749,10 +749,6 @@ llvm::StructType *CodeGenTypes::ConvertRecordDeclType(const RecordDecl *RD) {
     return Ty;
   }
 
-  assert((!Context.getLangOpts().SYCLIsDevice || !isa<CXXRecordDecl>(RD) ||
-          !dyn_cast<CXXRecordDecl>(RD)->isPolymorphic()) &&
-         "Types with virtual functions not allowed in SYCL");
-
   // Okay, this is a definition of a type.  Compile the implementation now.
   bool InsertResult = RecordsBeingLaidOut.insert(Key).second;
   (void)InsertResult;
diff --git a/clang/test/CodeGenSYCL/virtual-types.cpp b/clang/test/CodeGenSYCL/virtual-types.cpp
new file mode 100644
index 000000000000..a8d149b9dbe7
--- /dev/null
+++ b/clang/test/CodeGenSYCL/virtual-types.cpp
@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+struct Struct {
+  virtual void foo() {}
+  void bar() {}
+};
+
+int main() {
+  kernel_single_task<class kernel_function>([]() {
+                                            Struct S;
+                                            S.bar(); });
+  return 0;
+}
+
+
+// Struct layout big enough for vtable.
+// CHECK: %struct._ZTS6Struct.Struct = type { i32 (...)** }
+// VTable:
+// CHECK: @_ZTV6Struct = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* null, i8* bitcast ({ i8*, i8* }* @_ZTI6Struct to i8*), i8* bitcast (void (%struct._ZTS6Struct.Struct addrspace(4)*)* @_ZN6Struct3fooEv to i8*)] }, comdat, align 8
+// CHECK: @[[TYPEINFO:.+]] = external global i8*
+// TypeInfo Name:
+// CHECK: @_ZTS6Struct = linkonce_odr constant [8 x i8] c"6Struct\00", comdat, align 1
+// TypeInfo:
+// CHECK: @_ZTI6Struct = linkonce_odr constant { i8*, i8* } { i8* bitcast (i8** getelementptr inbounds (i8*, i8** @[[TYPEINFO]], i64 2) to i8*), i8* getelementptr inbounds ([8 x i8], [8 x i8]* @_ZTS6Struct, i32 0, i32 0) }, comdat, align 8
diff --git a/sycl/source/context.cpp b/sycl/source/context.cpp
index 6540c6fc9246..e54bd97c7f79 100644
--- a/sycl/source/context.cpp
+++ b/sycl/source/context.cpp
@@ -21,7 +21,7 @@
 
 // 4.6.2 Context class
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 context::context(const async_handler &AsyncHandler)
     : context(default_selector().select_device(), AsyncHandler) {}
diff --git a/sycl/source/detail/accessor_impl.cpp b/sycl/source/detail/accessor_impl.cpp
index e1d9062fcc00..9df50b1b9c1d 100644
--- a/sycl/source/detail/accessor_impl.cpp
+++ b/sycl/source/detail/accessor_impl.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/detail/accessor_impl.hpp>
 #include <CL/sycl/detail/scheduler/scheduler.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/builtins_common.cpp b/sycl/source/detail/builtins_common.cpp
index ebd2f00d93a6..b60af85cd4b5 100644
--- a/sycl/source/detail/builtins_common.cpp
+++ b/sycl/source/detail/builtins_common.cpp
@@ -22,7 +22,7 @@
 namespace s = cl::sycl;
 namespace d = s::detail;
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace __host_std {
 namespace {
 
diff --git a/sycl/source/detail/builtins_geometric.cpp b/sycl/source/detail/builtins_geometric.cpp
index efe0f3d9aa6c..fe6466c5609e 100644
--- a/sycl/source/detail/builtins_geometric.cpp
+++ b/sycl/source/detail/builtins_geometric.cpp
@@ -16,7 +16,7 @@
 namespace s = cl::sycl;
 namespace d = s::detail;
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace __host_std {
 
 s::cl_float Dot(s::cl_float2, s::cl_float2);
diff --git a/sycl/source/detail/builtins_integer.cpp b/sycl/source/detail/builtins_integer.cpp
index 49f83becfb6d..73ef8e73bce9 100644
--- a/sycl/source/detail/builtins_integer.cpp
+++ b/sycl/source/detail/builtins_integer.cpp
@@ -17,7 +17,7 @@
 namespace s = cl::sycl;
 namespace d = s::detail;
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace __host_std {
 namespace {
 
diff --git a/sycl/source/detail/builtins_math.cpp b/sycl/source/detail/builtins_math.cpp
index 2fa0014a7840..b7cfe5b982db 100644
--- a/sycl/source/detail/builtins_math.cpp
+++ b/sycl/source/detail/builtins_math.cpp
@@ -22,7 +22,7 @@
 namespace s = cl::sycl;
 namespace d = s::detail;
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace __host_std {
 
 namespace {
diff --git a/sycl/source/detail/builtins_relational.cpp b/sycl/source/detail/builtins_relational.cpp
index 3660e8291113..bbcd9fe0c1c9 100644
--- a/sycl/source/detail/builtins_relational.cpp
+++ b/sycl/source/detail/builtins_relational.cpp
@@ -16,7 +16,7 @@
 namespace s = cl::sycl;
 namespace d = s::detail;
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace __host_std {
 namespace {
 
diff --git a/sycl/source/detail/common.cpp b/sycl/source/detail/common.cpp
index 6758bcd4925a..dc6e19b83aac 100644
--- a/sycl/source/detail/common.cpp
+++ b/sycl/source/detail/common.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/common_info.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/config.cpp b/sycl/source/detail/config.cpp
index 0db83f5540a7..dc8db099069d 100644
--- a/sycl/source/detail/config.cpp
+++ b/sycl/source/detail/config.cpp
@@ -17,7 +17,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 {
 
diff --git a/sycl/source/detail/config.hpp b/sycl/source/detail/config.hpp
index 4319e00c9a98..5c006a8874d4 100644
--- a/sycl/source/detail/config.hpp
+++ b/sycl/source/detail/config.hpp
@@ -10,7 +10,7 @@
 
 #include <cstdlib>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp
index 7a6603001f70..2f468cf64718 100644
--- a/sycl/source/detail/context_impl.cpp
+++ b/sycl/source/detail/context_impl.cpp
@@ -17,7 +17,7 @@
 #include <CL/sycl/platform.hpp>
 #include <CL/sycl/stl.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp
index 06fedcf58037..c49ac7b6d2e4 100644
--- a/sycl/source/detail/device_impl.cpp
+++ b/sycl/source/detail/device_impl.cpp
@@ -11,7 +11,7 @@
 
 #include <algorithm>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/device_info.cpp b/sycl/source/detail/device_info.cpp
index c50d58ee51cf..52aaedcf8319 100644
--- a/sycl/source/detail/device_info.cpp
+++ b/sycl/source/detail/device_info.cpp
@@ -19,7 +19,7 @@
   (__GNUC__ * 10000 + __GNUC_MINOR__ * 100 + __GNUC_PATCHLEVEL__)
 #endif
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/error_handling/enqueue_kernel.cpp b/sycl/source/detail/error_handling/enqueue_kernel.cpp
index 1ab7a11578f1..2d72f516f3a9 100644
--- a/sycl/source/detail/error_handling/enqueue_kernel.cpp
+++ b/sycl/source/detail/error_handling/enqueue_kernel.cpp
@@ -14,7 +14,7 @@
 
 #include <CL/sycl/detail/pi.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/error_handling/error_handling.hpp b/sycl/source/detail/error_handling/error_handling.hpp
index 5bdf0cb90ffa..be6edba17f2d 100644
--- a/sycl/source/detail/error_handling/error_handling.hpp
+++ b/sycl/source/detail/error_handling/error_handling.hpp
@@ -11,7 +11,7 @@
 #include <CL/sycl/detail/pi.h>
 #include <CL/sycl/detail/cg.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp
index 7b28e93f04ee..e9838614e334 100644
--- a/sycl/source/detail/event_impl.cpp
+++ b/sycl/source/detail/event_impl.cpp
@@ -13,7 +13,7 @@
 
 #include <chrono>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/force_device.cpp b/sycl/source/detail/force_device.cpp
index d7fa4778964f..95335687ad3b 100644
--- a/sycl/source/detail/force_device.cpp
+++ b/sycl/source/detail/force_device.cpp
@@ -11,7 +11,7 @@
 #include <CL/sycl/stl.hpp>
 #include <cstdlib>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/helpers.cpp b/sycl/source/detail/helpers.cpp
index f6b9cfc83005..e93d425aa838 100644
--- a/sycl/source/detail/helpers.cpp
+++ b/sycl/source/detail/helpers.cpp
@@ -14,7 +14,7 @@
 
 #include <memory>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 using ContextImplPtr = std::shared_ptr<cl::sycl::detail::context_impl>;
 namespace detail {
diff --git a/sycl/source/detail/image_accessor_util.cpp b/sycl/source/detail/image_accessor_util.cpp
index 11df1306d750..de6524dea30a 100644
--- a/sycl/source/detail/image_accessor_util.cpp
+++ b/sycl/source/detail/image_accessor_util.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/builtins.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/image_impl.cpp b/sycl/source/detail/image_impl.cpp
index c32553ab6b63..2fb088e1edd5 100644
--- a/sycl/source/detail/image_impl.cpp
+++ b/sycl/source/detail/image_impl.cpp
@@ -8,7 +8,7 @@
 
 #include <CL/sycl/image.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 63e978c6abd6..4bb9e56bcb8a 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -15,7 +15,7 @@
 
 #include <memory>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/kernel_info.cpp b/sycl/source/detail/kernel_info.cpp
index 32dde9e87b7b..aa38ad505683 100644
--- a/sycl/source/detail/kernel_info.cpp
+++ b/sycl/source/detail/kernel_info.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/detail/kernel_info.hpp>
 #include <CL/sycl/device.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 template <>
diff --git a/sycl/source/detail/kernel_program_cache.cpp b/sycl/source/detail/kernel_program_cache.cpp
index 8a780841f707..03d780653738 100644
--- a/sycl/source/detail/kernel_program_cache.cpp
+++ b/sycl/source/detail/kernel_program_cache.cpp
@@ -8,7 +8,7 @@
 
 #include <CL/sycl/detail/kernel_program_cache.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 KernelProgramCache::~KernelProgramCache() {
diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp
index 1e97e1f162e7..254b63592abe 100644
--- a/sycl/source/detail/memory_manager.cpp
+++ b/sycl/source/detail/memory_manager.cpp
@@ -17,7 +17,7 @@
 #include <cstring>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/os_util.cpp b/sycl/source/detail/os_util.cpp
index e6d29fda9c85..9fcc7ccd0394 100644
--- a/sycl/source/detail/os_util.cpp
+++ b/sycl/source/detail/os_util.cpp
@@ -44,7 +44,7 @@
 
 #endif // SYCL_RT_OS
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp
index abb177c81347..10df59f5d496 100644
--- a/sycl/source/detail/pi.cpp
+++ b/sycl/source/detail/pi.cpp
@@ -15,7 +15,7 @@
 #include <stddef.h>
 #include <string>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 namespace pi {
diff --git a/sycl/source/detail/platform_impl.cpp b/sycl/source/detail/platform_impl.cpp
index d4f3b4da0806..cc55e98df85f 100644
--- a/sycl/source/detail/platform_impl.cpp
+++ b/sycl/source/detail/platform_impl.cpp
@@ -15,7 +15,7 @@
 #include <cstring>
 #include <regex>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/platform_info.cpp b/sycl/source/detail/platform_info.cpp
index 27669f36ae35..c06be8d42467 100644
--- a/sycl/source/detail/platform_info.cpp
+++ b/sycl/source/detail/platform_info.cpp
@@ -8,7 +8,7 @@
 
 #include <CL/sycl/detail/platform_info.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/platform_util.cpp b/sycl/source/detail/platform_util.cpp
index 480579d6c4f3..6027eec91c39 100644
--- a/sycl/source/detail/platform_util.cpp
+++ b/sycl/source/detail/platform_util.cpp
@@ -16,7 +16,7 @@
 #include <intrin.h>
 #endif
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/posix_pi.cpp b/sycl/source/detail/posix_pi.cpp
index dd1a865f53c3..77080dd40a6f 100644
--- a/sycl/source/detail/posix_pi.cpp
+++ b/sycl/source/detail/posix_pi.cpp
@@ -1,7 +1,17 @@
+//==---------------- posix_pi.cpp ------------------------------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl/detail/defines.hpp>
+
 #include <dlfcn.h>
 #include <string>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 namespace pi {
diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp
index b4fab517b92f..24ed7d4787cb 100644
--- a/sycl/source/detail/program_impl.cpp
+++ b/sycl/source/detail/program_impl.cpp
@@ -15,7 +15,7 @@
 #include <fstream>
 #include <memory>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index d2ef945232ba..4df570273d4a 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -27,7 +27,7 @@
 #include <mutex>
 #include <sstream>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index 7c343c96ec38..238f5e835b4c 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/sycl/source/detail/queue_impl.cpp
@@ -16,7 +16,7 @@
 
 #include <cstring>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 template <> cl_uint queue_impl::get_info<info::queue::reference_count>() const {
diff --git a/sycl/source/detail/sampler_impl.cpp b/sycl/source/detail/sampler_impl.cpp
index db8a9e6d029f..808208f6eb54 100644
--- a/sycl/source/detail/sampler_impl.cpp
+++ b/sycl/source/detail/sampler_impl.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/detail/context_impl.hpp>
 #include <CL/sycl/detail/sampler_impl.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 1212b84bad98..fae6ed19e88f 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -32,7 +32,7 @@
 #include <memory>
 #endif
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp
index 41ca557d3c49..492cf34a1c5f 100644
--- a/sycl/source/detail/scheduler/graph_builder.cpp
+++ b/sycl/source/detail/scheduler/graph_builder.cpp
@@ -22,7 +22,7 @@
 #include <set>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/scheduler/graph_processor.cpp b/sycl/source/detail/scheduler/graph_processor.cpp
index 78f4ec26adf0..b3cb7f80438c 100644
--- a/sycl/source/detail/scheduler/graph_processor.cpp
+++ b/sycl/source/detail/scheduler/graph_processor.cpp
@@ -13,7 +13,7 @@
 #include <memory>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp
index 28a77adc77c7..8e3fe9a4cbcf 100644
--- a/sycl/source/detail/scheduler/scheduler.cpp
+++ b/sycl/source/detail/scheduler/scheduler.cpp
@@ -16,7 +16,7 @@
 #include <set>
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/stream_impl.cpp b/sycl/source/detail/stream_impl.cpp
index e66cd11ea6cb..dcb5a75748cb 100644
--- a/sycl/source/detail/stream_impl.cpp
+++ b/sycl/source/detail/stream_impl.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/detail/stream_impl.hpp>
 #include <cstdio>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/usm/clusm.cpp b/sycl/source/detail/usm/clusm.cpp
index cac788320fe9..87de620a3f05 100644
--- a/sycl/source/detail/usm/clusm.cpp
+++ b/sycl/source/detail/usm/clusm.cpp
@@ -19,7 +19,7 @@
 
 cl::sycl::detail::usm::CLUSM *gCLUSM = nullptr;
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 namespace usm {
diff --git a/sycl/source/detail/usm/usm_dispatch.cpp b/sycl/source/detail/usm/usm_dispatch.cpp
index e995f10629ed..599aeeffa0d8 100644
--- a/sycl/source/detail/usm/usm_dispatch.cpp
+++ b/sycl/source/detail/usm/usm_dispatch.cpp
@@ -9,7 +9,7 @@
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/detail/usm_dispatch.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 namespace usm {
diff --git a/sycl/source/detail/usm/usm_impl.cpp b/sycl/source/detail/usm/usm_impl.cpp
index 23bb2b97c316..39b1f212ae5f 100644
--- a/sycl/source/detail/usm/usm_impl.cpp
+++ b/sycl/source/detail/usm/usm_impl.cpp
@@ -15,7 +15,7 @@
 
 #include <cstdlib>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 using alloc = cl::sycl::usm::alloc;
diff --git a/sycl/source/detail/util.cpp b/sycl/source/detail/util.cpp
index 3a712009a868..7b98f1a7dab8 100644
--- a/sycl/source/detail/util.cpp
+++ b/sycl/source/detail/util.cpp
@@ -8,7 +8,7 @@
 
 #include <CL/sycl/detail/util.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/detail/windows_pi.cpp b/sycl/source/detail/windows_pi.cpp
index 69817ab23c97..dee857644da8 100644
--- a/sycl/source/detail/windows_pi.cpp
+++ b/sycl/source/detail/windows_pi.cpp
@@ -1,8 +1,18 @@
+//==---------------- windows_pi.cpp ----------------------------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl/detail/defines.hpp>
+
 #include <windows.h>
 #include <winreg.h>
 #include <string>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 namespace pi {
diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp
index 907dc0fc7fdb..52e05d72b059 100644
--- a/sycl/source/device.cpp
+++ b/sycl/source/device.cpp
@@ -12,7 +12,7 @@
 #include <CL/sycl/device_selector.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 void force_type(info::device_type &t, const info::device_type &ft) {
diff --git a/sycl/source/device_selector.cpp b/sycl/source/device_selector.cpp
index 21f4ba8d4c36..6f3e1c636bee 100644
--- a/sycl/source/device_selector.cpp
+++ b/sycl/source/device_selector.cpp
@@ -12,7 +12,7 @@
 #include <CL/sycl/stl.hpp>
 // 4.6.1 Device selection class
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 device device_selector::select_device() const {
   vector_class<device> devices = device::get_devices();
diff --git a/sycl/source/event.cpp b/sycl/source/event.cpp
index 0b24dbf62dd2..569d61340f1f 100644
--- a/sycl/source/event.cpp
+++ b/sycl/source/event.cpp
@@ -19,7 +19,7 @@
 #include <memory>
 #include <unordered_set>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 event::event() : impl(std::make_shared<detail::event_impl>()) {}
diff --git a/sycl/source/exception.cpp b/sycl/source/exception.cpp
index 7948fedee711..756571ef65fd 100644
--- a/sycl/source/exception.cpp
+++ b/sycl/source/exception.cpp
@@ -10,7 +10,7 @@
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/exception.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 const char *exception::what() const noexcept { return MMsg.c_str(); }
diff --git a/sycl/source/exception_list.cpp b/sycl/source/exception_list.cpp
index 9da0593fb019..4db32e0c8955 100644
--- a/sycl/source/exception_list.cpp
+++ b/sycl/source/exception_list.cpp
@@ -11,7 +11,7 @@
 
 #include <utility>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 exception_list::size_type exception_list::size() const { return MList.size(); }
diff --git a/sycl/source/half_type.cpp b/sycl/source/half_type.cpp
index 6f14ec47b072..9a5e8565a809 100644
--- a/sycl/source/half_type.cpp
+++ b/sycl/source/half_type.cpp
@@ -11,7 +11,7 @@
 #include <CL/sycl/detail/platform_util.hpp>
 #include <cstring>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace detail {
 
diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
index 87c737010082..24f28b7fca4c 100644
--- a/sycl/source/kernel.cpp
+++ b/sycl/source/kernel.cpp
@@ -10,7 +10,7 @@
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/program.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 kernel::kernel(cl_kernel ClKernel, const context &SyclContext)
diff --git a/sycl/source/ordered_queue.cpp b/sycl/source/ordered_queue.cpp
index 9cbcf5627810..85dd1763dadc 100644
--- a/sycl/source/ordered_queue.cpp
+++ b/sycl/source/ordered_queue.cpp
@@ -11,7 +11,7 @@
 
 #include <algorithm>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 ordered_queue::ordered_queue(const context &syclContext,
                              const device_selector &deviceSelector,
diff --git a/sycl/source/platform.cpp b/sycl/source/platform.cpp
index 5e56a7ed506e..020637c5551a 100644
--- a/sycl/source/platform.cpp
+++ b/sycl/source/platform.cpp
@@ -13,7 +13,7 @@
 #include <CL/sycl/info/info_desc.hpp>
 #include <CL/sycl/platform.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 platform::platform() : impl(std::make_shared<detail::platform_impl>()) {}
diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp
index cb6582eb95b2..e45b91740d70 100644
--- a/sycl/source/queue.cpp
+++ b/sycl/source/queue.cpp
@@ -11,7 +11,7 @@
 
 #include <algorithm>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 queue::queue(const context &syclContext, const device_selector &deviceSelector,
              const async_handler &asyncHandler, const property_list &propList) {
diff --git a/sycl/source/sampler.cpp b/sycl/source/sampler.cpp
index 13163dce2dd0..d06222a13cf4 100644
--- a/sycl/source/sampler.cpp
+++ b/sycl/source/sampler.cpp
@@ -8,7 +8,7 @@
 
 #include <CL/sycl/sampler.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 sampler::sampler(coordinate_normalization_mode normalizationMode,
                  addressing_mode addressingMode, filtering_mode filteringMode)
diff --git a/sycl/source/stream.cpp b/sycl/source/stream.cpp
index 4a67826b58bb..84b5ee60a952 100644
--- a/sycl/source/stream.cpp
+++ b/sycl/source/stream.cpp
@@ -8,7 +8,7 @@
 
 #include <CL/sycl/stream.hpp>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 stream::stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH)
diff --git a/sycl/test/basic_tests/sycl-namespace.cpp b/sycl/test/basic_tests/sycl-namespace.cpp
new file mode 100644
index 000000000000..ead6b0dc8248
--- /dev/null
+++ b/sycl/test/basic_tests/sycl-namespace.cpp
@@ -0,0 +1,20 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+int main() {
+  ::sycl::queue q;
+  int r = 0;
+  {
+    sycl::buffer<int, 1> b(&r, 1);
+    q.submit([&](sycl::handler &h) {
+      auto a = b.get_access<sycl::access::mode::write>(h);
+      h.single_task<class T>([=]() { a[0] = 42; });
+    });
+  }
+  return r - 42;
+}