22# See https://llvm.org/LICENSE.txt for license information.
33# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
44
5+ import re
56import gdb
67import gdb .xmethod
78import gdb .printing
8- import itertools
9- import re
109
1110### XMethod implementations ###
1211
13- """
14- Generalized base class for buffer index calculation
15- """
1612class Accessor :
13+ """Generalized base class for buffer index calculation"""
14+
1715 def memory_range (self , dim ):
1816 pass
1917
@@ -43,10 +41,9 @@ def value(self, arg):
4341 return self .data ().cast (self .result_type .pointer ())[self .index (arg )]
4442
4543
46- """
47- For Host device memory layout
48- """
4944class HostAccessor (Accessor ):
45+ """For Host device memory layout"""
46+
5047 def payload (self ):
5148 return self .obj ['impl' ]['_M_ptr' ].dereference ()
5249
@@ -59,10 +56,9 @@ def offset(self, dim):
5956 def data (self ):
6057 return self .payload ()['MData' ]
6158
62- """
63- For CPU/GPU memory layout
64- """
6559class DeviceAccessor (Accessor ):
60+ """For CPU/GPU memory layout"""
61+
6662 def memory_range (self , dim ):
6763 return self .obj ['impl' ]['MemRange' ]['common_array' ][dim ]
6864
@@ -73,10 +69,9 @@ def data(self):
7369 return self .obj ['MData' ]
7470
7571
76- """
77- Generic implementation for N-dimensional ID
78- """
7972class AccessorOpIndex (gdb .xmethod .XMethodWorker ):
73+ """Generic implementation for N-dimensional ID"""
74+
8075 def __init__ (self , class_type , result_type , depth ):
8176 self .class_type = class_type
8277 self .result_type = result_type
@@ -106,25 +101,25 @@ def __call__(self, obj, arg):
106101 return None
107102
108103
109- """
110- Introduces an extra overload for 1D case that takes plain size_t
111- """
112104class AccessorOpIndex1D (AccessorOpIndex ):
105+ """Introduces an extra overload for 1D case that takes plain size_t"""
106+
113107 def get_arg_types (self ):
114- assert ( self .depth == 1 )
108+ assert self .depth == 1
115109 return gdb .lookup_type ('size_t' )
116110
117111
118- class AccessorOpIndexMatcher (gdb .xmethod .XMethodMatcher ):
112+ class AccessorMatcher (gdb .xmethod .XMethodMatcher ):
113+ """Entry point for cl::sycl::accessor"""
119114 def __init__ (self ):
120- gdb .xmethod .XMethodMatcher .__init__ (self , 'AccessorOpIndexMatcher ' )
115+ gdb .xmethod .XMethodMatcher .__init__ (self , 'AccessorMatcher ' )
121116
122117 def match (self , class_type , method_name ):
123118 if method_name != 'operator[]' :
124119 return None
125120
126121 result = re .match ('^cl::sycl::accessor<.+>$' , class_type .tag )
127- if ( result == None ) :
122+ if result is None :
128123 return None
129124
130125 depth = int (class_type .template_argument (1 ))
@@ -137,15 +132,86 @@ def match(self, class_type, method_name):
137132 methods .append (AccessorOpIndex1D (class_type , result_type , depth ))
138133 return methods
139134
135+ class PrivateMemoryOpCall (gdb .xmethod .XMethodWorker ):
136+ """Provides operator() overload for h_item argument"""
137+
138+ class ItemBase :
139+ """Wrapper for cl::sycl::detail::ItemBase which reimplements index calculation"""
140+
141+ def __init__ (self , obj , ):
142+ result = re .match ('^cl::sycl::detail::ItemBase<(.+), (.+)>$' , str (obj .type ))
143+ self .dim = int (result [1 ])
144+ self .with_offset = (result [2 ] == 'true' )
145+ self .obj = obj
146+
147+ def get_linear_id (self ):
148+ index = self .obj ['MIndex' ]['common_array' ]
149+ extent = self .obj ['MExtent' ]['common_array' ]
150+
151+ if self .with_offset :
152+ offset = self .obj ['MOffset' ]['common_array' ]
153+ if self .dim == 1 :
154+ return index [0 ] - offset [0 ]
155+ elif self .dim == 2 :
156+ return (index [0 ] - offset [0 ]) * extent [1 ] + (index [1 ] - offset [1 ])
157+ else :
158+ return ((index [0 ] - offset [0 ]) * extent [1 ] * extent [2 ]) + \
159+ ((index [1 ] - offset [1 ]) * extent [2 ]) + (index [2 ] - offset [2 ])
160+ else :
161+ if self .dim == 1 :
162+ return index [0 ]
163+ elif self .dim == 2 :
164+ return index [0 ] * extent [1 ] + index [1 ]
165+ else :
166+ return (index [0 ] * extent [1 ] * extent [2 ]) + (index [1 ] * extent [2 ]) + index [2 ]
167+
168+ def __init__ (self , result_type , dim ):
169+ self .result_type = result_type
170+ self .dim = dim
171+
172+ def get_arg_types (self ):
173+ return gdb .lookup_type ("cl::sycl::h_item<%s>" % self .dim )
174+
175+ def get_result_type (self , * args ):
176+ return self .result_type
177+
178+ def __call__ (self , obj , * args ):
179+ if obj ['Val' ].type .tag == self .result_type :
180+ # On device private_memory is a simple wrapper over actual value
181+ return obj ['Val' ]
182+ else :
183+ # On host it wraps a unique_ptr to an array of items
184+ item_base = args [0 ]['localItem' ]['MImpl' ]
185+ item_base = self .ItemBase (item_base )
186+ index = item_base .get_linear_id ()
187+ return obj ['Val' ]['_M_t' ]['_M_t' ]['_M_head_impl' ][index ]
140188
141- gdb .xmethod .register_xmethod_matcher (None , AccessorOpIndexMatcher (), replace = True )
189+ class PrivateMemoryMatcher (gdb .xmethod .XMethodMatcher ):
190+ """Entry point for cl::sycl::private_memory"""
191+
192+ def __init__ (self ):
193+ gdb .xmethod .XMethodMatcher .__init__ (self , 'PrivateMemoryMatcher' )
194+
195+ def match (self , class_type , method_name ):
196+ if method_name != 'operator()' :
197+ return None
198+
199+ result = re .match ('^cl::sycl::private_memory<(cl::sycl::id<.+>), (.+)>$' , class_type .tag )
200+ if result is None :
201+ return None
202+
203+ return PrivateMemoryOpCall (result [1 ], result [2 ])
204+
205+
206+
207+ gdb .xmethod .register_xmethod_matcher (None , AccessorMatcher (), replace = True )
208+ gdb .xmethod .register_xmethod_matcher (None , PrivateMemoryMatcher (), replace = True )
142209
143210### Pretty-printer implementations ###
144211
145- """
146- Print an object deriving from cl::sycl::detail::array
147- """
148212class SyclArrayPrinter :
213+ """Print an object deriving from cl::sycl::detail::array"""
214+
149215 class ElementIterator :
150216 def __init__ (self , data , size ):
151217 self .data = data
@@ -168,7 +234,7 @@ def __next__(self):
168234
169235 def __init__ (self , value ):
170236 if value .type .code == gdb .TYPE_CODE_REF :
171- if hasattr (gdb .Value ,"referenced_value" ):
237+ if hasattr (gdb .Value , "referenced_value" ):
172238 value = value .referenced_value ()
173239
174240 self .value = value
@@ -182,7 +248,7 @@ def children(self):
182248 # There is no way to return an error from this method. Return an
183249 # empty iterable to make GDB happy and rely on to_string method
184250 # to take care of formatting.
185- return [ ]
251+ return []
186252
187253 def to_string (self ):
188254 try :
@@ -197,10 +263,9 @@ def to_string(self):
197263 def display_hint (self ):
198264 return 'array'
199265
200- """
201- Print a cl::sycl::buffer
202- """
203266class SyclBufferPrinter :
267+ """Print a cl::sycl::buffer"""
268+
204269 def __init__ (self , value ):
205270 self .value = value
206271 self .type = value .type .unqualified ().strip_typedefs ()
@@ -217,8 +282,7 @@ def to_string(self):
217282 self .value ['impl' ].address ))
218283
219284sycl_printer = gdb .printing .RegexpCollectionPrettyPrinter ("SYCL" )
220- sycl_printer .add_printer ("cl::sycl::id" , '^cl::sycl::id<.*$' , SyclArrayPrinter )
221- sycl_printer .add_printer ("cl::sycl::range" , '^cl::sycl::range<.*$' , SyclArrayPrinter )
285+ sycl_printer .add_printer ("cl::sycl::id" , '^cl::sycl::id<.*$' , SyclArrayPrinter )
286+ sycl_printer .add_printer ("cl::sycl::range" , '^cl::sycl::range<.*$' , SyclArrayPrinter )
222287sycl_printer .add_printer ("cl::sycl::buffer" , '^cl::sycl::buffer<.*$' , SyclBufferPrinter )
223288gdb .printing .register_pretty_printer (None , sycl_printer , True )
224-
0 commit comments