Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/lib/CodeGen/CGCall.cpp

@@ -4971,10 +4971,21 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
 
   // 4. Finish the call.
 
+  // SYCL does not support C++ exceptions or termination in device code, so all
+  // functions have to return.
+  bool SyclSkipNoReturn = false;
+  if (getLangOpts().SYCLIsDevice && CI->doesNotReturn()) {
+    if (auto *F = CI->getCalledFunction())
+      F->removeFnAttr(llvm::Attribute::NoReturn);
+    CI->removeAttribute(llvm::AttributeList::FunctionIndex,
+                        llvm::Attribute::NoReturn);
+    SyclSkipNoReturn = true;
+  }
+
   // If the call doesn't return for non-sycl devices, finish the basic block and
   // clear the insertion point; this allows the rest of IRGen to discard
   // unreachable code.
-  if (CI->doesNotReturn() && !getLangOpts().SYCLIsDevice) {
+  if (!SyclSkipNoReturn && CI->doesNotReturn()) {
     if (UnusedReturnSizePtr)
       PopCleanupBlock();
 

clang/test/CodeGenSYCL/remove-ur-inst.cpp

@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-enable-optimizations -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
 SYCL_EXTERNAL void doesNotReturn() throw() __attribute__((__noreturn__));
 
@@ -11,6 +12,7 @@ int main() {
   kernel<class test>([]() {
     doesNotReturn();
     // CHECK-NOT: unreachable
+    // CHECK-NOT: noreturn
   });
   return 0;
-}
+}

sycl/CMakeLists.txt

@@ -14,7 +14,7 @@ include(AddSYCLExecutable)
 set(SYCL_MAJOR_VERSION 2)
 set(SYCL_MINOR_VERSION 1)
 set(SYCL_PATCH_VERSION 0)
-set(SYCL_DEV_ABI_VERSION 2)
+set(SYCL_DEV_ABI_VERSION 3)
 if (SYCL_ADD_DEV_VERSION_POSTFIX)
   set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}")
 endif()
@@ -248,7 +248,7 @@ add_subdirectory( source )
 # Auxilliary extras for SYCL headers/library
 if (NOT WIN32)
   install(FILES
-      "${CMAKE_CURRENT_SOURCE_DIR}/xmethods/libsycl.so-gdb.py"
+      "${CMAKE_CURRENT_SOURCE_DIR}/gdb/libsycl.so-gdb.py"
       RENAME "libsycl.so.${SYCL_VERSION_STRING}-gdb.py"
       DESTINATION "lib${LLVM_LIBDIR_SUFFIX}/"
       COMPONENT sycl-headers-extras)

sycl/doc/GetStartedGuide.md

@@ -27,7 +27,7 @@ and a wide range of compute accelerators such as GPU and FPGA.
 * `ninja` -
 [Download](https://github.com/ninja-build/ninja/wiki/Pre-built-Ninja-packages)
 * C++ compiler
-  * Linux: `GCC` version 5.1.0 or later (including libstdc++) -
+  * Linux: `GCC` version 7.1.0 or later (including libstdc++) -
     [Download](https://gcc.gnu.org/install/)
   * Windows: `Visual Studio` version 15.7 preview 4 or later -
     [Download](https://visualstudio.microsoft.com/downloads/)

sycl/doc/extensions/IntelFPGA/FPGALsu.md

@@ -0,0 +1,116 @@
+
+# FPGA lsu
+
+The Intel FPGA `lsu` class is implemented in `CL/sycl/intel/fpga_lsu.hpp` which
+is included in `CL/sycl/intel/fpga_extensions.hpp`.
+
+The class `cl::sycl::intel::lsu` allows users to explicitly request that the
+implementation of a global memory access is configured in a certain way. The
+class has two member functions, `load()` and `store()` which allow loading from
+and storing to a `global_ptr`, respectively, and is templated on the following
+4 optional paremeters:
+
+1.  **`cl::sycl::intel::burst_coalesce<B>`, where `B` is a boolean**: request,
+to the extent possible, that a dynamic burst coalescer be implemented when
+`load` or `store` are called. The default value of this parameter is `false`.
+2. **`cl::sycl::intel::cache<N>`, where `N` is an integer greater or equal to
+0**: request, to the extent possible, that a read-only cache of the specified
+size in bytes be implemented when when `load` is called. It is not allowed to
+use that parameter for `store`. The default value of this parameter is `0`.
+3. **`cl::sycl::intel::statically_coalesce<N>`, where `B` is a boolean**:
+request, to the extent possible, that `load` or `store` accesses, is allowed to
+be statically coalesced with other memory accesses at compile time. The default
+value of this parameter is `true`.
+4. **`cl::sycl::intel::prefetch<B>`, where `N` is a boolean**: request, to the
+extent possible, that a prefetcher be implemented when `load` is called. It is
+not allowed to use that parameter for `store`. The default value of this
+parameter is `false`.
+
+Currently, not every combination of parameters is allowed due to limitations in
+the backend. The following rules apply:
+1. For `store`, `cl::sycl::intel::cache` must be `0` and
+`cl::sycl::intel::prefetch` must be `false`.
+2. For `load`, if `cl::sycl::intel::cache` is set to a value greater than `0`,
+then `cl::sycl::intel::burst_coalesce` must be set to `true`.
+3. For `load`, exactly one of `cl::sycl::intel::prefetch` and
+`cl::sycl::intel::burst_coalesce` is allowed to be `true`.
+4. For `load`, exactly one of `cl::sycl::intel::prefetch` and
+`cl::sycl::intel::cache` is allowed to be `true`.
+
+## Implementation
+
+The implementation relies on the Clang built-in `__builtin_intel_fpga_mem` when
+parsing the SYCL device code. The built-in uses the LLVM `ptr.annotation`
+intrinsic under the hood to annotate the pointer that is being accessed.
+```c++
+template <class... mem_access_params> class lsu final {
+public:
+  lsu() = delete;
+
+  template <typename T> static T &load(sycl::global_ptr<T> Ptr) {
+    check_load();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    return *__builtin_intel_fpga_mem((T *)Ptr,
+                                     _burst_coalesce | _cache |
+                                         _dont_statically_coalesce | _prefetch,
+                                     _cache_val);
+#else
+    return *Ptr;
+#endif
+  }
+
+  template <typename T> static void store(sycl::global_ptr<T> Ptr, T Val) {
+    check_store();
+#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem)
+    *__builtin_intel_fpga_mem((T *)Ptr,
+                              _burst_coalesce | _cache |
+                                  _dont_statically_coalesce | _prefetch,
+                              _cache_val) = Val;
+#else
+    *Ptr = Val;
+#endif
+  }
+  ...
+}
+```
+
+## Usage
+
+```c++
+#include <CL/sycl/intel/fpga_extensions.hpp>
+...
+cl::sycl::buffer<int, 1> output_buffer(output_data, 1);
+cl::sycl::buffer<int, 1> input_buffer(input_data, 1);
+
+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();
+    auto output_ptr = output_accessor.get_pointer();
+
+    using PrefetchingLSU =
+        cl::sycl::intel::lsu<cl::sycl::intel::prefetch<true>,
+                             cl::sycl::intel::statically_coalesce<false>>;
+
+    using BurstCoalescedLSU =
+        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<false>,
+                             cl::sycl::intel::statically_coalesce<false>>;
+
+    using CachingLSU =
+        cl::sycl::intel::lsu<cl::sycl::intel::burst_coalesce<true>,
+                             cl::sycl::intel::cache<1024>,
+                             cl::sycl::intel::statically_coalesce<true>>;
+
+    using PipelinedLSU = cl::sycl::intel::lsu<>;
+
+    int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0]
+    int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1]
+
+    BurstCoalescedLSU::store(output_ptr, X); // output_ptr[0] = X
+    PipelinedLSU::store(output_ptr + 1, Y);  // output_ptr[1] = Y
+  });
+});
+...
+```

sycl/doc/extensions/SubGroup/SYCL_INTEL_sub_group.asciidoc

@@ -68,7 +68,7 @@ Providing a generic group abstraction encapsulating the shared functionality of
 
 === Attributes
 
-The +[[intel::reqd_sub_group_size(n)]]+ attribute indicates that the kernel must be compiled and executed with a sub-group of size _n_.  The value of _n_ must be set to a sub-group size that is both supported by the device and compatible with all language features used by the kernel, or device compilation will fail.  The set of valid sub-group sizes can be queried as described below.
+The +[[intel::reqd_sub_group_size(n)]]+ attribute indicates that the kernel must be compiled and executed with a sub-group of size _n_.  The value of _n_ must be a compile-time integral constant expression.  The value of _n_ must be set to a sub-group size that is both supported by the device and compatible with all language features used by the kernel, or device compilation will fail.  The set of valid sub-group sizes can be queried as described below.
 
 In addition to device functions, the required sub-group size attribute may also be specified in the definition of a named functor object, as in the example below:
 
@@ -314,6 +314,7 @@ Yes, the four shuffles in this extension are a defining feature of sub-groups.
 |4|2020-04-21|John Pennycook|*Restore missing barrier function*
 |5|2020-04-21|John Pennycook|*Restore sub-group shuffles as member functions*
 |6|2020-04-22|John Pennycook|*Align with SYCL_INTEL_device_specific_kernel_queries*
+|7|2020-07-13|John Pennycook|*Clarify that reqd_sub_group_size must be a compile-time constant*
 |========================================
 
 //************************************************************************

sycl/xmethods/libsycl.so-gdb.py renamed to sycl/gdb/libsycl.so-gdb.py

@@ -2,9 +2,14 @@
 # See https://llvm.org/LICENSE.txt for license information.
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+import gdb
 import gdb.xmethod
+import gdb.printing
+import itertools
 import re
 
+### XMethod implementations ###
+
 """
 Generalized base class for buffer index calculation
 """
@@ -134,3 +139,86 @@ def match(self, class_type, method_name):
 
 
 gdb.xmethod.register_xmethod_matcher(None, AccessorOpIndexMatcher(), replace=True)
+
+### Pretty-printer implementations ###
+
+"""
+Print an object deriving from cl::sycl::detail::array
+"""
+class SyclArrayPrinter:
+    class ElementIterator:
+        def __init__(self, data, size):
+            self.data = data
+            self.size = size
+            self.count = 0
+
+        def __iter__(self):
+            return self
+
+        def __next__(self):
+            if self.count == self.size:
+                raise StopIteration
+            count = self.count
+            self.count = self.count + 1
+            try:
+                elt = self.data[count]
+            except:
+                elt = "<error reading variable>"
+            return ('[%d]' % count, elt)
+
+    def __init__(self, value):
+        if value.type.code == gdb.TYPE_CODE_REF:
+            if hasattr(gdb.Value,"referenced_value"):
+                value = value.referenced_value()
+
+        self.value = value
+        self.type = value.type.unqualified().strip_typedefs()
+        self.dimensions = self.type.template_argument(0)
+
+    def children(self):
+        try:
+            return self.ElementIterator(self.value['common_array'], self.dimensions)
+        except:
+            # There is no way to return an error from this method. Return an
+            # empty iterable to make GDB happy and rely on to_string method
+            # to take care of formatting.
+            return [ ]
+
+    def to_string(self):
+        try:
+            # Check if accessing array value will succeed and resort to
+            # error message otherwise. Individual array element access failures
+            # will be caught by iterator itself.
+            _ = self.value['common_array']
+            return self.type.tag
+        except:
+            return "<error reading variable>"
+
+    def display_hint(self):
+        return 'array'
+
+"""
+Print a cl::sycl::buffer
+"""
+class SyclBufferPrinter:
+    def __init__(self, value):
+        self.value = value
+        self.type = value.type.unqualified().strip_typedefs()
+        self.elt_type = value.type.template_argument(0)
+        self.dimensions = value.type.template_argument(1)
+        self.typeregex = re.compile('^([a-zA-Z0-9_:]+)(<.*>)?$')
+
+    def to_string(self):
+        match = self.typeregex.match(self.type.tag)
+        if not match:
+            return "<error parsing type>"
+        return ('%s<%s, %s> = {impl=%s}'
+                % (match.group(1), self.elt_type, self.dimensions,
+                   self.value['impl'].address))
+
+sycl_printer = gdb.printing.RegexpCollectionPrettyPrinter("SYCL")
+sycl_printer.add_printer("cl::sycl::id",     '^cl::sycl::id<.*$',     SyclArrayPrinter)
+sycl_printer.add_printer("cl::sycl::range",  '^cl::sycl::range<.*$',  SyclArrayPrinter)
+sycl_printer.add_printer("cl::sycl::buffer", '^cl::sycl::buffer<.*$', SyclBufferPrinter)
+gdb.printing.register_pretty_printer(None, sycl_printer, True)
+