diff --git a/sycl/gdb/libsycl.so-gdb.py b/sycl/gdb/libsycl.so-gdb.py index 46c79f2b9c661..b98e2c80949bf 100644 --- a/sycl/gdb/libsycl.so-gdb.py +++ b/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) - diff --git a/sycl/test/gdb/private-memory-device.cpp b/sycl/test/gdb/private-memory-device.cpp new file mode 100644 index 0000000000000..e02d0f140f05c --- /dev/null +++ b/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 +#include + +typedef cl::sycl::private_memory, 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' diff --git a/sycl/test/gdb/private-memory.cpp b/sycl/test/gdb/private-memory.cpp new file mode 100644 index 0000000000000..8016a3d310871 --- /dev/null +++ b/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 +#include + +typedef cl::sycl::private_memory, 1> dummy; + +// private_memory must have Val field of unique_ptr type + +// CHECK: CXXRecordDecl {{.*}} class private_memory definition +// CHECK: FieldDecl {{.*}} referenced Val {{.*}}:'unique_ptr'