From 27e7b6ab5f07e66fc019d2abff2e53a7a88ca2a8 Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Wed, 10 Jun 2020 13:17:11 +0300
Subject: [PATCH 1/7] [SYCL] Add device_ptr and host_ptr

Justification:
Currently a device backend can't trace from where a pointer allocated by
USM comes: it can be either allocated on host or on device (it's just a
pointer in OpenCL global address space). On FPGAs at least we can generate
more efficient hardware code if the user tells us where the pointer can
point. With this change users can create multi_ptr with specialized
address space global_host or global_device that will proved to the compiler
additional information to process load-store optimizations. Accessor
pointers shall be also moved to global_device address spaces - otherwise
backend would assume, that a pointer in global address space can access
both host and device memory.

Previously there were added global_device in global_host address spaces
for OpenCL/SYCL in clang. With this patch device_space and host_space
were added in the SYCL headers and are mapped into the new address spaces
and aliases to multi_ptr instantiated with the space: device_ptr and
host_ptr.

Added explicit conversion operator that allows to convert
device_ptr/host_ptr to global_ptr. Conversion in the opposite direction is
disallowed.

Also accessor to global_buffer pointer was moved to global_device address
spaces.

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 clang/test/CodeGenSYCL/Inputs/sycl.hpp        |  9 ++--
 .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 12 ++---
 clang/test/CodeGenSYCL/kernel-metadata.cpp    |  2 +-
 clang/test/SemaSYCL/Inputs/sycl.hpp           |  6 ++-
 clang/test/SemaSYCL/accessors-targets.cpp     |  2 +-
 clang/test/SemaSYCL/basic-kernel-wrapper.cpp  |  6 +--
 clang/test/SemaSYCL/fake-accessors.cpp        |  6 +--
 clang/test/SemaSYCL/wrapped-accessor.cpp      |  6 +--
 sycl/include/CL/sycl/access/access.hpp        | 33 +++++++++++-
 sycl/include/CL/sycl/atomic.hpp               | 14 +++--
 .../CL/sycl/detail/generic_type_lists.hpp     | 29 +++++------
 sycl/include/CL/sycl/handler.hpp              |  4 +-
 sycl/include/CL/sycl/multi_ptr.hpp            | 51 +++++++++++++------
 sycl/include/CL/sycl/pointers.hpp             |  6 +++
 .../check_device_code/kernel_arguments_as.cpp |  2 +-
 sycl/test/check_device_code/usm_pointers.cpp  | 41 +++++++++++++++
 sycl/test/multi_ptr/multi_ptr.cpp             | 17 +++++++
 17 files changed, 185 insertions(+), 61 deletions(-)
 create mode 100644 sycl/test/check_device_code/usm_pointers.cpp

diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 3184c58edcbfc..f1ad520ee248a 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -57,7 +57,9 @@ enum class address_space : int {
   private_space = 0,
   global_space,
   constant_space,
-  local_space
+  local_space,
+  device_space,
+  host_space
 };
 } // namespace access
 
@@ -139,8 +141,9 @@ class accessor {
   _ImplT<dimensions> impl;
 
 private:
-  void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
-              range<dimensions> MemRange, id<dimensions> Offset) {}
+  void __init(__attribute__((opencl_global_device)) dataT *Ptr,
+              range<dimensions> AccessRange, range<dimensions> MemRange,
+              id<dimensions> Offset) {}
 };
 
 template <int dimensions, access::mode accessmode, access::target accesstarget>
diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
index 31795fc73b776..ceab1dee68606 100644
--- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
@@ -20,12 +20,12 @@ int main() {
 }
 
 // CHECK: define spir_kernel void @{{.*}}kernel_function
-// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 addrspace(11)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
 // Check alloca for pointer argument
-// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
+// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(11)*
 // Check lambda object alloca
 // CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon"
 // Check allocas for ranges
@@ -34,7 +34,7 @@ int main() {
 // CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id"
 //
 // Check store of kernel pointer argument to alloca
-// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8
+// CHECK: store i32 addrspace(11)* [[MEM_ARG]], i32 addrspace(11)** [[MEM_ARG]].addr, align 8
 
 // Check for default constructor of accessor
 // CHECK: call spir_func {{.*}}accessor
@@ -43,12 +43,12 @@ int main() {
 // CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0
 
 // Check load from kernel pointer argument alloca
-// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr
+// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(11)*, i32 addrspace(11)** [[MEM_ARG]].addr
 
 // Check accessor __init method call
-// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
+// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
 // CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
 
 // Check lambda "()" operator call
 // CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]])
diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp
index dd502fa1dd844..d7f1f7cfe59f8 100644
--- a/clang/test/CodeGenSYCL/kernel-metadata.cpp
+++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_addr_space ![[MDAS:[0-9]+]] !kernel_arg_access_qual ![[MDAC:[0-9]+]] !kernel_arg_type ![[MDAT:[0-9]+]] !kernel_arg_base_type ![[MDAT:[0-9]+]] !kernel_arg_type_qual ![[MDATQ:[0-9]+]]
-// CHECK: ![[MDAS]] = !{i32 1, i32 0, i32 0, i32 0}
+// CHECK: ![[MDAS]] = !{i32 11, i32 0, i32 0, i32 0}
 // CHECK: ![[MDAC]] = !{!"none", !"none", !"none", !"none"}
 // CHECK: ![[MDAT]] = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"}
 // CHECK: ![[MDATQ]] = !{!"", !"", !"", !""}
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 5bd37447ce814..670c57c65c898 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -33,7 +33,9 @@ enum class address_space : int {
   private_space = 0,
   global_space,
   constant_space,
-  local_space
+  local_space,
+  device_space,
+  host_space
 };
 } // namespace access
 
@@ -57,7 +59,7 @@ struct DeviceValueType;
 
 template <typename dataT>
 struct DeviceValueType<dataT, access::target::global_buffer> {
-  using type = __attribute__((opencl_global)) dataT;
+  using type = __attribute__((opencl_global_device)) dataT;
 };
 
 template <typename dataT>
diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp
index dbaab2664e95c..b03f4dbaa11ba 100644
--- a/clang/test/SemaSYCL/accessors-targets.cpp
+++ b/clang/test/SemaSYCL/accessors-targets.cpp
@@ -37,5 +37,5 @@ int main() {
       });
 }
 // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
-// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_global{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 // CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
index 1f500eff0a888..07420f7f296ff 100644
--- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
@@ -23,11 +23,11 @@ int main() {
 
 // Check declaration of the kernel
 
-// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 
 // Check parameters of the kernel
 
-// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global_device int *'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
@@ -47,7 +47,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var
 
 // CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *'
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global_device int *'
 
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp
index 24d36a6ba54b6..1c911ae2f1233 100644
--- a/clang/test/SemaSYCL/fake-accessors.cpp
+++ b/clang/test/SemaSYCL/fake-accessors.cpp
@@ -51,6 +51,6 @@ int main() {
         });
   return 0;
 }
-// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
-// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
-// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: fake_accessors{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_typedef{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_alias{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index 83bb3ff2448fb..18fe7e70d9591 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -23,11 +23,11 @@ int main() {
 }
 
 // Check declaration of the kernel
-// CHECK: wrapped_access{{.*}} 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: wrapped_access{{.*}} 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>, __global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 
 // Check parameters of the kernel
 // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>'
-// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *'
+// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global_device int *'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
@@ -49,7 +49,7 @@ int main() {
 
 // Parameters of the _init method
 // CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *'
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global_device int *'
 
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp
index 10101d02435f5..7e8348d274fef 100644
--- a/sycl/include/CL/sycl/access/access.hpp
+++ b/sycl/include/CL/sycl/access/access.hpp
@@ -45,7 +45,9 @@ enum class address_space : int {
   private_space = 0,
   global_space,
   constant_space,
-  local_space
+  local_space,
+  device_space = 11,
+  host_space = 12
 };
 
 } // namespace access
@@ -103,11 +105,15 @@ constexpr bool modeWritesNewData(access::mode m) {
 
 #ifdef __SYCL_DEVICE_ONLY__
 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
+#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
+#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
 #else
 #define __OPENCL_GLOBAL_AS__
+#define __OPENCL_GLOBAL_DEVICE_AS__
+#define __OPENCL_GLOBAL_HOST_AS__
 #define __OPENCL_LOCAL_AS__
 #define __OPENCL_CONSTANT_AS__
 #define __OPENCL_PRIVATE_AS__
@@ -118,6 +124,11 @@ template <access::target accessTarget> struct TargetToAS {
       access::address_space::global_space;
 };
 
+template <> struct TargetToAS<access::target::global_buffer> {
+  constexpr static access::address_space AS =
+      access::address_space::device_space;
+};
+
 template <> struct TargetToAS<access::target::local> {
   constexpr static access::address_space AS =
       access::address_space::local_space;
@@ -141,6 +152,16 @@ struct PtrValueType<ElementType, access::address_space::global_space> {
   using type = __OPENCL_GLOBAL_AS__ ElementType;
 };
 
+template <typename ElementType>
+struct PtrValueType<ElementType, access::address_space::device_space> {
+  using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
+};
+
+template <typename ElementType>
+struct PtrValueType<ElementType, access::address_space::host_space> {
+  using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
+};
+
 template <typename ElementType>
 struct PtrValueType<ElementType, access::address_space::constant_space> {
   // Current implementation of address spaces handling leads to possibility
@@ -171,6 +192,14 @@ struct remove_AS<__OPENCL_GLOBAL_AS__ T> {
   typedef T type;
 };
 
+template <class T> struct remove_AS<__OPENCL_GLOBAL_DEVICE_AS__ T> {
+  typedef T type;
+};
+
+template <class T> struct remove_AS<__OPENCL_GLOBAL_HOST_AS__ T> {
+  typedef T type;
+};
+
 template <class T>
 struct remove_AS<__OPENCL_PRIVATE_AS__ T> {
   typedef T type;
@@ -188,6 +217,8 @@ struct remove_AS<__OPENCL_CONSTANT_AS__ T> {
 #endif
 
 #undef __OPENCL_GLOBAL_AS__
+#undef __OPENCL_GLOBAL_DEVICE_AS__
+#undef __OPENCL_GLOBAL_HOST_AS__
 #undef __OPENCL_LOCAL_AS__
 #undef __OPENCL_CONSTANT_AS__
 #undef __OPENCL_PRIVATE_AS__
diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp
index da9daa465efdd..5244ebb637dd8 100644
--- a/sycl/include/CL/sycl/atomic.hpp
+++ b/sycl/include/CL/sycl/atomic.hpp
@@ -47,7 +47,8 @@ template <typename T> struct IsValidAtomicType {
 
 template <cl::sycl::access::address_space AS> struct IsValidAtomicAddressSpace {
   static constexpr bool value = (AS == access::address_space::global_space ||
-                                 AS == access::address_space::local_space);
+                                 AS == access::address_space::local_space ||
+                                 AS == access::address_space::device_space);
 };
 
 // Type trait to translate a cl::sycl::access::address_space to
@@ -56,6 +57,9 @@ template <access::address_space AS> struct GetSpirvMemoryScope {};
 template <> struct GetSpirvMemoryScope<access::address_space::global_space> {
   static constexpr auto scope = __spv::Scope::Device;
 };
+template <> struct GetSpirvMemoryScope<access::address_space::device_space> {
+  static constexpr auto scope = __spv::Scope::Device;
+};
 template <> struct GetSpirvMemoryScope<access::address_space::local_space> {
   static constexpr auto scope = __spv::Scope::Workgroup;
 };
@@ -168,12 +172,12 @@ template <typename T, access::address_space addressSpace =
                           access::address_space::global_space>
 class atomic {
   static_assert(detail::IsValidAtomicType<T>::value,
-                "Invalid SYCL atomic type.  Valid types are: int, "
-                "unsigned int, long, unsigned long, long long,  unsigned "
+                "Invalid SYCL atomic type. Valid types are: int, "
+                "unsigned int, long, unsigned long, long long, unsigned "
                 "long long, float");
   static_assert(detail::IsValidAtomicAddressSpace<addressSpace>::value,
-                "Invalid SYCL atomic address_space.  Valid address spaces are: "
-                "global_space, local_space");
+                "Invalid SYCL atomic address_space. Valid address spaces are: "
+                "global_space, local_space, device_space");
   static constexpr auto SpirvScope =
       detail::GetSpirvMemoryScope<addressSpace>::scope;
 
diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
index 191b52765c524..64e28b4695669 100644
--- a/sycl/include/CL/sycl/detail/generic_type_lists.hpp
+++ b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
@@ -361,21 +361,20 @@ using nan_list = type_list<gtl::unsigned_short_list, gtl::unsigned_int_list,
 } // namespace gtl
 namespace gvl {
 // address spaces
-using all_address_space_list =
-    address_space_list<access::address_space::local_space,
-                       access::address_space::global_space,
-                       access::address_space::private_space,
-                       access::address_space::constant_space>;
-
-using nonconst_address_space_list =
-    address_space_list<access::address_space::local_space,
-                       access::address_space::global_space,
-                       access::address_space::private_space>;
-
-using nonlocal_address_space_list =
-    address_space_list<access::address_space::global_space,
-                       access::address_space::private_space,
-                       access::address_space::constant_space>;
+using all_address_space_list = address_space_list<
+    access::address_space::local_space, access::address_space::global_space,
+    access::address_space::private_space, access::address_space::constant_space,
+    access::address_space::device_space, access::address_space::host_space>;
+
+using nonconst_address_space_list = address_space_list<
+    access::address_space::local_space, access::address_space::global_space,
+    access::address_space::private_space, access::address_space::device_space,
+    access::address_space::host_space>;
+
+using nonlocal_address_space_list = address_space_list<
+    access::address_space::global_space, access::address_space::private_space,
+    access::address_space::constant_space, access::address_space::device_space,
+    access::address_space::host_space>;
 } // namespace gvl
 } // namespace detail
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 5ff72f711f0d2..c88ca284ab50b 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -505,7 +505,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, T>
   readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
-    atomic<T, access::address_space::global_space> AtomicSrc = Src;
+    atomic<T, access::address_space::device_space> AtomicSrc = Src;
     return AtomicSrc.load();
   }
 
@@ -528,7 +528,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, void>
   writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
-    atomic<T, access::address_space::global_space> AtomicDst = Dst;
+    atomic<T, access::address_space::device_space> AtomicDst = Dst;
     AtomicDst.store(V);
   }
 
diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp
index 4495c654ecb3a..48577933dce14 100644
--- a/sycl/include/CL/sycl/multi_ptr.hpp
+++ b/sycl/include/CL/sycl/multi_ptr.hpp
@@ -108,17 +108,18 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
   }
 
-  // Only if Space == global_space
+  // Only if Space == global_space || device_space
   template <int dimensions, access::mode Mode,
             access::placeholder isPlaceholder,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
-                Space == access::address_space::global_space>::type>
+                (Space == access::address_space::global_space ||
+                 Space == access::address_space::device_space)>::type>
   multi_ptr(accessor<ElementType, dimensions, Mode,
                      access::target::global_buffer, isPlaceholder>
                 Accessor) {
-    m_Pointer = (pointer_t)(Accessor.get_pointer().m_Pointer);
+    m_Pointer = (pointer_t)(Accessor.get_pointer().get());
   }
 
   // Only if Space == local_space
@@ -152,14 +153,16 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   //    2. from multi_ptr<ElementType, Space> to multi_ptr<const ElementType,
   //    Space>
 
-  // Only if Space == global_space and element type is const
-  template <
-      int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-      access::address_space _Space = Space, typename ET = ElementType,
-      typename = typename std::enable_if<
-          _Space == Space && Space == access::address_space::global_space &&
-          std::is_const<ET>::value &&
-          std::is_same<ET, ElementType>::value>::type>
+  // Only if Space == global_space || device_space and element type is const
+  template <int dimensions, access::mode Mode,
+            access::placeholder isPlaceholder,
+            access::address_space _Space = Space, typename ET = ElementType,
+            typename = typename std::enable_if<
+                _Space == Space &&
+                (Space == access::address_space::global_space ||
+                 Space == access::address_space::device_space) &&
+                std::is_const<ET>::value &&
+                std::is_same<ET, ElementType>::value>::type>
   multi_ptr(accessor<typename std::remove_const<ET>::type, dimensions, Mode,
                      access::target::global_buffer, isPlaceholder>
                 Accessor)
@@ -271,6 +274,22 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return multi_ptr(m_Pointer - r);
   }
 
+  // Explicit conversion to global_space
+  // Only available if Space == address_space::device_space ||
+  // Space == address_space::host_space
+  template <access::address_space _Space = Space,
+            typename = typename std::enable_if<
+                _Space == Space &&
+                (Space == access::address_space::device_space ||
+                 Space == access::address_space::host_space)>::type>
+  explicit
+  operator multi_ptr<ElementType, access::address_space::global_space>() const {
+    using global_pointer_t = typename detail::PtrValueType<
+        ElementType, access::address_space::global_space>::type *;
+    return multi_ptr<ElementType, access::address_space::global_space>(
+        (global_pointer_t)m_Pointer);
+  }
+
   // Only if Space == global_space
   template <access::address_space _Space = Space,
             typename = typename std::enable_if<
@@ -345,12 +364,13 @@ template <access::address_space Space> class multi_ptr<void, Space> {
     return *this;
   }
 
-  // Only if Space == global_space
+  // Only if Space == global_space || device_space
   template <typename ElementType, int dimensions, access::mode Mode,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
-                Space == access::address_space::global_space>::type>
+                (Space == access::address_space::global_space ||
+                 Space == access::address_space::device_space)>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::global_buffer,
                access::placeholder::false_t>
@@ -466,12 +486,13 @@ class multi_ptr<const void, Space> {
     return *this;
   }
 
-  // Only if Space == global_space
+  // Only if Space == global_space || device_space
   template <typename ElementType, int dimensions, access::mode Mode,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
-                Space == access::address_space::global_space>::type>
+                (Space == access::address_space::global_space ||
+                 Space == access::address_space::device_space)>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::global_buffer,
                access::placeholder::false_t>
diff --git a/sycl/include/CL/sycl/pointers.hpp b/sycl/include/CL/sycl/pointers.hpp
index 9f91ba70ee6b7..3a4fb3beda66d 100644
--- a/sycl/include/CL/sycl/pointers.hpp
+++ b/sycl/include/CL/sycl/pointers.hpp
@@ -19,6 +19,12 @@ template <typename ElementType, access::address_space Space> class multi_ptr;
 template <typename ElementType>
 using global_ptr = multi_ptr<ElementType, access::address_space::global_space>;
 
+template <typename ElementType>
+using device_ptr = multi_ptr<ElementType, access::address_space::device_space>;
+
+template <typename ElementType>
+using host_ptr = multi_ptr<ElementType, access::address_space::host_space>;
+
 template <typename ElementType>
 using local_ptr = multi_ptr<ElementType, access::address_space::local_space>;
 
diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index 0c4c4a1dd2b49..42d7f6cf4649e 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -4,7 +4,7 @@
 // Check the address space of the pointer in accessor class.
 //
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
-// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(1)* }
+// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(11)* }
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
 //
diff --git a/sycl/test/check_device_code/usm_pointers.cpp b/sycl/test/check_device_code/usm_pointers.cpp
new file mode 100644
index 0000000000000..0bf07764d2420
--- /dev/null
+++ b/sycl/test/check_device_code/usm_pointers.cpp
@@ -0,0 +1,41 @@
+// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning
+// RUN: FileCheck %s --input-file %t.ll
+//
+// Check the address space of the pointer in multi_ptr class
+//
+// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(11)* }
+// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(12)* }
+//
+// CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
+// CHECK: %m_Pointer = getelementptr inbounds %[[DEVPTR_T]]
+// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(11)*, i8 addrspace(11)* addrspace(4)* %m_Pointer
+// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(11)* %[[DEVLOAD]] to i8 addrspace(4)*
+// ret i8 addrspace(4)* %[[DEVCAST]]
+//
+// CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
+// CHECK: %m_Pointer = getelementptr inbounds %[[HOSTPTR_T]]
+// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(12)*, i8 addrspace(12)* addrspace(4)* %m_Pointer
+// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(12)* %[[HOSTLOAD]] to i8 addrspace(4)*
+// ret i8 addrspace(4)* %[[HOSTCAST]]
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+int main() {
+  cl::sycl::queue queue;
+  {
+    queue.submit([&](cl::sycl::handler &cgh) {
+      cgh.single_task<class check_adress_space>([=]() {
+        void *Ptr = nullptr;
+        device_ptr<void> DevPtr(Ptr);
+        host_ptr<void> HostPtr(Ptr);
+        global_ptr<void> GlobPtr = global_ptr<void>(DevPtr);
+        GlobPtr = global_ptr<void>(HostPtr);
+      });
+    });
+    queue.wait();
+  }
+
+  return 0;
+}
diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp
index c2e44f461e1b7..04d978ca8e696 100644
--- a/sycl/test/multi_ptr/multi_ptr.cpp
+++ b/sycl/test/multi_ptr/multi_ptr.cpp
@@ -82,6 +82,7 @@ template <typename T> void testMultPtr() {
         auto local_ptr = make_ptr<T, access::address_space::local_space>(
           localAccessor.get_pointer());
 
+        // General conversions in multi_ptr class
         T *RawPtr = nullptr;
         global_ptr<T> ptr_4(RawPtr);
         ptr_4 = RawPtr;
@@ -92,6 +93,12 @@ template <typename T> void testMultPtr() {
 
         ptr_6 = (void *)RawPtr;
 
+        // Explicit conversions for device_ptr/host_ptr to global_ptr
+        device_ptr<void> ptr_7((void *)RawPtr);
+        global_ptr<void> ptr_8 = global_ptr<void>(ptr_7);
+        host_ptr<void> ptr_9((void *)RawPtr);
+        global_ptr<void> ptr_10 = global_ptr<void>(ptr_9);
+
         innerFunc<T>(wiID.get(0), ptr_1, ptr_2, local_ptr);
       });
     });
@@ -109,12 +116,14 @@ void testMultPtrArrowOperator() {
   point<T> data_1[1] = {1};
   point<T> data_2[1] = {2};
   point<T> data_3[1] = {3};
+  point<T> data_4[1] = {4};
 
   {
     range<1> numOfItems{1};
     buffer<point<T>, 1> bufferData_1(data_1, numOfItems);
     buffer<point<T>, 1> bufferData_2(data_2, numOfItems);
     buffer<point<T>, 1> bufferData_3(data_3, numOfItems);
+    buffer<point<T>, 1> bufferData_4(data_4, numOfItems);
     queue myQueue;
     myQueue.submit([&](handler &cgh) {
       accessor<point<T>, 1, access::mode::read, access::target::global_buffer,
@@ -126,6 +135,9 @@ void testMultPtrArrowOperator() {
       accessor<point<T>, 1, access::mode::read_write, access::target::local,
                access::placeholder::false_t>
           accessorData_3(1, cgh);
+      accessor<point<T>, 1, access::mode::read, access::target::global_buffer,
+               access::placeholder::false_t>
+          accessorData_4(bufferData_4, cgh);
 
       cgh.single_task<class testMultPtrArrowOperatorKernel<T>>([=]() {
         auto ptr_1 = make_ptr<point<T>, access::address_space::global_space>(
@@ -134,10 +146,13 @@ void testMultPtrArrowOperator() {
             accessorData_2.get_pointer());
         auto ptr_3 = make_ptr<point<T>, access::address_space::local_space>(
             accessorData_3.get_pointer());
+        auto ptr_4 = make_ptr<point<T>, access::address_space::device_space>(
+            accessorData_4.get_pointer());
 
         auto x1 = ptr_1->x;
         auto x2 = ptr_2->x;
         auto x3 = ptr_3->x;
+        auto x4 = ptr_4 -> x;
 
         static_assert(std::is_same<decltype(x1), T>::value,
                       "Expected decltype(ptr_1->x) == T");
@@ -145,6 +160,8 @@ void testMultPtrArrowOperator() {
                       "Expected decltype(ptr_2->x) == T");
         static_assert(std::is_same<decltype(x3), T>::value,
                       "Expected decltype(ptr_3->x) == T");
+        static_assert(std::is_same<decltype(x4), T>::value,
+                      "Expected decltype(ptr_4->x) == T");
       });
     });
   }

From 23a8adafaff1a2a28d896f76d18fd1d6f68c5450 Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Fri, 19 Jun 2020 00:54:09 +0300
Subject: [PATCH 2/7] Remove assigned values

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 sycl/include/CL/sycl/access/access.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp
index 7e8348d274fef..dab04441cba0a 100644
--- a/sycl/include/CL/sycl/access/access.hpp
+++ b/sycl/include/CL/sycl/access/access.hpp
@@ -46,8 +46,8 @@ enum class address_space : int {
   global_space,
   constant_space,
   local_space,
-  device_space = 11,
-  host_space = 12
+  device_space,
+  host_space
 };
 
 } // namespace access

From 5aa55f390d6624176a99fde35d42ee3b72e4b49c Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Mon, 22 Jun 2020 18:35:19 +0300
Subject: [PATCH 3/7] Rename device/host_space to global_device/host_space

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 sycl/include/CL/sycl/access/access.hpp        | 10 +++----
 sycl/include/CL/sycl/atomic.hpp               | 12 +++++----
 .../CL/sycl/detail/generic_type_lists.hpp     | 27 +++++++++++--------
 sycl/include/CL/sycl/handler.hpp              |  4 +--
 sycl/include/CL/sycl/multi_ptr.hpp            | 25 ++++++++---------
 sycl/include/CL/sycl/pointers.hpp             |  6 +++--
 sycl/test/multi_ptr/multi_ptr.cpp             |  2 +-
 7 files changed, 48 insertions(+), 38 deletions(-)

diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp
index dab04441cba0a..0aaf4ebd0e89e 100644
--- a/sycl/include/CL/sycl/access/access.hpp
+++ b/sycl/include/CL/sycl/access/access.hpp
@@ -46,8 +46,8 @@ enum class address_space : int {
   global_space,
   constant_space,
   local_space,
-  device_space,
-  host_space
+  global_device_space,
+  global_host_space
 };
 
 } // namespace access
@@ -126,7 +126,7 @@ template <access::target accessTarget> struct TargetToAS {
 
 template <> struct TargetToAS<access::target::global_buffer> {
   constexpr static access::address_space AS =
-      access::address_space::device_space;
+      access::address_space::global_device_space;
 };
 
 template <> struct TargetToAS<access::target::local> {
@@ -153,12 +153,12 @@ struct PtrValueType<ElementType, access::address_space::global_space> {
 };
 
 template <typename ElementType>
-struct PtrValueType<ElementType, access::address_space::device_space> {
+struct PtrValueType<ElementType, access::address_space::global_device_space> {
   using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType;
 };
 
 template <typename ElementType>
-struct PtrValueType<ElementType, access::address_space::host_space> {
+struct PtrValueType<ElementType, access::address_space::global_host_space> {
   using type = __OPENCL_GLOBAL_HOST_AS__ ElementType;
 };
 
diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp
index 5244ebb637dd8..6c0be13b2c523 100644
--- a/sycl/include/CL/sycl/atomic.hpp
+++ b/sycl/include/CL/sycl/atomic.hpp
@@ -46,9 +46,10 @@ template <typename T> struct IsValidAtomicType {
 };
 
 template <cl::sycl::access::address_space AS> struct IsValidAtomicAddressSpace {
-  static constexpr bool value = (AS == access::address_space::global_space ||
-                                 AS == access::address_space::local_space ||
-                                 AS == access::address_space::device_space);
+  static constexpr bool value =
+      (AS == access::address_space::global_space ||
+       AS == access::address_space::local_space ||
+       AS == access::address_space::global_device_space);
 };
 
 // Type trait to translate a cl::sycl::access::address_space to
@@ -57,7 +58,8 @@ template <access::address_space AS> struct GetSpirvMemoryScope {};
 template <> struct GetSpirvMemoryScope<access::address_space::global_space> {
   static constexpr auto scope = __spv::Scope::Device;
 };
-template <> struct GetSpirvMemoryScope<access::address_space::device_space> {
+template <>
+struct GetSpirvMemoryScope<access::address_space::global_device_space> {
   static constexpr auto scope = __spv::Scope::Device;
 };
 template <> struct GetSpirvMemoryScope<access::address_space::local_space> {
@@ -177,7 +179,7 @@ class atomic {
                 "long long, float");
   static_assert(detail::IsValidAtomicAddressSpace<addressSpace>::value,
                 "Invalid SYCL atomic address_space. Valid address spaces are: "
-                "global_space, local_space, device_space");
+                "global_space, local_space, global_device_space");
   static constexpr auto SpirvScope =
       detail::GetSpirvMemoryScope<addressSpace>::scope;
 
diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
index 64e28b4695669..9965ea66eee9a 100644
--- a/sycl/include/CL/sycl/detail/generic_type_lists.hpp
+++ b/sycl/include/CL/sycl/detail/generic_type_lists.hpp
@@ -364,17 +364,22 @@ namespace gvl {
 using all_address_space_list = address_space_list<
     access::address_space::local_space, access::address_space::global_space,
     access::address_space::private_space, access::address_space::constant_space,
-    access::address_space::device_space, access::address_space::host_space>;
-
-using nonconst_address_space_list = address_space_list<
-    access::address_space::local_space, access::address_space::global_space,
-    access::address_space::private_space, access::address_space::device_space,
-    access::address_space::host_space>;
-
-using nonlocal_address_space_list = address_space_list<
-    access::address_space::global_space, access::address_space::private_space,
-    access::address_space::constant_space, access::address_space::device_space,
-    access::address_space::host_space>;
+    access::address_space::global_device_space,
+    access::address_space::global_host_space>;
+
+using nonconst_address_space_list =
+    address_space_list<access::address_space::local_space,
+                       access::address_space::global_space,
+                       access::address_space::private_space,
+                       access::address_space::global_device_space,
+                       access::address_space::global_host_space>;
+
+using nonlocal_address_space_list =
+    address_space_list<access::address_space::global_space,
+                       access::address_space::private_space,
+                       access::address_space::constant_space,
+                       access::address_space::global_device_space,
+                       access::address_space::global_host_space>;
 } // namespace gvl
 } // namespace detail
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index c88ca284ab50b..b06cbc412a0b1 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -505,7 +505,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, T>
   readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
-    atomic<T, access::address_space::device_space> AtomicSrc = Src;
+    atomic<T, access::address_space::global_device_space> AtomicSrc = Src;
     return AtomicSrc.load();
   }
 
@@ -528,7 +528,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, void>
   writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
-    atomic<T, access::address_space::device_space> AtomicDst = Dst;
+    atomic<T, access::address_space::global_device_space> AtomicDst = Dst;
     AtomicDst.store(V);
   }
 
diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp
index 48577933dce14..764efac6b0e4a 100644
--- a/sycl/include/CL/sycl/multi_ptr.hpp
+++ b/sycl/include/CL/sycl/multi_ptr.hpp
@@ -108,14 +108,14 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
   }
 
-  // Only if Space == global_space || device_space
+  // Only if Space == global_space || global_device_space
   template <int dimensions, access::mode Mode,
             access::placeholder isPlaceholder,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
-                 Space == access::address_space::device_space)>::type>
+                 Space == access::address_space::global_device_space)>::type>
   multi_ptr(accessor<ElementType, dimensions, Mode,
                      access::target::global_buffer, isPlaceholder>
                 Accessor) {
@@ -153,14 +153,15 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   //    2. from multi_ptr<ElementType, Space> to multi_ptr<const ElementType,
   //    Space>
 
-  // Only if Space == global_space || device_space and element type is const
+  // Only if Space == global_space || global_device_space and element type is
+  // const
   template <int dimensions, access::mode Mode,
             access::placeholder isPlaceholder,
             access::address_space _Space = Space, typename ET = ElementType,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
-                 Space == access::address_space::device_space) &&
+                 Space == access::address_space::global_device_space) &&
                 std::is_const<ET>::value &&
                 std::is_same<ET, ElementType>::value>::type>
   multi_ptr(accessor<typename std::remove_const<ET>::type, dimensions, Mode,
@@ -275,13 +276,13 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   }
 
   // Explicit conversion to global_space
-  // Only available if Space == address_space::device_space ||
-  // Space == address_space::host_space
+  // Only available if Space == address_space::global_device_space ||
+  // Space == address_space::global_host_space
   template <access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
-                (Space == access::address_space::device_space ||
-                 Space == access::address_space::host_space)>::type>
+                (Space == access::address_space::global_device_space ||
+                 Space == access::address_space::global_host_space)>::type>
   explicit
   operator multi_ptr<ElementType, access::address_space::global_space>() const {
     using global_pointer_t = typename detail::PtrValueType<
@@ -364,13 +365,13 @@ template <access::address_space Space> class multi_ptr<void, Space> {
     return *this;
   }
 
-  // Only if Space == global_space || device_space
+  // Only if Space == global_space || global_device_space
   template <typename ElementType, int dimensions, access::mode Mode,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
-                 Space == access::address_space::device_space)>::type>
+                 Space == access::address_space::global_device_space)>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::global_buffer,
                access::placeholder::false_t>
@@ -486,13 +487,13 @@ class multi_ptr<const void, Space> {
     return *this;
   }
 
-  // Only if Space == global_space || device_space
+  // Only if Space == global_space || global_device_space
   template <typename ElementType, int dimensions, access::mode Mode,
             access::address_space _Space = Space,
             typename = typename std::enable_if<
                 _Space == Space &&
                 (Space == access::address_space::global_space ||
-                 Space == access::address_space::device_space)>::type>
+                 Space == access::address_space::global_device_space)>::type>
   multi_ptr(
       accessor<ElementType, dimensions, Mode, access::target::global_buffer,
                access::placeholder::false_t>
diff --git a/sycl/include/CL/sycl/pointers.hpp b/sycl/include/CL/sycl/pointers.hpp
index 3a4fb3beda66d..efec74e0fd3a6 100644
--- a/sycl/include/CL/sycl/pointers.hpp
+++ b/sycl/include/CL/sycl/pointers.hpp
@@ -20,10 +20,12 @@ template <typename ElementType>
 using global_ptr = multi_ptr<ElementType, access::address_space::global_space>;
 
 template <typename ElementType>
-using device_ptr = multi_ptr<ElementType, access::address_space::device_space>;
+using device_ptr =
+    multi_ptr<ElementType, access::address_space::global_device_space>;
 
 template <typename ElementType>
-using host_ptr = multi_ptr<ElementType, access::address_space::host_space>;
+using host_ptr =
+    multi_ptr<ElementType, access::address_space::global_host_space>;
 
 template <typename ElementType>
 using local_ptr = multi_ptr<ElementType, access::address_space::local_space>;
diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp
index 04d978ca8e696..bd394dbd559d6 100644
--- a/sycl/test/multi_ptr/multi_ptr.cpp
+++ b/sycl/test/multi_ptr/multi_ptr.cpp
@@ -146,7 +146,7 @@ void testMultPtrArrowOperator() {
             accessorData_2.get_pointer());
         auto ptr_3 = make_ptr<point<T>, access::address_space::local_space>(
             accessorData_3.get_pointer());
-        auto ptr_4 = make_ptr<point<T>, access::address_space::device_space>(
+        auto ptr_4 = make_ptr<point<T>, access::address_space::global_device_space>(
             accessorData_4.get_pointer());
 
         auto x1 = ptr_1->x;

From 2c832ee9466b3ab990c670f3619ff2542198c4f6 Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Wed, 24 Jun 2020 20:50:51 +0300
Subject: [PATCH 4/7] Remove clang testing

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 clang/test/CodeGenSYCL/Inputs/sycl.hpp          |  9 +++------
 clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp | 12 ++++++------
 clang/test/CodeGenSYCL/kernel-metadata.cpp      |  2 +-
 clang/test/SemaSYCL/Inputs/sycl.hpp             |  6 ++----
 clang/test/SemaSYCL/accessors-targets.cpp       |  2 +-
 clang/test/SemaSYCL/basic-kernel-wrapper.cpp    |  6 +++---
 clang/test/SemaSYCL/fake-accessors.cpp          |  6 +++---
 clang/test/SemaSYCL/wrapped-accessor.cpp        |  6 +++---
 8 files changed, 22 insertions(+), 27 deletions(-)

diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index f1ad520ee248a..3184c58edcbfc 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -57,9 +57,7 @@ enum class address_space : int {
   private_space = 0,
   global_space,
   constant_space,
-  local_space,
-  device_space,
-  host_space
+  local_space
 };
 } // namespace access
 
@@ -141,9 +139,8 @@ class accessor {
   _ImplT<dimensions> impl;
 
 private:
-  void __init(__attribute__((opencl_global_device)) dataT *Ptr,
-              range<dimensions> AccessRange, range<dimensions> MemRange,
-              id<dimensions> Offset) {}
+  void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
+              range<dimensions> MemRange, id<dimensions> Offset) {}
 };
 
 template <int dimensions, access::mode accessmode, access::target accesstarget>
diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
index ceab1dee68606..31795fc73b776 100644
--- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
@@ -20,12 +20,12 @@ int main() {
 }
 
 // CHECK: define spir_kernel void @{{.*}}kernel_function
-// CHECK-SAME: i32 addrspace(11)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
 // Check alloca for pointer argument
-// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(11)*
+// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
 // Check lambda object alloca
 // CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon"
 // Check allocas for ranges
@@ -34,7 +34,7 @@ int main() {
 // CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id"
 //
 // Check store of kernel pointer argument to alloca
-// CHECK: store i32 addrspace(11)* [[MEM_ARG]], i32 addrspace(11)** [[MEM_ARG]].addr, align 8
+// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8
 
 // Check for default constructor of accessor
 // CHECK: call spir_func {{.*}}accessor
@@ -43,12 +43,12 @@ int main() {
 // CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0
 
 // Check load from kernel pointer argument alloca
-// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(11)*, i32 addrspace(11)** [[MEM_ARG]].addr
+// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr
 
 // Check accessor __init method call
-// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
+// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
 // CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
 
 // Check lambda "()" operator call
 // CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]])
diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp
index d7f1f7cfe59f8..dd502fa1dd844 100644
--- a/clang/test/CodeGenSYCL/kernel-metadata.cpp
+++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_addr_space ![[MDAS:[0-9]+]] !kernel_arg_access_qual ![[MDAC:[0-9]+]] !kernel_arg_type ![[MDAT:[0-9]+]] !kernel_arg_base_type ![[MDAT:[0-9]+]] !kernel_arg_type_qual ![[MDATQ:[0-9]+]]
-// CHECK: ![[MDAS]] = !{i32 11, i32 0, i32 0, i32 0}
+// CHECK: ![[MDAS]] = !{i32 1, i32 0, i32 0, i32 0}
 // CHECK: ![[MDAC]] = !{!"none", !"none", !"none", !"none"}
 // CHECK: ![[MDAT]] = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"}
 // CHECK: ![[MDATQ]] = !{!"", !"", !"", !""}
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 670c57c65c898..5bd37447ce814 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -33,9 +33,7 @@ enum class address_space : int {
   private_space = 0,
   global_space,
   constant_space,
-  local_space,
-  device_space,
-  host_space
+  local_space
 };
 } // namespace access
 
@@ -59,7 +57,7 @@ struct DeviceValueType;
 
 template <typename dataT>
 struct DeviceValueType<dataT, access::target::global_buffer> {
-  using type = __attribute__((opencl_global_device)) dataT;
+  using type = __attribute__((opencl_global)) dataT;
 };
 
 template <typename dataT>
diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp
index b03f4dbaa11ba..dbaab2664e95c 100644
--- a/clang/test/SemaSYCL/accessors-targets.cpp
+++ b/clang/test/SemaSYCL/accessors-targets.cpp
@@ -37,5 +37,5 @@ int main() {
       });
 }
 // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
-// CHECK: {{.*}}use_global{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 // CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
index 07420f7f296ff..1f500eff0a888 100644
--- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
@@ -23,11 +23,11 @@ int main() {
 
 // Check declaration of the kernel
 
-// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 
 // Check parameters of the kernel
 
-// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global_device int *'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
@@ -47,7 +47,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var
 
 // CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global_device int *'
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *'
 
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp
index 1c911ae2f1233..24d36a6ba54b6 100644
--- a/clang/test/SemaSYCL/fake-accessors.cpp
+++ b/clang/test/SemaSYCL/fake-accessors.cpp
@@ -51,6 +51,6 @@ int main() {
         });
   return 0;
 }
-// CHECK: fake_accessors{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
-// CHECK: accessor_typedef{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
-// CHECK: accessor_alias{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index 18fe7e70d9591..83bb3ff2448fb 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -23,11 +23,11 @@ int main() {
 }
 
 // Check declaration of the kernel
-// CHECK: wrapped_access{{.*}} 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>, __global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: wrapped_access{{.*}} 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 
 // Check parameters of the kernel
 // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>'
-// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global_device int *'
+// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
@@ -49,7 +49,7 @@ int main() {
 
 // Parameters of the _init method
 // CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global_device int *'
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *'
 
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>

From 2ef86a0c0c86fed2e887724798fb39defca9dffd Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Thu, 25 Jun 2020 17:42:30 +0300
Subject: [PATCH 5/7] Update ASi values

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 sycl/test/check_device_code/kernel_arguments_as.cpp |  2 +-
 sycl/test/check_device_code/usm_pointers.cpp        | 12 ++++++------
 2 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index 42d7f6cf4649e..7faef5ec72444 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -4,7 +4,7 @@
 // Check the address space of the pointer in accessor class.
 //
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
-// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(11)* }
+// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(5)* }
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
 //
diff --git a/sycl/test/check_device_code/usm_pointers.cpp b/sycl/test/check_device_code/usm_pointers.cpp
index 0bf07764d2420..aa0a0ed58045d 100644
--- a/sycl/test/check_device_code/usm_pointers.cpp
+++ b/sycl/test/check_device_code/usm_pointers.cpp
@@ -3,19 +3,19 @@
 //
 // Check the address space of the pointer in multi_ptr class
 //
-// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(11)* }
-// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(12)* }
+// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(5)* }
+// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(6)* }
 //
 // CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
 // CHECK: %m_Pointer = getelementptr inbounds %[[DEVPTR_T]]
-// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(11)*, i8 addrspace(11)* addrspace(4)* %m_Pointer
-// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(11)* %[[DEVLOAD]] to i8 addrspace(4)*
+// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(5)*, i8 addrspace(5)* addrspace(4)* %m_Pointer
+// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(5)* %[[DEVLOAD]] to i8 addrspace(4)*
 // ret i8 addrspace(4)* %[[DEVCAST]]
 //
 // CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
 // CHECK: %m_Pointer = getelementptr inbounds %[[HOSTPTR_T]]
-// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(12)*, i8 addrspace(12)* addrspace(4)* %m_Pointer
-// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(12)* %[[HOSTLOAD]] to i8 addrspace(4)*
+// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(6)*, i8 addrspace(6)* addrspace(4)* %m_Pointer
+// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(6)* %[[HOSTLOAD]] to i8 addrspace(4)*
 // ret i8 addrspace(4)* %[[HOSTCAST]]
 
 #include <CL/sycl.hpp>

From fc5aa11379a3e97a6d26318a09d29f1d96bc84b9 Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Wed, 1 Jul 2020 16:48:03 +0300
Subject: [PATCH 6/7] Move accessor pointers back to global space

Otherwise it breaks atomics.

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 sycl/include/CL/sycl/access/access.hpp           |  5 -----
 sycl/include/CL/sycl/handler.hpp                 |  4 ++--
 sycl/include/CL/sycl/multi_ptr.hpp               | 16 ----------------
 .../check_device_code/kernel_arguments_as.cpp    |  2 +-
 4 files changed, 3 insertions(+), 24 deletions(-)

diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp
index 0aaf4ebd0e89e..9187a972bd6fb 100644
--- a/sycl/include/CL/sycl/access/access.hpp
+++ b/sycl/include/CL/sycl/access/access.hpp
@@ -124,11 +124,6 @@ template <access::target accessTarget> struct TargetToAS {
       access::address_space::global_space;
 };
 
-template <> struct TargetToAS<access::target::global_buffer> {
-  constexpr static access::address_space AS =
-      access::address_space::global_device_space;
-};
-
 template <> struct TargetToAS<access::target::local> {
   constexpr static access::address_space AS =
       access::address_space::local_space;
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index b06cbc412a0b1..5ff72f711f0d2 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -505,7 +505,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, T>
   readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
-    atomic<T, access::address_space::global_device_space> AtomicSrc = Src;
+    atomic<T, access::address_space::global_space> AtomicSrc = Src;
     return AtomicSrc.load();
   }
 
@@ -528,7 +528,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, void>
   writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
-    atomic<T, access::address_space::global_device_space> AtomicDst = Dst;
+    atomic<T, access::address_space::global_space> AtomicDst = Dst;
     AtomicDst.store(V);
   }
 
diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp
index 764efac6b0e4a..1a59113d9fc18 100644
--- a/sycl/include/CL/sycl/multi_ptr.hpp
+++ b/sycl/include/CL/sycl/multi_ptr.hpp
@@ -275,22 +275,6 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return multi_ptr(m_Pointer - r);
   }
 
-  // Explicit conversion to global_space
-  // Only available if Space == address_space::global_device_space ||
-  // Space == address_space::global_host_space
-  template <access::address_space _Space = Space,
-            typename = typename std::enable_if<
-                _Space == Space &&
-                (Space == access::address_space::global_device_space ||
-                 Space == access::address_space::global_host_space)>::type>
-  explicit
-  operator multi_ptr<ElementType, access::address_space::global_space>() const {
-    using global_pointer_t = typename detail::PtrValueType<
-        ElementType, access::address_space::global_space>::type *;
-    return multi_ptr<ElementType, access::address_space::global_space>(
-        (global_pointer_t)m_Pointer);
-  }
-
   // Only if Space == global_space
   template <access::address_space _Space = Space,
             typename = typename std::enable_if<
diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index 7faef5ec72444..0c4c4a1dd2b49 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -4,7 +4,7 @@
 // Check the address space of the pointer in accessor class.
 //
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
-// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(5)* }
+// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(1)* }
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
 //

From 1a1237ca8aa3c4cdaff9afbe846274dcaa7ae309 Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Wed, 1 Jul 2020 16:52:31 +0300
Subject: [PATCH 7/7] Ignore clang-format concern

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
---
 sycl/test/multi_ptr/multi_ptr.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp
index bd394dbd559d6..9ebb33046a459 100644
--- a/sycl/test/multi_ptr/multi_ptr.cpp
+++ b/sycl/test/multi_ptr/multi_ptr.cpp
@@ -152,7 +152,7 @@ void testMultPtrArrowOperator() {
         auto x1 = ptr_1->x;
         auto x2 = ptr_2->x;
         auto x3 = ptr_3->x;
-        auto x4 = ptr_4 -> x;
+        auto x4 = ptr_4->x;
 
         static_assert(std::is_same<decltype(x1), T>::value,
                       "Expected decltype(ptr_1->x) == T");