Merge remote-tracking branch 'upstream/sycl' into stmt-attr-cleanups

Author: AaronBallman

clang/lib/Frontend/InitPreprocessor.cpp

@@ -1129,6 +1129,12 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
     Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1");
     Builder.defineMacro("SYCL_EXTERNAL", "__attribute__((sycl_device))");
 
+    // Enable __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ macro for
+    // all FPGA compilations.
+    if (TI.getTriple().getSubArch() == llvm::Triple::SPIRSubArch_fpga) {
+      Builder.defineMacro("__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__", "1");
+    }
+
     if (TI.getTriple().isNVPTX()) {
         Builder.defineMacro("__SYCL_NVPTX__", "1");
     }

clang/test/Preprocessor/predefined-macros.c

@@ -214,3 +214,16 @@
 // CHECK-HIP-DEV: #define __HIPCC__ 1
 // CHECK-HIP-DEV: #define __HIP_DEVICE_COMPILE__ 1
 // CHECK-HIP-DEV: #define __HIP__ 1
+
+// RUN: %clang_cc1 %s -E -dM -fsycl-is-device \
+// RUN:   -triple spir64_fpga-unknown-unknown-sycldevice -o - \
+// RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-RANGE
+// CHECK-RANGE: #define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1
+
+// RUN: %clang_cc1 %s -E -dM -fsycl-is-device -o - \
+// RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-NO-RANGE
+
+// RUN: %clang_cc1 %s -E -dM -o - \
+// RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-NO-RANGE
+
+// CHECK-NO-RANGE-NOT: #define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -322,12 +322,13 @@ extern SYCL_EXTERNAL ap_int<Wout> __spirv_ArbitraryFloatCastINTEL(
 
 template <int WA, int Wout>
 extern SYCL_EXTERNAL ap_int<Wout> __spirv_ArbitraryFloatCastFromIntINTEL(
-    ap_int<WA> A, int32_t Mout, int32_t EnableSubnormals = 0,
-    int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
+    ap_int<WA> A, int32_t Mout, bool FromSign = false,
+    int32_t EnableSubnormals = 0, int32_t RoundingMode = 0,
+    int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int Wout>
 extern SYCL_EXTERNAL ap_int<Wout> __spirv_ArbitraryFloatCastToIntINTEL(
-    ap_int<WA> A, int32_t MA, int32_t EnableSubnormals = 0,
+    ap_int<WA> A, int32_t MA, bool ToSign = false, int32_t EnableSubnormals = 0,
     int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept;
 
 template <int WA, int WB, int Wout>

sycl/include/CL/sycl/INTEL/fpga_lsu.hpp

@@ -47,7 +47,9 @@ template <class... _mem_access_params> class lsu final {
 public:
   lsu() = delete;
 
-  template <typename _T> static _T load(sycl::global_ptr<_T> Ptr) {
+  template <typename _T, access::address_space _space>
+  static _T load(sycl::multi_ptr<_T, _space> Ptr) {
+    check_space<_space>();
     check_load();
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
     return *__builtin_intel_fpga_mem((_T *)Ptr,
@@ -59,7 +61,9 @@ template <class... _mem_access_params> class lsu final {
 #endif
   }
 
-  template <typename _T> static void store(sycl::global_ptr<_T> Ptr, _T Val) {
+  template <typename _T, access::address_space _space>
+  static void store(sycl::multi_ptr<_T, _space> Ptr, _T Val) {
+    check_space<_space>();
     check_store();
 #if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
     *__builtin_intel_fpga_mem((_T *)Ptr,
@@ -92,6 +96,14 @@ template <class... _mem_access_params> class lsu final {
 
   static_assert(_cache_val >= 0, "cache size parameter must be non-negative");
 
+  template <access::address_space _space> static void check_space() {
+    static_assert(_space == access::address_space::global_space ||
+                      _space == access::address_space::global_device_space ||
+                      _space == access::address_space::global_host_space,
+                  "lsu controls are only supported for global_ptr, "
+                  "device_ptr, and host_ptr objects");
+  }
+
   static void check_load() {
     static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE,
                   "unable to implement a cache without a burst coalescer");

sycl/test/check_device_code/fpga_ihs_float.cpp

@@ -11,7 +11,11 @@
 
 #include "CL/__spirv/spirv_ops.hpp"
 
-constexpr int32_t Subnorm = 0, RndMode = 2, RndAcc = 1;
+constexpr int32_t Subnorm = 0;
+constexpr int32_t RndMode = 2;
+constexpr int32_t RndAcc = 1;
+constexpr bool FromSign = false;
+constexpr bool ToSign = true;
 
 template <int EA, int MA, int Eout, int Mout>
 void ap_float_cast() {
@@ -27,17 +31,17 @@ void ap_float_cast_from_int() {
   ap_int<WA> A;
   ap_int<1 + Eout + Mout> cast_from_int_res =
       __spirv_ArbitraryFloatCastFromIntINTEL<WA, 1 + Eout + Mout>(
-          A, Mout, Subnorm, RndMode, RndAcc);
-  // CHECK: call spir_func signext i25 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastFromIntINTEL{{.*}}(i43 {{[%a-z0-9.]+}}, i32 16, i32 0, i32 2, i32 1)
+          A, Mout, FromSign, Subnorm, RndMode, RndAcc);
+  // CHECK: call spir_func signext i25 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastFromIntINTEL{{.*}}(i43 {{[%a-z0-9.]+}}, i32 16, i1 zeroext false, i32 0, i32 2, i32 1)
 }
 
 template <int EA, int MA, int Wout>
 void ap_float_cast_to_int() {
   ap_int<1 + EA + MA> A;
   ap_int<Wout> cast_to_int_res =
       __spirv_ArbitraryFloatCastToIntINTEL<1 + EA + MA, Wout>(
-          A, MA, Subnorm, RndMode, RndAcc);
-  // CHECK: call spir_func signext i30 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastToIntINTEL{{.*}}(i23 signext {{[%a-z0-9.]+}}, i32 15, i32 0, i32 2, i32 1)
+          A, MA, ToSign, Subnorm, RndMode, RndAcc);
+  // CHECK: call spir_func signext i30 @_Z{{[0-9]+}}__spirv_ArbitraryFloatCastToIntINTEL{{.*}}(i23 signext {{[%a-z0-9.]+}}, i32 15, i1 zeroext true, i32 0, i32 2, i32 1)
 }
 
 template <int EA, int MA, int EB, int MB, int Eout, int Mout>

sycl/test/extensions/fpga.cpp

@@ -59,16 +59,14 @@ int main() {
   /*Check LSU interface*/
   {
     cl::sycl::buffer<int, 1> output_buffer(1);
-    cl::sycl::buffer<int, 1> input_buffer(1);
+    auto *in_ptr = cl::sycl::malloc_host<int>(1, Queue.get_context());
 
     Queue.submit([&](cl::sycl::handler &cgh) {
       auto output_accessor =
           output_buffer.get_access<cl::sycl::access::mode::write>(cgh);
-      auto input_accessor =
-          input_buffer.get_access<cl::sycl::access::mode::read>(cgh);
 
       cgh.single_task<class kernel>([=] {
-        auto input_ptr = input_accessor.get_pointer();
+        cl::sycl::host_ptr<int> input_ptr(in_ptr);
         auto output_ptr = output_accessor.get_pointer();
 
         using PrefetchingLSU =