Merge pull request #220 from frasercrmck/spirv-ll-opencl-dbg

[spirv-ll] Nominally support the DebugInfo/OpenCL.DebugInfo.100 sets
uxlfoundation · Nov 20, 2023 · b7b0874 · b7b0874
2 parents 336052e + 23d0dff
commit b7b0874
Show file tree

Hide file tree

Showing 9 changed files with 540 additions and 4 deletions.
diff --git a/modules/compiler/spirv-ll/CMakeLists.txt b/modules/compiler/spirv-ll/CMakeLists.txt
@@ -47,6 +47,7 @@ endif()
 add_ca_library(spirv-ll STATIC
   ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/assert.h
   ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/builder.h
+  ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/builder_debug_info.h
   ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/builder_glsl.h
   ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/builder_group_async_copies.h
   ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/builder_opencl.h
@@ -55,9 +56,10 @@ add_ca_library(spirv-ll STATIC
   ${CMAKE_CURRENT_SOURCE_DIR}/include/spirv-ll/opcodes.h
   ${CMAKE_CURRENT_SOURCE_DIR}/source/builder.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/source/builder_core.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/source/builder_debug_info.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/source/builder_glsl.cpp
-  ${CMAKE_CURRENT_SOURCE_DIR}/source/builder_opencl.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/source/builder_group_async_copies.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/source/builder_opencl.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/source/context.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/source/module.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/source/opcodes.cpp)

diff --git a/modules/compiler/spirv-ll/include/spirv-ll/builder_debug_info.h b/modules/compiler/spirv-ll/include/spirv-ll/builder_debug_info.h
@@ -0,0 +1,41 @@
+// Copyright (C) Codeplay Software Limited
+//
+// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+// Exceptions; you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+// License for the specific language governing permissions and limitations
+// under the License.
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef SPIRV_LL_SPV_BUILDER_DEBUG_INFO_H_INCLUDED
+#define SPIRV_LL_SPV_BUILDER_DEBUG_INFO_H_INCLUDED
+
+#include <spirv-ll/builder.h>
+
+namespace spirv_ll {
+
+/// @brief Combined builder for the DebugInfo and OpenCLDebugInfo100 extended
+/// instruction sets.
+class DebugInfoBuilder : public ExtInstSetHandler {
+ public:
+  /// @brief Constructor.
+  ///
+  /// @param[in] builder spirv_ll::Builder object that will own this object.
+  /// @param[in] module The module being translated.
+  DebugInfoBuilder(Builder &builder, Module &module)
+      : ExtInstSetHandler(builder, module) {}
+
+  /// @see ExtInstSetHandler::create
+  virtual llvm::Error create(OpExtInst const &opc) override;
+};
+
+}  // namespace spirv_ll
+
+#endif  // SPIRV_LL_SPV_BUILDER_DEBUG_INFO_H_INCLUDED
diff --git a/modules/compiler/spirv-ll/include/spirv-ll/module.h b/modules/compiler/spirv-ll/include/spirv-ll/module.h
@@ -43,9 +43,11 @@
 namespace spirv_ll {
 /// @brief Enum class used to represent an Extended Instruction Set.
 enum class ExtendedInstrSet {
-  GLSL450,           ///< The "GLSL.std.450" instruction set.
-  OpenCL,            ///< The "OpenCL.std" instruction set.
-  GroupAsyncCopies,  ///< The "Codeplay.GroupAsyncCopies" instruction set.
+  GLSL450,             ///< The "GLSL.std.450" instruction set.
+  OpenCL,              ///< The "OpenCL.std" instruction set.
+  GroupAsyncCopies,    ///< The "Codeplay.GroupAsyncCopies" instruction set.
+  DebugInfo,           ///< The "DebugInfo" instruction set.
+  OpenCLDebugInfo100,  ///< The "OpenCL.DebugInfo.100" instruction set.
 };
 
 /// @brief Interface to a binary SPIR-V module's header.

diff --git a/modules/compiler/spirv-ll/source/builder_core.cpp b/modules/compiler/spirv-ll/source/builder_core.cpp
@@ -32,6 +32,7 @@
 #include <multi_llvm/vector_type_helper.h>
 #include <spirv-ll/assert.h>
 #include <spirv-ll/builder.h>
+#include <spirv-ll/builder_debug_info.h>
 #include <spirv-ll/builder_glsl.h>
 #include <spirv-ll/builder_group_async_copies.h>
 #include <spirv-ll/builder_opencl.h>
@@ -317,6 +318,15 @@ llvm::Error Builder::create<OpExtInstImport>(const OpExtInstImport *op) {
         ExtendedInstrSet::GroupAsyncCopies);
     module.associateExtendedInstrSet(op->IdResult(),
                                      ExtendedInstrSet::GroupAsyncCopies);
+  } else if (name == "DebugInfo") {
+    registerExtInstHandler<DebugInfoBuilder>(ExtendedInstrSet::DebugInfo);
+    module.associateExtendedInstrSet(op->IdResult(),
+                                     ExtendedInstrSet::DebugInfo);
+  } else if (name == "OpenCL.DebugInfo.100") {
+    registerExtInstHandler<DebugInfoBuilder>(
+        ExtendedInstrSet::OpenCLDebugInfo100);
+    module.associateExtendedInstrSet(op->IdResult(),
+                                     ExtendedInstrSet::OpenCLDebugInfo100);
   } else {
     return makeStringError(llvm::Twine(name.data()) +
                            " extended instruction set is not supported!\n");

diff --git a/modules/compiler/spirv-ll/source/builder_debug_info.cpp b/modules/compiler/spirv-ll/source/builder_debug_info.cpp
@@ -0,0 +1,44 @@
+
+// Copyright (C) Codeplay Software Limited
+//
+// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+// Exceptions; you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+// License for the specific language governing permissions and limitations
+// under the License.
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <spirv-ll/builder_debug_info.h>
+
+namespace spirv_ll {
+
+llvm::Error DebugInfoBuilder::create(OpExtInst const &) {
+  // We currently let all of these instructions through without question. From
+  // the OpenCL.DebugInfo.100 instruction set:
+  // 2.1 Removing Instructions
+  //   All instructions in this extended set have no semantic impact and can be
+  //   safely removed. This is easily done if all debug instructions are removed
+  //   together, at once. However, when removing a subset, for example, inlining
+  //   a function, there may be dangling references to <id> that have been
+  //   removed. These can be replaced with the Result <id> of the DebugInfoNone
+  //   instruction.
+
+  // Note that this does still assume that none of these instructions are
+  // expected to produce LLVM values for anything other than non-semantic
+  // instructions in these same extended instruction sets (mixing and matching
+  // DebugInfo OpenCL.DebugInfo.100 is fine). However, it's an unlikely
+  // scenario that anything produced by instructions in these sets is used by
+  // another instruction set we support: these instructions are all debug info,
+  // and all instructions in this set return 'OpTypeVoid' so can't really be
+  // used/referenced by most ops in a meaningful way anyway.
+  return llvm::Error::success();
+}
+
+}  // namespace spirv_ll
diff --git a/modules/compiler/spirv-ll/source/context.cpp b/modules/compiler/spirv-ll/source/context.cpp
@@ -176,6 +176,9 @@ cargo::expected<spirv_ll::Module, spirv_ll::Error> spirv_ll::Context::translate(
       case spv::OpSourceExtension:
         error = builder.create<OpSourceExtension>(op);
         break;
+      case spv::OpModuleProcessed:
+        // Ignore this.
+        break;
       case spv::OpName:
         error = builder.create<OpName>(op);
         break;

diff --git a/modules/compiler/spirv-ll/test/spvasm/CMakeLists.txt b/modules/compiler/spirv-ll/test/spvasm/CMakeLists.txt
@@ -2197,6 +2197,8 @@ set(SPVASM_V1_1_FILES
   op_execution_mode_subgroups_per_workgroup.spvasm
   op_get_default_queue.spvasm
   op_get_kernel_max_num_subgroups.spvasm
+  opencl_debug_info_100.spvasm
+  opencl_debug_info_100_clean.spvasm
 )
 
 set(UNVERIFIABLE_SPVASM_FILES
@@ -2207,6 +2209,9 @@ set(UNVERIFIABLE_SPVASM_FILES
   codeplay/opencl_usm_generic_address_space.spvasm
   codeplay/opencl_group_async_copy_2d2d.spvasm
   codeplay/opencl_group_async_copy_3d3d.spvasm
+  # Unclear whether or not the SPIR-V is valid, but the output of llvm-spirv
+  # produces this and the verifier doesn't like it.
+  opencl_debug_info_100.spvasm
 )
 
 # Remove obsolete lit test inputs from the binary directory.

diff --git a/modules/compiler/spirv-ll/test/spvasm/opencl_debug_info_100.spvasm b/modules/compiler/spirv-ll/test/spvasm/opencl_debug_info_100.spvasm
@@ -0,0 +1,212 @@
+; Copyright (C) Codeplay Software Limited
+;
+; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+; Exceptions; you may not use this file except in compliance with the License.
+; You may obtain a copy of the License at
+;
+;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+;
+; Unless required by applicable law or agreed to in writing, software
+; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+; License for the specific language governing permissions and limitations
+; under the License.
+;
+; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+; This test was created by compiling the following OpenCL kernel:
+;
+;  kernel void foo(global int *a, global float *b, global char *c) {
+;    size_t id = get_global_id(0);
+;    int a0 = a[id];
+;    int b0 = convert_int(b[id]);
+;    a0 /= 2;
+;    c[id] = a0 * b0;
+;  }
+;
+; with:
+;   clang-17 -target spir64-unknown-unknown -x cl -cl-kernel-arg-info -g \
+;            -emit-llvm -o kernel.bc kernel.cl
+; then:
+;   llvm-spirv-18 kernel.bc -o kernel.spv --spirv-debug-info-version=ocl-100 \
+;                 --debugify-level='location+variables'
+; then:
+;   spirv-dis kernel.spv
+;
+; And then the following line was removed because it was using an extension
+; without declaring it, and it's unnecessary:
+;   OpDecorate %mul NoSignedWrap
+;
+; It is classed as unverifiable because even current versions of spirv-val
+; (v2023.5 at the time of writing) fail to verify this program. There are
+; various issues depending on the version, so they aren't listed.
+; A verifier-clean version is located elsewhere in this test suite.
+
+; RUN: %if online-spirv-as %{ spirv-as --target-env spv1.1 -o %spv_file_s %s %}
+; RUN: spirv-ll-tool -a OpenCL -b 64 -c Int64 %spv_file_s | FileCheck %s
+
+; SPIR-V
+; Version: 1.4
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 84
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Linkage
+               OpCapability Kernel
+               OpCapability Int64
+               OpCapability Int8
+          %1 = OpExtInstImport "OpenCL.std"
+          %2 = OpExtInstImport "OpenCL.DebugInfo.100"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %15 "foo" %__spirv_BuiltInGlobalInvocationId
+         %40 = OpString "/tmp/kernel.cl"
+         %41 = OpString "//__CSK_MD5:778eaa9cd957459f44d042d2dd1bf574"
+         %44 = OpString "int"
+         %48 = OpString "float"
+         %51 = OpString "char"
+         %56 = OpString "size_t"
+         %57 = OpString "unsigned long"
+         %60 = OpString "/include/opencl-c-base.h"
+         %61 = OpString "//__CSK_MD5:8040a97cda029467f3f64c25e932a46e"
+         %64 = OpString "foo"
+         %65 = OpString ""
+         %67 = OpString "a"
+         %69 = OpString "b"
+         %71 = OpString "c"
+         %73 = OpString "id"
+         %75 = OpString "a0"
+         %77 = OpString "b0"
+               OpSource OpenCL_C 102000
+               OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
+               OpName %a "a"
+               OpName %b "b"
+               OpName %c "c"
+               OpName %entry "entry"
+               OpName %call "call"
+               OpName %arrayidx "arrayidx"
+               OpName %arrayidx1 "arrayidx1"
+               OpName %call2 "call2"
+               OpName %div "div"
+               OpName %mul "mul"
+               OpName %conv "conv"
+               OpName %arrayidx3 "arrayidx3"
+               OpModuleProcessed "Debug info producer: clang version 17.0.3 ([email protected]:ComputeAorta/llvm-project.git 37b79e779f447f1c714af7f907e7a2ec846d1da0)"
+               OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
+               OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
+               OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
+               OpDecorate %a FuncParamAttr NoCapture
+               OpDecorate %a FuncParamAttr NoWrite
+               OpDecorate %a Alignment 4
+               OpDecorate %b FuncParamAttr NoCapture
+               OpDecorate %b FuncParamAttr NoWrite
+               OpDecorate %b Alignment 4
+               OpDecorate %c FuncParamAttr NoCapture
+               OpDecorate %c Alignment 1
+      %ulong = OpTypeInt 64 0
+       %uint = OpTypeInt 32 0
+      %uchar = OpTypeInt 8 0
+     %uint_2 = OpConstant %uint 2
+    %uint_32 = OpConstant %uint 32
+     %uint_8 = OpConstant %uint 8
+    %uint_64 = OpConstant %uint 64
+    %v3ulong = OpTypeVector %ulong 3
+%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
+       %void = OpTypeVoid
+%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
+      %float = OpTypeFloat 32
+%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
+%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
+         %14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_float %_ptr_CrossWorkgroup_uchar
+%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
+         %20 = OpExtInst %void %2 DebugInfoNone
+         %42 = OpExtInst %void %2 DebugSource %40 %41
+         %43 = OpExtInst %void %2 DebugCompilationUnit 65536 5 %42 OpenCL_C
+         %46 = OpExtInst %void %2 DebugTypeBasic %44 %uint_32 Signed
+         %47 = OpExtInst %void %2 DebugTypePointer %46 CrossWorkgroup None
+         %49 = OpExtInst %void %2 DebugTypeBasic %48 %uint_32 Float
+         %50 = OpExtInst %void %2 DebugTypePointer %49 CrossWorkgroup None
+         %53 = OpExtInst %void %2 DebugTypeBasic %51 %uint_8 SignedChar
+         %54 = OpExtInst %void %2 DebugTypePointer %53 CrossWorkgroup None
+         %55 = OpExtInst %void %2 DebugTypeFunction None %20 %47 %50 %54
+         %59 = OpExtInst %void %2 DebugTypeBasic %57 %uint_64 Unsigned
+         %62 = OpExtInst %void %2 DebugSource %60 %61
+         %63 = OpExtInst %void %2 DebugTypedef %56 %59 %62 0 0 %43
+         %66 = OpExtInst %void %2 DebugFunction %64 %55 %42 1 0 %43 %65 FlagIsDefinition|FlagPrototyped|FlagIsOptimized 1 %15 %20
+         %68 = OpExtInst %void %2 DebugLocalVariable %67 %47 %42 1 0 %66 None 1
+         %70 = OpExtInst %void %2 DebugLocalVariable %69 %50 %42 1 0 %66 None 2
+         %72 = OpExtInst %void %2 DebugLocalVariable %71 %54 %42 1 0 %66 None 3
+         %74 = OpExtInst %void %2 DebugLocalVariable %73 %63 %42 2 0 %66 None
+         %76 = OpExtInst %void %2 DebugLocalVariable %75 %46 %42 3 0 %66 None
+         %78 = OpExtInst %void %2 DebugLocalVariable %77 %46 %42 4 0 %66 None
+         %79 = OpExtInst %void %2 DebugOperation Constu 0
+         %80 = OpExtInst %void %2 DebugOperation Swap
+         %81 = OpExtInst %void %2 DebugOperation Xderef
+         %82 = OpExtInst %void %2 DebugExpression %79 %80 %81
+
+; CHECK: define spir_kernel void @foo(ptr addrspace(1){{.*}} %a, ptr addrspace(1){{.*}} %b, ptr addrspace(1){{.*}} %c)
+         %15 = OpFunction %void None %14
+          %a = OpFunctionParameter %_ptr_CrossWorkgroup_uint
+          %b = OpFunctionParameter %_ptr_CrossWorkgroup_float
+          %c = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
+      %entry = OpLabel
+         %83 = OpExtInst %void %2 DebugScope %66
+               OpLine %40 0 0
+         %21 = OpExtInst %void %2 DebugValue %68 %a %82
+         %22 = OpExtInst %void %2 DebugValue %70 %b %82
+         %23 = OpExtInst %void %2 DebugValue %72 %c %82
+               OpLine %40 2 15
+; CHECK: [[ID:%.*]] = call spir_func i64 @_Z13get_global_idj(i32 {{.*}}), !dbg [[IDLoc:![0-9]+]]
+         %24 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
+       %call = OpCompositeExtract %ulong %24 0
+               OpLine %40 0 0
+         %26 = OpExtInst %void %2 DebugValue %74 %call %82
+               OpLine %40 3 12
+; CHECK: [[A0IDX:%.*]] = getelementptr inbounds i32, ptr addrspace(1) %a, i64 [[ID]], !dbg [[A0Loc:![0-9]+]]
+   %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %a %call
+; CHECK: [[A0:%.*]] = load i32, ptr addrspace(1) [[A0IDX]], align 4, !dbg [[A0Loc]]
+         %28 = OpLoad %uint %arrayidx Aligned 4
+               OpLine %40 0 0
+         %29 = OpExtInst %void %2 DebugValue %76 %28 %82
+               OpLine %40 4 24
+; CHECK: [[B0IDX:%.*]] = getelementptr inbounds float, ptr addrspace(1) %b, i64 [[ID]], !dbg [[B0LDLoc:![0-9]+]]
+  %arrayidx1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %b %call
+; CHECK: [[B0LD:%.*]] = load float, ptr addrspace(1) [[B0IDX]], align 4, !dbg [[B0LDLoc]]
+         %31 = OpLoad %float %arrayidx1 Aligned 4
+               OpLine %40 4 12
+; CHECK: [[B0:%.*]] = call spir_func i32 @_Z11convert_intf(float [[B0LD]]), !dbg [[B0Loc:![0-9]+]]
+      %call2 = OpConvertFToS %uint %31
+               OpLine %40 0 0
+         %33 = OpExtInst %void %2 DebugValue %78 %call2 %82
+               OpLine %40 5 6
+; CHECK: [[DIV:%.*]] = sdiv i32 [[A0]], 2, !dbg [[DIVLoc:![0-9]+]]
+        %div = OpSDiv %uint %28 %uint_2
+               OpLine %40 0 0
+         %36 = OpExtInst %void %2 DebugValue %76 %div %82
+               OpLine %40 6 14
+; CHECK: [[MUL:%.*]] = mul i32 [[B0]], [[DIV]], !dbg [[MULLoc:![0-9]+]]
+        %mul = OpIMul %uint %call2 %div
+               OpLine %40 6 11
+; CHECK: [[TRUNC:%.*]] = trunc i32 [[MUL]] to i8, !dbg [[TRUNCLoc:![0-9]+]]
+       %conv = OpUConvert %uchar %mul
+               OpLine %40 6 3
+; CHECK: [[C0IDX:%.*]] = getelementptr inbounds i8, ptr addrspace(1) %c, i64 [[ID]], !dbg [[C0IDXLoc:![0-9]+]]
+  %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %c %call
+               OpLine %40 6 9
+; CHECK: store i8 [[TRUNC]], ptr addrspace(1) [[C0IDX]], align 1, !dbg [[STORELoc:![0-9]+]]
+               OpStore %arrayidx3 %conv Aligned 1
+               OpLine %40 7 1
+; CHECK: ret void, !dbg [[RETLoc:![0-9]+]]
+               OpReturn
+               OpFunctionEnd
+
+; CHECK-DAG: [[SCOPE:![0-9]+]] = distinct !DILexicalBlock(
+; CHECK-DAG: [[IDLoc]] = !DILocation(line: 2, column: 15, scope: [[SCOPE]])
+; CHECK-DAG: [[A0Loc]] = !DILocation(line: 3, column: 12, scope: [[SCOPE]])
+; CHECK-DAG: [[B0LDLoc]] = !DILocation(line: 4, column: 24, scope: [[SCOPE]])
+; CHECK-DAG: [[B0Loc]] = !DILocation(line: 4, column: 12, scope: [[SCOPE]])
+; CHECK-DAG: [[DIVLoc]] = !DILocation(line: 5, column: 6, scope: [[SCOPE]])
+; CHECK-DAG: [[MULLoc]] = !DILocation(line: 6, column: 14, scope: [[SCOPE]])
+; CHECK-DAG: [[TRUNCLoc]] = !DILocation(line: 6, column: 11, scope: [[SCOPE]])
+; CHECK-DAG: [[C0IDXLoc]] = !DILocation(line: 6, column: 3, scope: [[SCOPE]])
+; CHECK-DAG: [[STORELoc]] = !DILocation(line: 6, column: 9, scope: [[SCOPE]])
+; CHECK-DAG: [[RETLoc]] = !DILocation(line: 7, column: 1, scope: [[SCOPE]])