Add private_memory::operator() support to xmethods

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 []>'