Add basic xmethod impl for accessor.operator[]

sycl/CMakeLists.txt

@@ -51,6 +51,7 @@ endif()
 set(CLANG_VERSION "${CLANG_VERSION_MAJOR}.${CLANG_VERSION_MINOR}.${CLANG_VERSION_PATCHLEVEL}")
 
 set(LLVM_INST_INC_DIRECTORY "lib${LLVM_LIBDIR_SUFFIX}/clang/${CLANG_VERSION}/include")
+set(LLVM_PYTHON_DIRECTORY "lib${LLVM_LIBDIR_SUFFIX}/clang/${CLANG_VERSION}/python")
 set(dst_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include)
 
 # Find OpenCL headers and libraries installed in the system and use them to
@@ -150,6 +151,12 @@ COMMENT "Copying SYCL headers ...")
 # Configure SYCL headers
 install(DIRECTORY "${sycl_inc_dir}/." DESTINATION "${LLVM_INST_INC_DIRECTORY}" COMPONENT sycl-headers)
 
+# Auxilliary extras for SYCL headers
+add_custom_target(sycl-headers-extras ALL
+COMMAND ${CMAKE_COMMAND} -E copy_directory ${CMAKE_CURRENT_SOURCE_DIR}/xmethods ${dst_dir}
+COMMENT "Copying SYCL header xmethod scripts ...")
+install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/xmethods" DESTINATION ${LLVM_PYTHON_DIRECTORY} COMPONENT sycl-headers-extras)
+
 set(SYCL_RT_LIBS sycl)
 if (MSVC)
   list(APPEND SYCL_RT_LIBS sycld)
@@ -244,6 +251,7 @@ set( SYCL_TOOLCHAIN_DEPLOY_COMPONENTS
      clang-resource-headers
      opencl-headers
      sycl-headers
+     sycl-headers-extras
      sycl
      pi_opencl
 )

sycl/test/xmethods/accessors-device.cpp

@@ -0,0 +1,19 @@
+// RUN: %clangxx -fsycl-device-only -c -fno-color-diagnostics -Xclang -ast-dump %s | FileCheck %s
+// UNSUPPORTED: windows
+#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,25 @@
+// RUN: %clangxx -c -fno-color-diagnostics -Xclang -ast-dump %s | FileCheck %s
+// UNSUPPORTED: windows
+#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())