[SYCL] Enable USM address spaces heneration upder opt (#2127)

This patch serves two purposes:
1. It reverts shift of accessors to global buffers from __global to
   __global_device address space, which was considered as not a good
   idea until proven otherwise. Still this shift is enabled under
   -fsycl_enable_usm_address_spaces option for testing;
2. It fixes a mangling issue of functions that accepts arguments
   of multi_ptr<> type for non-SPIR targets (see:
   #2039)

Signed-off-by: Dmitry Sidorov <[email protected]>

Author: MrSidims

clang/include/clang/Driver/Options.td

@@ -1842,7 +1842,7 @@ def fintelfpga : Flag<["-"], "fintelfpga">, Group<f_Group>,
   Flags<[CC1Option, CoreOption]>, HelpText<"Perform ahead of time compilation for FPGA">;
 def fsycl_enable_usm_address_spaces : Flag<["-"], "fsycl-enable-usm-address-spaces">,
  Group<f_Group>, Flags<[CC1Option, CoreOption]>,
- HelpText<"Enable SPV_INTEL_usm_storage_classes extension">;
+ HelpText<"Enable USM address spaces">;
 def fsycl_device_only : Flag<["-"], "fsycl-device-only">, Flags<[CoreOption]>,
   HelpText<"Compile SYCL kernels for device">;
 def fsycl_targets_EQ : CommaJoined<["-"], "fsycl-targets=">, Flags<[DriverOption, CC1Option, CoreOption]>,

clang/lib/Driver/ToolChains/Clang.cpp

@@ -6183,6 +6183,13 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     }
     if (Args.hasArg(options::OPT_fsycl_unnamed_lambda))
       CmdArgs.push_back("-fsycl-unnamed-lambda");
+
+    // Enable generation of USM address spaces as opt-in.
+    // __ENABLE_USM_ADDR_SPACE__ will be used during compilation of SYCL headers
+    if (getToolChain().getTriple().getSubArch() ==
+            llvm::Triple::SPIRSubArch_fpga &&
+        Args.hasArg(options::OPT_fsycl_enable_usm_address_spaces))
+      CmdArgs.push_back("-D__ENABLE_USM_ADDR_SPACE__");
   }
 
   if (IsHIP)

clang/test/Driver/sycl-offload.c

@@ -668,7 +668,8 @@
 // CHK-TOOLS-GEN: clang-offload-wrapper{{.*}} "-o=[[OUTPUT5:.+\.bc]]" "-host=x86_64-unknown-linux-gnu" "-target=spir64_gen{{.*}}" "-kind=sycl" "[[OUTPUT4]]"
 // CHK-TOOLS-CPU: clang-offload-wrapper{{.*}} "-o=[[OUTPUT5:.+\.bc]]" "-host=x86_64-unknown-linux-gnu" "-target=spir64_x86_64{{.*}}" "-kind=sycl" "[[OUTPUT4]]"
 // CHK-TOOLS-AOT: llc{{.*}} "-filetype=obj" "-o" "[[OUTPUT6:.+\.o]]" "[[OUTPUT5]]"
-// CHK-TOOLS-FPGA: clang{{.*}} "-triple" "spir64_fpga-unknown-unknown-sycldevice" {{.*}} "-fsycl-int-header=[[INPUT1:.+\.h]]" "-faddrsig"
+// CHK-TOOLS-FPGA-USM-DISABLE: clang{{.*}} "-triple" "spir64_fpga-unknown-unknown-sycldevice" {{.*}} "-fsycl-int-header=[[INPUT1:.+\.h]]" "-faddrsig"
+// CHK-TOOLS-FPGA-USM-ENABLE: clang{{.*}} "-triple" "spir64_fpga-unknown-unknown-sycldevice" {{.*}} "-fsycl-int-header=[[INPUT1:.+\.h]]" "-D__ENABLE_USM_ADDR_SPACE__" "-faddrsig"
 // CHK-TOOLS-GEN: clang{{.*}} "-triple" "spir64_gen-unknown-unknown-sycldevice" {{.*}} "-fsycl-int-header=[[INPUT1:.+\.h]]" "-faddrsig"
 // CHK-TOOLS-CPU: clang{{.*}} "-triple" "spir64_x86_64-unknown-unknown-sycldevice" {{.*}} "-fsycl-int-header=[[INPUT1:.+\.h]]" "-faddrsig"
 // CHK-TOOLS-AOT: clang{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-include" "[[INPUT1]]" {{.*}} "-o" "[[OUTPUT7:.+\.o]]"

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

@@ -105,8 +105,13 @@ constexpr bool modeWritesNewData(access::mode m) {
 
 #ifdef __SYCL_DEVICE_ONLY__
 #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global))
+#ifdef __ENABLE_USM_ADDR_SPACE__
 #define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device))
 #define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host))
+#else
+#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global))
+#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global))
+#endif // __ENABLE_USM_ADDR_SPACE__
 #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
 #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
 #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
@@ -124,10 +129,12 @@ template <access::target accessTarget> struct TargetToAS {
       access::address_space::global_space;
 };
 
+#ifdef __ENABLE_USM_ADDR_SPACE__
 template <> struct TargetToAS<access::target::global_buffer> {
   constexpr static access::address_space AS =
       access::address_space::global_device_space;
 };
+#endif // __ENABLE_USM_ADDR_SPACE__
 
 template <> struct TargetToAS<access::target::local> {
   constexpr static access::address_space AS =
@@ -192,13 +199,15 @@ struct remove_AS<__OPENCL_GLOBAL_AS__ T> {
   typedef T type;
 };
 
+#ifdef __ENABLE_USM_ADDR_SPACE__
 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;
 };
+#endif // __ENABLE_USM_ADDR_SPACE__
 
 template <class T>
 struct remove_AS<__OPENCL_PRIVATE_AS__ T> {

sycl/include/CL/sycl/atomic.hpp

@@ -198,6 +198,7 @@ class atomic {
                   "T and pointerT must be same size");
   }
 
+#ifdef __ENABLE_USM_ADDR_SPACE__
   // Create atomic in global_space with one from global_device_space
   template <access::address_space _Space = addressSpace,
             typename = typename std::enable_if<
@@ -214,6 +215,7 @@ class atomic {
   atomic(atomic<T, access::address_space::global_device_space> &&RHS) {
     Ptr = RHS.Ptr;
   }
+#endif // __ENABLE_USM_ADDR_SPACE__
 
   void store(T Operand, memory_order Order = memory_order::relaxed) {
     __spirv_AtomicStore(

sycl/include/CL/sycl/handler.hpp

@@ -506,7 +506,11 @@ 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 {
+#ifdef __ENABLE_USM_ADDR_SPACE__
     atomic<T, access::address_space::global_device_space> AtomicSrc = Src;
+#else
+    atomic<T, access::address_space::global_space> AtomicSrc = Src;
+#endif // __ENABLE_USM_ADDR_SPACE__
     return AtomicSrc.load();
   }
 
@@ -529,7 +533,11 @@ 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 {
+#ifdef __ENABLE_USM_ADDR_SPACE__
     atomic<T, access::address_space::global_device_space> AtomicDst = Dst;
+#else
+    atomic<T, access::address_space::global_space> AtomicDst = Dst;
+#endif // __ENABLE_USM_ADDR_SPACE__
     AtomicDst.store(V);
   }
 

sycl/include/CL/sycl/multi_ptr.hpp

@@ -275,6 +275,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return multi_ptr(m_Pointer - r);
   }
 
+#ifdef __ENABLE_USM_ADDR_SPACE__
   // Explicit conversion to global_space
   // Only available if Space == address_space::global_device_space ||
   // Space == address_space::global_host_space
@@ -290,6 +291,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return multi_ptr<ElementType, access::address_space::global_space>(
         reinterpret_cast<global_pointer_t>(m_Pointer));
   }
+#endif // __ENABLE_USM_ADDR_SPACE__
 
   // Only if Space == global_space
   template <access::address_space _Space = Space,

sycl/test/check_device_code/kernel_arguments_as.cpp

@@ -1,11 +1,14 @@
 // 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 -Xclang -disable-llvm-passes
-// RUN: FileCheck %s --input-file %t.ll
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
+// 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 -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
 //
 // Check the address space of the pointer in accessor class.
 //
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
-// CHECK: %[[UNION]] = type { i32 addrspace(5)* }
+// CHECK-DISABLE: %[[UNION]] = type { i32 addrspace(1)* }
+// CHECK-ENABLE: %[[UNION]] = type { i32 addrspace(5)* }
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
 //

sycl/test/check_device_code/usm_pointers.cpp

@@ -1,21 +1,29 @@
 // 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 -Xclang -disable-llvm-passes
-// RUN: FileCheck %s --input-file %t.ll
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
+// 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 -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
 //
 // Check the address space of the pointer in multi_ptr class
 //
-// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(5)* }
-// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(6)* }
+// CHECK-DISABLE: %[[DEVPTR_T:.*]] = type { i8 addrspace(1)* }
+// CHECK-DISABLE: %[[HOSTPTR_T:.*]] = type { i8 addrspace(1)* }
+// CHECK-ENABLE: %[[DEVPTR_T:.*]] = type { i8 addrspace(5)* }
+// CHECK-ENABLE: %[[HOSTPTR_T:.*]] = type { i8 addrspace(6)* }
 //
 // CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
 // CHECK: %[[M_PTR:.*]] = getelementptr inbounds %[[DEVPTR_T]]
-// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(5)*, i8 addrspace(5)* addrspace(4)* %[[M_PTR]]
-// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(5)* %[[DEVLOAD]] to i8 addrspace(4)*
+// CHECK-DISABLE-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* %[[M_PTR]]
+// CHECK-DISABLE-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[DEVLOAD]] to i8 addrspace(4)*
+// CHECK-ENABLE-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(5)*, i8 addrspace(5)* addrspace(4)* %[[M_PTR]]
+// CHECK-ENABLE-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_PTR]] = getelementptr inbounds %[[HOSTPTR_T]]
-// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(6)*, i8 addrspace(6)* addrspace(4)* %[[M_PTR]]
-// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(6)* %[[HOSTLOAD]] to i8 addrspace(4)*
+// CHECK-DISABLE-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* %[[M_PTR]]
+// CHECK-DISABLE-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[HOSTLOAD]] to i8 addrspace(4)*
+// CHECK-ENABLE-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(6)*, i8 addrspace(6)* addrspace(4)* %[[M_PTR]]
+// CHECK-ENABLE-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(6)* %[[HOSTLOAD]] to i8 addrspace(4)*
 // ret i8 addrspace(4)* %[[HOSTCAST]]
 
 #include <CL/sycl.hpp>