Merge remote-tracking branch 'upstream/sycl' into glyons-name

Author: glyons-intel

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/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)
+