Add basic xmethod impl for accessor.operator[]

Intended to address templated function/operator debugging problem in
GDB. The added script is the currently used implementation in our gdb
- it only covers single-index `accessor.operator[]` with multiple
dimensions, that may be extended in the future depending on user
experience requirements.

It is much desirable to distribute this as part of compiler package
because it depends on private implementation details of SYCL headers and
must always be strictly in sync.

Signed-off-by: Mihails Strasuns <[email protected]>

Author: mihails-strasuns-intel

sycl/test/xmethods/accessors-device.cpp

@@ -0,0 +1,18 @@
+// RUN: %clangxx -fsycl-device-only -c -fno-color-diagnostics -Xclang -ast-dump %s | FileCheck %s
+#include <CL/sycl/accessor.hpp>
+
+typedef cl::sycl::accessor<int, 1, cl::sycl::access::mode::read> dummy;
+
+// AccessorImplDevice must have MemRange and Offset fields
+
+// CHECK: CXXRecordDecl {{.*}} class AccessorImplDevice definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced Offset
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced MemRange
+
+// accessor.impl must be present and of AccessorImplDevice type
+
+// CHECK: CXXRecordDecl {{.*}} class accessor definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced impl 'detail::AccessorImplDevice<AdjustedDim>'

sycl/test/xmethods/accessors.cpp

@@ -0,0 +1,24 @@
+// RUN: %clangxx -c -fno-color-diagnostics -Xclang -ast-dump %s | FileCheck %s
+#include <CL/sycl/accessor.hpp>
+
+typedef cl::sycl::accessor<int, 1, cl::sycl::access::mode::read> dummy;
+
+// AccessorImplHost must have MMemoryRange, MOffset and MData fields
+
+// CHECK: CXXRecordDecl {{.*}} class AccessorImplHost definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced MOffset
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced MMemoryRange
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced MData
+
+// accessor.impl must be present and of shared_ptr<AccessorImplHost> type
+
+// CHECK: CXXRecordDecl {{.*}} class AccessorBaseHost definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: FieldDecl {{.*}} referenced impl {{.*}}:'std::shared_ptr<cl::sycl::detail::AccessorImplHost>'
+// CHECK: CXXRecordDecl {{.*}} class accessor definition
+// CHECK-NOT: CXXRecordDecl {{.*}} definition
+// CHECK: public {{.*}}:'cl::sycl::detail::AccessorBaseHost'
+

sycl/xmethods/accessors.py

@@ -0,0 +1,136 @@
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+import gdb.xmethod
+import re
+
+"""
+Generalized base class for buffer index calculation
+"""
+class Accessor:
+    def memory_range(self, dim):
+        pass
+
+    def offset(self, dim):
+        pass
+
+    def data(self):
+        pass
+
+    def __init__(self, obj, result_type, depth):
+        self.obj = obj
+        self.result_type = result_type
+        self.depth = depth
+
+    def index(self, arg):
+        if arg.type.code == gdb.TYPE_CODE_INT:
+            return int(arg)
+        # https://github.com/intel/llvm/blob/97272b7ebd569bfa13811913a31e30f926559217/sycl/include/CL/sycl/accessor.hpp#L678-L690
+        result = 0
+        for dim in range(self.depth):
+            result = result * self.memory_range(dim) + \
+                self.offset(dim) + \
+                arg['common_array'][dim]
+        return result
+
+    def value(self, arg):
+        return self.data().cast(self.result_type.pointer())[self.index(arg)]
+
+
+"""
+For Host device memory layout
+"""
+class HostAccessor(Accessor):
+    def payload(self):
+        return self.obj['impl']['_M_ptr'].dereference()
+
+    def memory_range(self, dim):
+        return self.payload()['MMemoryRange']['common_array'][dim]
+
+    def offset(self, dim):
+        return self.payload()['MOffset']['common_array'][dim]
+
+    def data(self):
+        return self.payload()['MData']
+
+"""
+For CPU/GPU memory layout
+"""
+class DeviceAccessor(Accessor):
+    def memory_range(self, dim):
+        return self.obj['impl']['MemRange']['common_array'][dim]
+
+    def offset(self, dim):
+        return self.obj['impl']['Offset']['common_array'][dim]
+
+    def data(self):
+        return self.obj['MData']
+
+
+"""
+Generic implementation for N-dimensional ID
+"""
+class AccessorOpIndex(gdb.xmethod.XMethodWorker):
+    def __init__(self, class_type, result_type, depth):
+        self.class_type = class_type
+        self.result_type = result_type
+        self.depth = depth
+
+    def get_arg_types(self):
+        return gdb.lookup_type("cl::sycl::id<%s>" % self.depth)
+
+    def get_result_type(self):
+        return self.result_type
+
+    def __call__(self, obj, arg):
+        # No way to wasily figure out which devices is currently being used,
+        # try all accessor implementations until one of them works:
+        accessors = [
+            DeviceAccessor(obj, self.result_type, self.depth),
+            HostAccessor(obj, self.result_type, self.depth)
+        ]
+        for accessor in accessors:
+            try:
+                return accessor.value(arg)
+            except:
+                pass
+
+        print("Failed to call '%s.operator[](%s)" % (obj.type, arg.type))
+
+        return None
+
+
+"""
+Introduces an extra overload for 1D case that takes plain size_t
+"""
+class AccessorOpIndex1D(AccessorOpIndex):
+    def get_arg_types(self):
+        assert(self.depth == 1)
+        return gdb.lookup_type('size_t')
+
+
+class AccessorOpIndexMatcher(gdb.xmethod.XMethodMatcher):
+    def __init__(self):
+        gdb.xmethod.XMethodMatcher.__init__(self, 'AccessorOpIndexMatcher')
+
+    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):
+            return None
+
+        depth = int(class_type.template_argument(1))
+        result_type = class_type.template_argument(0)
+
+        methods = [
+            AccessorOpIndex(class_type, result_type, depth)
+        ]
+        if depth == 1:
+            methods.append(AccessorOpIndex1D(class_type, result_type, depth))
+        return methods
+
+
+gdb.xmethod.register_xmethod_matcher(None, AccessorOpIndexMatcher())