|
| 1 | +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
| 2 | +# See https://llvm.org/LICENSE.txt for license information. |
| 3 | +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
| 4 | + |
| 5 | +import gdb.xmethod |
| 6 | +import re |
| 7 | + |
| 8 | +""" |
| 9 | +Generalized base class for buffer index calculation |
| 10 | +""" |
| 11 | +class Accessor: |
| 12 | + def memory_range(self, dim): |
| 13 | + pass |
| 14 | + |
| 15 | + def offset(self, dim): |
| 16 | + pass |
| 17 | + |
| 18 | + def data(self): |
| 19 | + pass |
| 20 | + |
| 21 | + def __init__(self, obj, result_type, depth): |
| 22 | + self.obj = obj |
| 23 | + self.result_type = result_type |
| 24 | + self.depth = depth |
| 25 | + |
| 26 | + def index(self, arg): |
| 27 | + if arg.type.code == gdb.TYPE_CODE_INT: |
| 28 | + return int(arg) |
| 29 | + # https://github.com/intel/llvm/blob/97272b7ebd569bfa13811913a31e30f926559217/sycl/include/CL/sycl/accessor.hpp#L678-L690 |
| 30 | + result = 0 |
| 31 | + for dim in range(self.depth): |
| 32 | + result = result * self.memory_range(dim) + \ |
| 33 | + self.offset(dim) + \ |
| 34 | + arg['common_array'][dim] |
| 35 | + return result |
| 36 | + |
| 37 | + def value(self, arg): |
| 38 | + return self.data().cast(self.result_type.pointer())[self.index(arg)] |
| 39 | + |
| 40 | + |
| 41 | +""" |
| 42 | +For Host device memory layout |
| 43 | +""" |
| 44 | +class HostAccessor(Accessor): |
| 45 | + def payload(self): |
| 46 | + return self.obj['impl']['_M_ptr'].dereference() |
| 47 | + |
| 48 | + def memory_range(self, dim): |
| 49 | + return self.payload()['MMemoryRange']['common_array'][dim] |
| 50 | + |
| 51 | + def offset(self, dim): |
| 52 | + return self.payload()['MOffset']['common_array'][dim] |
| 53 | + |
| 54 | + def data(self): |
| 55 | + return self.payload()['MData'] |
| 56 | + |
| 57 | +""" |
| 58 | +For CPU/GPU memory layout |
| 59 | +""" |
| 60 | +class DeviceAccessor(Accessor): |
| 61 | + def memory_range(self, dim): |
| 62 | + return self.obj['impl']['MemRange']['common_array'][dim] |
| 63 | + |
| 64 | + def offset(self, dim): |
| 65 | + return self.obj['impl']['Offset']['common_array'][dim] |
| 66 | + |
| 67 | + def data(self): |
| 68 | + return self.obj['MData'] |
| 69 | + |
| 70 | + |
| 71 | +""" |
| 72 | +Generic implementation for N-dimensional ID |
| 73 | +""" |
| 74 | +class AccessorOpIndex(gdb.xmethod.XMethodWorker): |
| 75 | + def __init__(self, class_type, result_type, depth): |
| 76 | + self.class_type = class_type |
| 77 | + self.result_type = result_type |
| 78 | + self.depth = depth |
| 79 | + |
| 80 | + def get_arg_types(self): |
| 81 | + return gdb.lookup_type("cl::sycl::id<%s>" % self.depth) |
| 82 | + |
| 83 | + def get_result_type(self): |
| 84 | + return self.result_type |
| 85 | + |
| 86 | + def __call__(self, obj, arg): |
| 87 | + # No way to wasily figure out which devices is currently being used, |
| 88 | + # try all accessor implementations until one of them works: |
| 89 | + accessors = [ |
| 90 | + DeviceAccessor(obj, self.result_type, self.depth), |
| 91 | + HostAccessor(obj, self.result_type, self.depth) |
| 92 | + ] |
| 93 | + for accessor in accessors: |
| 94 | + try: |
| 95 | + return accessor.value(arg) |
| 96 | + except: |
| 97 | + pass |
| 98 | + |
| 99 | + print("Failed to call '%s.operator[](%s)" % (obj.type, arg.type)) |
| 100 | + |
| 101 | + return None |
| 102 | + |
| 103 | + |
| 104 | +""" |
| 105 | +Introduces an extra overload for 1D case that takes plain size_t |
| 106 | +""" |
| 107 | +class AccessorOpIndex1D(AccessorOpIndex): |
| 108 | + def get_arg_types(self): |
| 109 | + assert(self.depth == 1) |
| 110 | + return gdb.lookup_type('size_t') |
| 111 | + |
| 112 | + |
| 113 | +class AccessorOpIndexMatcher(gdb.xmethod.XMethodMatcher): |
| 114 | + def __init__(self): |
| 115 | + gdb.xmethod.XMethodMatcher.__init__(self, 'AccessorOpIndexMatcher') |
| 116 | + |
| 117 | + def match(self, class_type, method_name): |
| 118 | + if method_name != 'operator[]': |
| 119 | + return None |
| 120 | + |
| 121 | + result = re.match('^cl::sycl::accessor<.+>$', class_type.tag) |
| 122 | + if (result == None): |
| 123 | + return None |
| 124 | + |
| 125 | + depth = int(class_type.template_argument(1)) |
| 126 | + result_type = class_type.template_argument(0) |
| 127 | + |
| 128 | + methods = [ |
| 129 | + AccessorOpIndex(class_type, result_type, depth) |
| 130 | + ] |
| 131 | + if depth == 1: |
| 132 | + methods.append(AccessorOpIndex1D(class_type, result_type, depth)) |
| 133 | + return methods |
| 134 | + |
| 135 | + |
| 136 | +gdb.xmethod.register_xmethod_matcher(None, AccessorOpIndexMatcher()) |
0 commit comments