[SYCL] Add private_memory::operator() support to xmethods (#2657)

Allows to inspect values wrapped into private_memory<T> by evaluating
`operator()` from GDB. Necessary for GPU devices where inferior calls
are not supported, however Host implementation is also provided for
completeness.

Also slightly tweaks existing class naming scheme in the GDB python
script.

Author: mihails-strasuns-intel

sycl/gdb/libsycl.so-gdb.py

@@ -2,18 +2,16 @@
 # See https://llvm.org/LICENSE.txt for license information.
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+import re
 import gdb
 import gdb.xmethod
 import gdb.printing
-import itertools
-import re
 
 ### XMethod implementations ###
 
-"""
-Generalized base class for buffer index calculation
-"""
 class Accessor:
+    """Generalized base class for buffer index calculation"""
+
     def memory_range(self, dim):
         pass
 
@@ -43,10 +41,9 @@ def value(self, arg):
         return self.data().cast(self.result_type.pointer())[self.index(arg)]
 
 
-"""
-For Host device memory layout
-"""
 class HostAccessor(Accessor):
+    """For Host device memory layout"""
+
     def payload(self):
         return self.obj['impl']['_M_ptr'].dereference()
 
@@ -59,10 +56,9 @@ def offset(self, dim):
     def data(self):
         return self.payload()['MData']
 
-"""
-For CPU/GPU memory layout
-"""
 class DeviceAccessor(Accessor):
+    """For CPU/GPU memory layout"""
+
     def memory_range(self, dim):
         return self.obj['impl']['MemRange']['common_array'][dim]
 
@@ -73,10 +69,9 @@ def data(self):
         return self.obj['MData']
 
 
-"""
-Generic implementation for N-dimensional ID
-"""
 class AccessorOpIndex(gdb.xmethod.XMethodWorker):
+    """Generic implementation for N-dimensional ID"""
+
     def __init__(self, class_type, result_type, depth):
         self.class_type = class_type
         self.result_type = result_type
@@ -106,25 +101,25 @@ def __call__(self, obj, arg):
         return None
 
 
-"""
-Introduces an extra overload for 1D case that takes plain size_t
-"""
 class AccessorOpIndex1D(AccessorOpIndex):
+    """Introduces an extra overload for 1D case that takes plain size_t"""
+
     def get_arg_types(self):
-        assert(self.depth == 1)
+        assert self.depth == 1
         return gdb.lookup_type('size_t')
 
 
-class AccessorOpIndexMatcher(gdb.xmethod.XMethodMatcher):
+class AccessorMatcher(gdb.xmethod.XMethodMatcher):
+    """Entry point for cl::sycl::accessor"""
     def __init__(self):
-        gdb.xmethod.XMethodMatcher.__init__(self, 'AccessorOpIndexMatcher')
+        gdb.xmethod.XMethodMatcher.__init__(self, 'AccessorMatcher')
 
     def match(self, class_type, method_name):
         if method_name != 'operator[]':
             return None
 
         result = re.match('^cl::sycl::accessor<.+>$', class_type.tag)
-        if (result == None):
+        if result is None:
             return None
 
         depth = int(class_type.template_argument(1))
@@ -137,15 +132,86 @@ def match(self, class_type, method_name):
             methods.append(AccessorOpIndex1D(class_type, result_type, depth))
         return methods
 
+class PrivateMemoryOpCall(gdb.xmethod.XMethodWorker):
+    """Provides operator() overload for h_item argument"""
+
+    class ItemBase:
+        """Wrapper for cl::sycl::detail::ItemBase which reimplements index calculation"""
+
+        def __init__(self, obj, ):
+            result = re.match('^cl::sycl::detail::ItemBase<(.+), (.+)>$', str(obj.type))
+            self.dim = int(result[1])
+            self.with_offset = (result[2] == 'true')
+            self.obj = obj
+
+        def get_linear_id(self):
+            index  = self.obj['MIndex']['common_array']
+            extent = self.obj['MExtent']['common_array']
+
+            if self.with_offset:
+                offset = self.obj['MOffset']['common_array']
+                if self.dim == 1:
+                    return index[0] - offset[0]
+                elif self.dim == 2:
+                    return (index[0] - offset[0]) * extent[1] + (index[1] - offset[1])
+                else:
+                    return ((index[0] - offset[0]) * extent[1] * extent[2]) + \
+                           ((index[1] - offset[1]) * extent[2]) + (index[2] - offset[2])
+            else:
+                if self.dim == 1:
+                    return index[0]
+                elif self.dim == 2:
+                    return index[0] * extent[1] + index[1]
+                else:
+                    return (index[0] * extent[1] * extent[2]) + (index[1] * extent[2]) + index[2]
+
+    def __init__(self, result_type, dim):
+        self.result_type = result_type
+        self.dim = dim
+
+    def get_arg_types(self):
+        return gdb.lookup_type("cl::sycl::h_item<%s>" % self.dim)
+
+    def get_result_type(self, *args):
+        return self.result_type
+
+    def __call__(self, obj, *args):
+        if obj['Val'].type.tag == self.result_type:
+            # On device private_memory is a simple wrapper over actual value
+            return obj['Val']
+        else:
+            # On host it wraps a unique_ptr to an array of items
+            item_base = args[0]['localItem']['MImpl']
+            item_base = self.ItemBase(item_base)
+            index = item_base.get_linear_id()
+            return obj['Val']['_M_t']['_M_t']['_M_head_impl'][index]
 
-gdb.xmethod.register_xmethod_matcher(None, AccessorOpIndexMatcher(), replace=True)
+class PrivateMemoryMatcher(gdb.xmethod.XMethodMatcher):
+    """Entry point for cl::sycl::private_memory"""
+
+    def __init__(self):
+        gdb.xmethod.XMethodMatcher.__init__(self, 'PrivateMemoryMatcher')
+
+    def match(self, class_type, method_name):
+        if method_name != 'operator()':
+            return None
+
+        result = re.match('^cl::sycl::private_memory<(cl::sycl::id<.+>), (.+)>$', class_type.tag)
+        if result is None:
+            return None
+
+        return PrivateMemoryOpCall(result[1], result[2])
+
+
+
+gdb.xmethod.register_xmethod_matcher(None, AccessorMatcher(), replace=True)
+gdb.xmethod.register_xmethod_matcher(None, PrivateMemoryMatcher(), replace=True)
 
 ### Pretty-printer implementations ###
 
-"""
-Print an object deriving from cl::sycl::detail::array
-"""
 class SyclArrayPrinter:
+    """Print an object deriving from cl::sycl::detail::array"""
+
     class ElementIterator:
         def __init__(self, data, size):
             self.data = data
@@ -168,7 +234,7 @@ def __next__(self):
 
     def __init__(self, value):
         if value.type.code == gdb.TYPE_CODE_REF:
-            if hasattr(gdb.Value,"referenced_value"):
+            if hasattr(gdb.Value, "referenced_value"):
                 value = value.referenced_value()
 
         self.value = value
@@ -182,7 +248,7 @@ def children(self):
             # 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 [ ]
+            return []
 
     def to_string(self):
         try:
@@ -197,10 +263,9 @@ def to_string(self):
     def display_hint(self):
         return 'array'
 
-"""
-Print a cl::sycl::buffer
-"""
 class SyclBufferPrinter:
+    """Print a cl::sycl::buffer"""
+
     def __init__(self, value):
         self.value = value
         self.type = value.type.unqualified().strip_typedefs()
@@ -217,8 +282,7 @@ def to_string(self):
                    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::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)
-

sycl/test/gdb/private-memory-device.cpp

@@ -0,0 +1,12 @@
+// RUN: %clangxx -fsycl-device-only -c -fno-color-diagnostics -Xclang -ast-dump %s -I %sycl_include -Wno-sycl-strict | FileCheck %s
+// UNSUPPORTED: windows
+#include <CL/sycl/group.hpp>
+#include <CL/sycl/id.hpp>
+
+typedef cl::sycl::private_memory<cl::sycl::id<1>, 1> dummy;
+
+// private_memory must have Val field of T type
+
+// CHECK: CXXRecordDecl {{.*}} class private_memory definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced Val 'T'

sycl/test/gdb/private-memory.cpp

@@ -0,0 +1,11 @@
+// RUN: %clangxx -c -fsycl -fno-color-diagnostics -Xclang -ast-dump %s | FileCheck %s
+// UNSUPPORTED: windows
+#include <CL/sycl/group.hpp>
+#include <CL/sycl/id.hpp>
+
+typedef cl::sycl::private_memory<cl::sycl::id<1>, 1> dummy;
+
+// private_memory must have Val field of unique_ptr<T [], ...> type
+
+// CHECK: CXXRecordDecl {{.*}} class private_memory definition
+// CHECK: FieldDecl {{.*}} referenced Val {{.*}}:'unique_ptr<T []>'