Skip to content

Instantly share code, notes, and snippets.

@benvanik

benvanik/0-input.mlir

Last active April 9, 2020 05:18
Show Gist options
  • Save benvanik/cc49403752e2002b872b7d3a17df3cc0 to your computer and use it in GitHub Desktop.
Save benvanik/cc49403752e2002b872b7d3a17df3cc0 to your computer and use it in GitHub Desktop.
Download ZIP
simple_mul target-specific interface
Raw
0-input.mlir
module {
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export} {
%0 = xla_hlo.multiply %arg0, %arg1 {name = "mul.1"} : tensor<4xf32>
return %0 : tensor<4xf32>
}
}
Raw
1-flow.mlir
module {
flow.executable @simple_mul_ex_dispatch_0 {
flow.dispatch.entry @simple_mul_ex_dispatch_0
module {
func @simple_mul_ex_dispatch_0(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
%0 = xla_hlo.multiply %arg0, %arg1 {name = "mul.1"} : tensor<4xf32>
return %0 : tensor<4xf32>
}
}
}
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export, iree.reflection = {f = "I11!B3!d4B3!d4R6!B3!d4", fv = "1"}} {
%c4 = constant 4 : index
%0 = flow.ex.stream.fragment(%arg2 = %c4 : index, %arg3 = %arg0 : tensor<4xf32>, %arg4 = %arg1 : tensor<4xf32>) -> tensor<4xf32> {
%1 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg2 : index](%arg3, %arg4) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
flow.return %1 : tensor<4xf32>
}
return %0 : tensor<4xf32>
}
}
Raw
2-hal-input.mlir
module {
hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} {
hal.interface @legacy_io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
}
hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]}
hal.executable.source "llvm-ir" {
module {
func @simple_mul_ex_dispatch_0() {
%c0 = constant 0 : index
%0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<4xf32>
%1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<4xf32>
%2 = call @simple_mul_ex_dispatch_0_impl(%0, %1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<4xf32>
return
}
func @simple_mul_ex_dispatch_0_impl(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> attributes {sym_visibility = "private"} {
%0 = xla_hlo.multiply %arg0, %arg1 {name = "mul.1"} : tensor<4xf32>
return %0 : tensor<4xf32>
}
hal.interface @legacy_io attributes {sym_visibility = "private"} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
}
}
}
hal.executable.source "vmla" {
module {
func @simple_mul_ex_dispatch_0() {
%c0 = constant 0 : index
%0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<4xf32>
%1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<4xf32>
%2 = call @simple_mul_ex_dispatch_0_impl(%0, %1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<4xf32>
return
}
func @simple_mul_ex_dispatch_0_impl(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> attributes {sym_visibility = "private"} {
%0 = xla_hlo.multiply %arg0, %arg1 {name = "mul.1"} : tensor<4xf32>
return %0 : tensor<4xf32>
}
hal.interface @legacy_io attributes {sym_visibility = "private"} {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
}
}
}
}
func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> attributes {iree.module.export, iree.reflection = {f = "I11!B3!d4B3!d4R6!B3!d4", fv = "1"}} {
%c4 = constant 4 : index
%0 = flow.ex.stream.fragment(%arg2 = %c4 : index, %arg3 = %arg0 : tensor<4xf32>, %arg4 = %arg1 : tensor<4xf32>) -> tensor<4xf32> {
%1 = flow.dispatch @simple_mul_ex_dispatch_0::@simple_mul_ex_dispatch_0[%arg2 : index](%arg3, %arg4) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
flow.return %1 : tensor<4xf32>
}
return %0 : tensor<4xf32>
}
}
Raw
3-hal-middle.mlir
module {
hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} {
hal.interface @legacy_io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
}
hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]}
hal.executable.target "llvm-ir" {
module {
llvm.func @simple_mul_ex_dispatch_0(%arg0: !llvm<"float*"> {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32}}, %arg1: !llvm<"float*"> {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32}}, %arg2: !llvm.i64 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32}}, %arg3: !llvm.i64 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32}}, %arg4: !llvm.i64 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32}}, %arg5: !llvm<"float*"> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32}}, %arg6: !llvm<"float*"> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32}}, %arg7: !llvm.i64 {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32}}, %arg8: !llvm.i64 {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32}}, %arg9: !llvm.i64 {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32}}, %arg10: !llvm<"float*"> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32}}, %arg11: !llvm<"float*"> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32}}, %arg12: !llvm.i64 {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32}}, %arg13: !llvm.i64 {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32}}, %arg14: !llvm.i64 {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32}}) attributes {iree.dispatch_fn_name = "simple_mul_ex_dispatch_0", llvm.emit_c_interface, spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} {
%0 = llvm.mlir.undef : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%1 = llvm.insertvalue %arg0, %0[0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%2 = llvm.insertvalue %arg1, %1[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%3 = llvm.insertvalue %arg2, %2[2] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%4 = llvm.insertvalue %arg3, %3[3, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%5 = llvm.insertvalue %arg4, %4[4, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%6 = llvm.insertvalue %arg5, %0[0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%7 = llvm.insertvalue %arg6, %6[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%8 = llvm.insertvalue %arg7, %7[2] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%9 = llvm.insertvalue %arg8, %8[3, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%10 = llvm.insertvalue %arg9, %9[4, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%11 = llvm.insertvalue %arg10, %0[0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%12 = llvm.insertvalue %arg11, %11[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%13 = llvm.insertvalue %arg12, %12[2] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%14 = llvm.insertvalue %arg13, %13[3, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%15 = llvm.insertvalue %arg14, %14[4, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%16 = llvm.mlir.constant(4 : index) : !llvm.i64
%17 = llvm.mlir.constant(1 : index) : !llvm.i64
%18 = llvm.mlir.constant(0 : index) : !llvm.i64
llvm.br ^bb1(%18 : !llvm.i64)
^bb1(%19: !llvm.i64): // 2 preds: ^bb0, ^bb2
%20 = llvm.icmp "slt" %19, %16 : !llvm.i64
llvm.cond_br %20, ^bb2, ^bb3
^bb2: // pred: ^bb1
%21 = llvm.extractvalue %5[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%22 = llvm.mul %19, %17 : !llvm.i64
%23 = llvm.add %18, %22 : !llvm.i64
%24 = llvm.getelementptr %21[%23] : (!llvm<"float*">, !llvm.i64) -> !llvm<"float*">
%25 = llvm.load %24 : !llvm<"float*">
%26 = llvm.extractvalue %10[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%27 = llvm.getelementptr %26[%23] : (!llvm<"float*">, !llvm.i64) -> !llvm<"float*">
%28 = llvm.load %27 : !llvm<"float*">
%29 = llvm.extractvalue %15[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%30 = llvm.getelementptr %29[%23] : (!llvm<"float*">, !llvm.i64) -> !llvm<"float*">
%31 = llvm.load %30 : !llvm<"float*">
%32 = llvm.fmul %25, %28 : !llvm.float
llvm.store %32, %30 : !llvm<"float*">
%33 = llvm.add %19, %17 : !llvm.i64
llvm.br ^bb1(%33 : !llvm.i64)
^bb3: // pred: ^bb1
llvm.return
}
llvm.func @_mlir_ciface_simple_mul_ex_dispatch_0(%arg0: !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }*"> {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32}}, %arg1: !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }*"> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32}}, %arg2: !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }*"> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32}}) attributes {iree.dispatch_fn_name = "simple_mul_ex_dispatch_0", llvm.emit_c_interface, spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} {
%0 = llvm.load %arg0 : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }*">
%1 = llvm.extractvalue %0[0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%2 = llvm.extractvalue %0[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%3 = llvm.extractvalue %0[2] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%4 = llvm.extractvalue %0[3, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%5 = llvm.extractvalue %0[4, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%6 = llvm.load %arg1 : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }*">
%7 = llvm.extractvalue %6[0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%8 = llvm.extractvalue %6[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%9 = llvm.extractvalue %6[2] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%10 = llvm.extractvalue %6[3, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%11 = llvm.extractvalue %6[4, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%12 = llvm.load %arg2 : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }*">
%13 = llvm.extractvalue %12[0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%14 = llvm.extractvalue %12[1] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%15 = llvm.extractvalue %12[2] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%16 = llvm.extractvalue %12[3, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
%17 = llvm.extractvalue %12[4, 0] : !llvm<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
llvm.call @simple_mul_ex_dispatch_0(%1, %2, %3, %4, %5, %7, %8, %9, %10, %11, %13, %14, %15, %16, %17) : (!llvm<"float*">, !llvm<"float*">, !llvm.i64, !llvm.i64, !llvm.i64, !llvm<"float*">, !llvm<"float*">, !llvm.i64, !llvm.i64, !llvm.i64, !llvm<"float*">, !llvm<"float*">, !llvm.i64, !llvm.i64, !llvm.i64) -> ()
llvm.return
}
}
}
hal.executable.target "vmla" {
module {
vm.module @module {
vm.func @simple_mul_ex_dispatch_0(%arg0: !vm.ref<!vmla.interface>) {
%zero = vm.const.i32.zero : i32
%c16 = vm.const.i32 16 : i32
%ref = vm.call @vmla.interface.binding(%arg0, %zero, %zero) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer>
%ref_0 = vm.call @vmla.buffer.view(%ref, %zero, %c16) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer>
%c1 = vm.const.i32 1 : i32
%ref_1 = vm.call @vmla.interface.binding(%arg0, %zero, %c1) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer>
%ref_2 = vm.call @vmla.buffer.view(%ref_1, %zero, %c16) : (!vm.ref<!vmla.buffer>, i32, i32) -> !vm.ref<!vmla.buffer>
%ref_3 = vm.call @vmla.buffer.alloc(%c16) : (i32) -> !vm.ref<!vmla.buffer>
vm.call @vmla.mul.f32(%ref_0, %ref_2, %ref_3) : (!vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>, !vm.ref<!vmla.buffer>) -> ()
%c2 = vm.const.i32 2 : i32
%ref_4 = vm.call @vmla.interface.binding(%arg0, %zero, %c2) : (!vm.ref<!vmla.interface>, i32, i32) -> !vm.ref<!vmla.buffer>
vm.call @vmla.buffer.copy(%ref_3, %zero, %ref_4, %zero, %c16) : (!vm.ref<!vmla.buffer>, i32, !vm.ref<!vmla.buffer>, i32, i32) -> ()
vm.return
}
vm.export @simple_mul_ex_dispatch_0
vm.import @vmla.interface.binding(%interface : !vm.ref<!vmla.interface>, %set : i32, %binding : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"}
vm.import @vmla.buffer.alloc(%byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"}
vm.import @vmla.buffer.view(%src : !vm.ref<!vmla.buffer>, %byte_offset : i32, %byte_length : i32) -> !vm.ref<!vmla.buffer> attributes {nosideeffects, sym_visibility = "private"}
vm.import @vmla.buffer.copy(%src : !vm.ref<!vmla.buffer>, %src_byte_offset : i32, %dst : !vm.ref<!vmla.buffer>, %dst_byte_offset : i32, %byte_length : i32) attributes {sym_visibility = "private"}
vm.import @vmla.mul.f32(%lhs : !vm.ref<!vmla.buffer>, %rhs : !vm.ref<!vmla.buffer>, %dst : !vm.ref<!vmla.buffer>) attributes {sym_visibility = "private"}
}
}
}
}
func @simple_mul(%arg0: !hal.buffer, %arg1: !hal.buffer) -> !hal.buffer attributes {iree.module.export = "simple_mul$raw"} {
%c0 = constant 0 : index
%c4 = constant 4 : index
%dev = hal.ex.shared_device : !hal.device
%allocator = hal.device.allocator %dev : !hal.allocator
%sz = hal.allocator.compute_size %allocator, shape = [%c4], element_type = 50331680
%buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
hal.ex.defer_release %buffer : !hal.buffer
%cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
hal.command_buffer.begin %cmd
%exe = hal.executable.lookup %dev, @simple_mul_ex_dispatch_0 : !hal.executable
%executable_layout = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
%allocator_0 = hal.buffer.allocator %arg0 : !hal.allocator
%sz_1 = hal.allocator.compute_size %allocator_0, shape = [%c4], element_type = 50331680
%allocator_2 = hal.buffer.allocator %arg1 : !hal.allocator
%sz_3 = hal.allocator.compute_size %allocator_2, shape = [%c4], element_type = 50331680
hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%arg0, %c0, %sz_1), 1 = (%arg1, %c0, %sz_3), 2 = (%buffer, %c0, %sz)]
hal.device.switch(%dev : !hal.device)
#hal.device.match.id<"llvm-ir">(%arg2 = %c4 : index, %arg3 = %cmd : !hal.command_buffer, %arg4 = %exe : !hal.executable) {
%c1 = constant 1 : index
%0 = addi %arg2, %c1 : index
%1 = subi %0, %c1 : index
%2 = divi_unsigned %1, %c1 : index
%3 = muli %2, %1 : index
%4 = subi %arg2, %3 : index
%5 = cmpi "sge", %4, %c1 : index
%6 = select %5, %4, %c1 : index
%7 = addi %6, %c1 : index
%8 = subi %7, %c1 : index
%9 = divi_unsigned %8, %c1 : index
%10 = muli %9, %8 : index
%11 = subi %6, %10 : index
%12 = cmpi "sge", %11, %c1 : index
%13 = select %12, %11, %c1 : index
%14 = addi %13, %c1 : index
%15 = subi %14, %c1 : index
%16 = divi_unsigned %15, %c1 : index
hal.command_buffer.dispatch %arg3, %arg4, entry_point = 0, workgroup_xyz = [%2, %9, %16]
hal.return
},
#hal.device.match.id<"vmla">(%arg2 = %c4 : index, %arg3 = %cmd : !hal.command_buffer, %arg4 = %exe : !hal.executable) {
%c1 = constant 1 : index
%0 = addi %arg2, %c1 : index
%1 = subi %0, %c1 : index
%2 = divi_unsigned %1, %c1 : index
%3 = muli %2, %1 : index
%4 = subi %arg2, %3 : index
%5 = cmpi "sge", %4, %c1 : index
%6 = select %5, %4, %c1 : index
%7 = addi %6, %c1 : index
%8 = subi %7, %c1 : index
%9 = divi_unsigned %8, %c1 : index
%10 = muli %9, %8 : index
%11 = subi %6, %10 : index
%12 = cmpi "sge", %11, %c1 : index
%13 = select %12, %11, %c1 : index
%14 = addi %13, %c1 : index
%15 = subi %14, %c1 : index
%16 = divi_unsigned %15, %c1 : index
hal.command_buffer.dispatch %arg3, %arg4, entry_point = 0, workgroup_xyz = [%2, %9, %16]
hal.return
}
%memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
hal.command_buffer.end %cmd
hal.ex.submit_and_wait %dev, %cmd
return %buffer : !hal.buffer
}
func @simple_mul$sync(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I11!B3!d4B3!d4R6!B3!d4", fv = "1"}} {
%buffer = hal.buffer_view.buffer %arg0 : !hal.buffer
%buffer_0 = hal.buffer_view.buffer %arg1 : !hal.buffer
%0 = call @simple_mul(%buffer, %buffer_0) : (!hal.buffer, !hal.buffer) -> !hal.buffer
%c4 = constant 4 : index
%view = hal.buffer_view.create %0, shape = [%c4], element_type = 50331680 : !hal.buffer_view
return %view : !hal.buffer_view
}
}
Raw
4-hal-output.mlir
module {
func @_device_match_id_0_initializer() -> i1 attributes {sym_visibility = "private"} {
%dev = hal.ex.shared_device : !hal.device
%0 = hal.device.match.id %dev, pattern = ["llvm-ir"] : (!hal.device) -> i1
return %0 : i1
}
hal.variable @_device_match_id_0 init(@_device_match_id_0_initializer) : i1
func @_device_match_id_1_initializer() -> i1 attributes {sym_visibility = "private"} {
%dev = hal.ex.shared_device : !hal.device
%0 = hal.device.match.id %dev, pattern = ["vmla"] : (!hal.device) -> i1
return %0 : i1
}
hal.variable @_device_match_id_1 init(@_device_match_id_1_initializer) : i1
hal.variable @_executable_simple_mul_ex_dispatch_0 mutable : !hal.executable
hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout
func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
%dev = hal.ex.shared_device : !hal.device
%descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
return %descriptor_set_layout : !hal.descriptor_set_layout
}
hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout
func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
%0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
%dev = hal.ex.shared_device : !hal.device
%executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
return %executable_layout : !hal.executable_layout
}
hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache
func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
%dev = hal.ex.shared_device : !hal.device
%executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
%0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
%executable_simple_mul_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @simple_mul_ex_dispatch_0 : !hal.executable
hal.variable.store %executable_simple_mul_ex_dispatch_0, @_executable_simple_mul_ex_dispatch_0 : !hal.executable
return %executable_cache_default : !hal.executable_cache
}
hal.executable @simple_mul_ex_dispatch_0 attributes {sym_visibility = "private"} {
hal.interface @legacy_io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
}
hal.executable.entry_point @simple_mul_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>, workgroup_size = [1 : index, 1 : index, 1 : index]}
hal.executable.binary attributes {data = dense<"0x100000004C4C564D08000C000400080008000000541B000004000000481B00003B204D6F64756C654944203D20274C4C564D4469616C6563744D6F64756C65270A736F757263655F66696C656E616D65203D20224C4C564D4469616C6563744D6F64756C65220A0A6465636C6172652069382A20406D616C6C6F6328693634290A0A6465636C61726520766F69642040667265652869382A290A0A646566696E6520766F6964204073696D706C655F6D756C5F65785F64697370617463685F3028666C6F61742A2025302C20666C6F61742A2025312C206936342025322C206936342025332C206936342025342C20666C6F61742A2025352C20666C6F61742A2025362C206936342025372C206936342025382C206936342025392C20666C6F61742A202531302C20666C6F61742A202531312C20693634202531322C20693634202531332C2069363420253134292021646267202133207B0A2020253136203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D20756E6465662C20666C6F61742A2025302C20300A2020253137203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20666C6F61742A2025312C20310A2020253138203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531372C206936342025322C20320A2020253139203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531382C206936342025332C20332C20300A2020253230203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531392C206936342025342C20342C20300A2020253231203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D20756E6465662C20666C6F61742A2025352C20300A2020253232203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532312C20666C6F61742A2025362C20310A2020253233203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532322C206936342025372C20320A2020253234203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532332C206936342025382C20332C20300A2020253235203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532342C206936342025392C20342C20300A2020253236203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D20756E6465662C20666C6F61742A202531302C20300A2020253237203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532362C20666C6F61742A202531312C20310A2020253238203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532372C20693634202531322C20320A2020253239203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532382C20693634202531332C20332C20300A2020253330203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532392C20693634202531342C20342C20300A20206272206C6162656C202533312C20216462672021370A0A33313A20202020202020202020202020202020202020202020202020202020202020202020202020202020202020202020203B207072656473203D202533342C202531350A2020253332203D2070686920693634205B202534372C20253334205D2C205B20302C20253135205D0A2020253333203D2069636D7020736C7420693634202533322C20342C20216462672021370A20206272206931202533332C206C6162656C202533342C206C6162656C202534382C20216462672021370A0A33343A20202020202020202020202020202020202020202020202020202020202020202020202020202020202020202020203B207072656473203D202533310A2020253335203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532302C20312C20216462672021370A2020253336203D206D756C20693634202533322C20312C20216462672021370A2020253337203D206164642069363420302C202533362C20216462672021370A2020253338203D20676574656C656D656E7470747220666C6F61742C20666C6F61742A202533352C20693634202533372C20216462672021370A2020253339203D206C6F616420666C6F61742C20666C6F61742A202533382C20216462672021370A2020253430203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532352C20312C20216462672021370A2020253431203D20676574656C656D656E7470747220666C6F61742C20666C6F61742A202534302C20693634202533372C20216462672021370A2020253432203D206C6F616420666C6F61742C20666C6F61742A202534312C20216462672021370A2020253433203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202533302C20312C20216462672021370A2020253434203D20676574656C656D656E7470747220666C6F61742C20666C6F61742A202534332C20693634202533372C20216462672021370A2020253435203D206C6F616420666C6F61742C20666C6F61742A202534342C20216462672021370A2020253436203D20666D756C20666C6F6174202533392C202534322C20216462672021370A202073746F726520666C6F6174202534362C20666C6F61742A202534342C20216462672021370A2020253437203D2061646420693634202533322C20312C20216462672021370A20206272206C6162656C202533312C20216462672021370A0A34383A20202020202020202020202020202020202020202020202020202020202020202020202020202020202020202020203B207072656473203D202533310A202072657420766F69642C20216462672021370A7D0A0A646566696E6520766F696420405F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30287B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025302C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025312C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A202532292021646267202139207B0A20202534203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025302C2021646267202131300A20202535203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20302C2021646267202131300A20202536203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20312C2021646267202131300A20202537203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20322C2021646267202131300A20202538203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20332C20302C2021646267202131300A20202539203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20342C20302C2021646267202131300A2020253130203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025312C2021646267202131300A2020253131203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20302C2021646267202131300A2020253132203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20312C2021646267202131300A2020253133203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20322C2021646267202131300A2020253134203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20332C20302C2021646267202131300A2020253135203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20342C20302C2021646267202131300A2020253136203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025322C2021646267202131300A2020253137203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20302C2021646267202131300A2020253138203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20312C2021646267202131300A2020253139203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20322C2021646267202131300A2020253230203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20332C20302C2021646267202131300A2020253231203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20342C20302C2021646267202131300A202063616C6C20766F6964204073696D706C655F6D756C5F65785F64697370617463685F3028666C6F61742A2025352C20666C6F61742A2025362C206936342025372C206936342025382C206936342025392C20666C6F61742A202531312C20666C6F61742A202531322C20693634202531332C20693634202531342C20693634202531352C20666C6F61742A202531372C20666C6F61742A202531382C20693634202531392C20693634202532302C2069363420253231292C2021646267202131300A202072657420766F69642C2021646267202131300A7D0A0A646566696E6520766F69642040696E766F6B655F5F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F302869382A2A20253029207B0A20202532203D20676574656C656D656E747074722069382A2C2069382A2A2025302C2069363420300A20202533203D206C6F61642069382A2C2069382A2A2025320A20202534203D20626974636173742069382A20253320746F207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A0A20202535203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A2025340A20202536203D20676574656C656D656E747074722069382A2C2069382A2A2025302C2069363420310A20202537203D206C6F61642069382A2C2069382A2A2025360A20202538203D20626974636173742069382A20253720746F207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A0A20202539203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A2025380A2020253130203D20676574656C656D656E747074722069382A2C2069382A2A2025302C2069363420320A2020253131203D206C6F61642069382A2C2069382A2A202531300A2020253132203D20626974636173742069382A2025313120746F207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A0A2020253133203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A202531320A202063616C6C20766F696420405F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30287B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025352C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025392C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A20253133290A202072657420766F69640A7D0A0A216C6C766D2E6462672E6375203D20217B21307D0A216C6C766D2E6D6F64756C652E666C616773203D20217B21327D0A0A2130203D2064697374696E637420214449436F6D70696C65556E6974286C616E67756167653A2044575F4C414E475F432C2066696C653A2021312C2070726F64756365723A20226D6C6972222C2069734F7074696D697A65643A20747275652C2072756E74696D6556657273696F6E3A20302C20656D697373696F6E4B696E643A2046756C6C4465627567290A2131203D2021444946696C652866696C656E616D653A20224C4C564D4469616C6563744D6F64756C65222C206469726563746F72793A20222F22290A2132203D20217B69333220322C202122446562756720496E666F2056657273696F6E222C2069333220337D0A2133203D2064697374696E63742021444953756270726F6772616D286E616D653A202273696D706C655F6D756C5F65785F64697370617463685F30222C206C696E6B6167654E616D653A202273696D706C655F6D756C5F65785F64697370617463685F30222C2073636F70653A206E756C6C2C2066696C653A2021342C206C696E653A20332C20747970653A2021352C2073636F70654C696E653A20332C207370466C6167733A2044495350466C6167446566696E6974696F6E207C2044495350466C61674F7074696D697A65642C20756E69743A2021302C2072657461696E65644E6F6465733A202136290A2134203D2021444946696C652866696C656E616D653A2022697265655C5C73616D706C65735C5C73696D706C655F656D62656464696E675C5C73696D706C655F656D62656464696E675F746573742E6D6C69722073706C6974206174206C696E65202331222C206469726563746F72793A2022443A5C5C4465765C5C6972656522290A2135203D20214449537562726F7574696E65547970652874797065733A202136290A2136203D20217B7D0A2137203D202144494C6F636174696F6E286C696E653A20332C20636F6C756D6E3A20382C2073636F70653A202138290A2138203D202144494C65786963616C426C6F636B46696C652873636F70653A2021332C2066696C653A2021342C206469736372696D696E61746F723A2030290A2139203D2064697374696E63742021444953756270726F6772616D286E616D653A20225F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30222C206C696E6B6167654E616D653A20225F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30222C2073636F70653A206E756C6C2C2066696C653A2021342C206C696E653A20332C20747970653A2021352C2073636F70654C696E653A20332C207370466C6167733A2044495350466C6167446566696E6974696F6E207C2044495350466C61674F7074696D697A65642C20756E69743A2021302C2072657461696E65644E6F6465733A202136290A213130203D202144494C6F636174696F6E286C696E653A20332C20636F6C756D6E3A20382C2073636F70653A20213131290A213131203D202144494C65786963616C426C6F636B46696C652873636F70653A2021392C2066696C653A2021342C206469736372696D696E61746F723A2030290A0100000004000000250000005F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30000000"> : vector<7068xi8>, format = 1280071245 : i32} {
}
hal.executable.binary attributes {data = dense<"0x10000000564D4C41000006000800040006000000040000005403000020000000424D4F44180020001C001800140010000C0000000000000008000400180000008C0200005C0000005000000044000000280000001400000004000000060000006D6F64756C65000003000000500200002C0200000C02000005000000B00100006C01000020010000D80000009C0000000100000058000000010000001400000001000000000000009F0000000500050078FEFFFF080000001400000006FEFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000B4FEFFFF080000001400000042FEFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000F0FEFFFF080000001C0000007EFEFFFF04000000030000000100000001000000010000000C000000766D6C612E6D756C2E6633320000000028FFFFFF0800000024000000B6FEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F7079000000006CFFFFFF080000002800000078FFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E7669657700000000B4FFFFFF0800000020000000C0FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000F4FFFFFF100000003000000008000C000800040008000000080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E6275666665720000060008000400060000000400000003000000693332009F00000008000009100000000100090100000002000902000000030052000000800300008000000000010001805202000080030001C0000001000100018052000000800300008000000200010002805202000080030002C00000010001000280520100008001000100010003805204000080030001C002C0038000005200000080030000C000000300010000805203000080050003C0000000C000000100000054000000"> : vector<880xi8>, format = 1447906369 : i32} {
}
}
func @simple_mul(%arg0: !hal.buffer, %arg1: !hal.buffer) -> !hal.buffer attributes {iree.module.export = "simple_mul$raw"} {
%c0 = constant 0 : index
%c4 = constant 4 : index
%dev = hal.ex.shared_device : !hal.device
%allocator = hal.device.allocator %dev : !hal.allocator
%sz = hal.allocator.compute_size %allocator, shape = [%c4], element_type = 50331680
%buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
hal.ex.defer_release %buffer : !hal.buffer
%cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
hal.command_buffer.begin %cmd
%0 = hal.variable.load @_executable_simple_mul_ex_dispatch_0 : !hal.executable
%1 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
%allocator_0 = hal.buffer.allocator %arg0 : !hal.allocator
%sz_1 = hal.allocator.compute_size %allocator_0, shape = [%c4], element_type = 50331680
%allocator_2 = hal.buffer.allocator %arg1 : !hal.allocator
%sz_3 = hal.allocator.compute_size %allocator_2, shape = [%c4], element_type = 50331680
hal.command_buffer.push_descriptor_set %cmd, %1, set=0, bindings=[0 = (%arg0, %c0, %sz_1), 1 = (%arg1, %c0, %sz_3), 2 = (%buffer, %c0, %sz)]
%2 = hal.variable.load @_device_match_id_0 : i1
cond_br %2, ^bb1, ^bb2
^bb1: // pred: ^bb0
call @simple_mul_switch_0_0(%c4, %cmd, %0) : (index, !hal.command_buffer, !hal.executable) -> ()
br ^bb5
^bb2: // pred: ^bb0
%3 = hal.variable.load @_device_match_id_1 : i1
cond_br %3, ^bb3, ^bb4
^bb3: // pred: ^bb2
call @simple_mul_switch_0_1(%c4, %cmd, %0) : (index, !hal.command_buffer, !hal.executable) -> ()
br ^bb5
^bb4: // pred: ^bb2
br ^bb5
^bb5: // 3 preds: ^bb1, ^bb3, ^bb4
%memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
hal.command_buffer.end %cmd
hal.ex.submit_and_wait %dev, %cmd
return %buffer : !hal.buffer
}
func @simple_mul_switch_0_0(%arg0: index, %arg1: !hal.command_buffer, %arg2: !hal.executable) attributes {sym_visibility = "private"} {
%c1 = constant 1 : index
%0 = addi %arg0, %c1 : index
%1 = subi %0, %c1 : index
%2 = divi_unsigned %1, %c1 : index
%3 = muli %2, %1 : index
%4 = subi %arg0, %3 : index
%5 = cmpi "sge", %4, %c1 : index
%6 = select %5, %4, %c1 : index
%7 = addi %6, %c1 : index
%8 = subi %7, %c1 : index
%9 = divi_unsigned %8, %c1 : index
%10 = muli %9, %8 : index
%11 = subi %6, %10 : index
%12 = cmpi "sge", %11, %c1 : index
%13 = select %12, %11, %c1 : index
%14 = addi %13, %c1 : index
%15 = subi %14, %c1 : index
%16 = divi_unsigned %15, %c1 : index
hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%2, %9, %16]
return
}
func @simple_mul_switch_0_1(%arg0: index, %arg1: !hal.command_buffer, %arg2: !hal.executable) attributes {sym_visibility = "private"} {
%c1 = constant 1 : index
%0 = addi %arg0, %c1 : index
%1 = subi %0, %c1 : index
%2 = divi_unsigned %1, %c1 : index
%3 = muli %2, %1 : index
%4 = subi %arg0, %3 : index
%5 = cmpi "sge", %4, %c1 : index
%6 = select %5, %4, %c1 : index
%7 = addi %6, %c1 : index
%8 = subi %7, %c1 : index
%9 = divi_unsigned %8, %c1 : index
%10 = muli %9, %8 : index
%11 = subi %6, %10 : index
%12 = cmpi "sge", %11, %c1 : index
%13 = select %12, %11, %c1 : index
%14 = addi %13, %c1 : index
%15 = subi %14, %c1 : index
%16 = divi_unsigned %15, %c1 : index
hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%2, %9, %16]
return
}
func @simple_mul$sync(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.module.export = "simple_mul", iree.reflection = {f = "I11!B3!d4B3!d4R6!B3!d4", fv = "1"}} {
%c4 = constant 4 : index
%buffer = hal.buffer_view.buffer %arg0 : !hal.buffer
%buffer_0 = hal.buffer_view.buffer %arg1 : !hal.buffer
%0 = call @simple_mul(%buffer, %buffer_0) : (!hal.buffer, !hal.buffer) -> !hal.buffer
%view = hal.buffer_view.create %0, shape = [%c4], element_type = 50331680 : !hal.buffer_view
return %view : !hal.buffer_view
}
}
Raw
5-vm.mlir
module {
vm.module @module {
vm.rodata @_utf8_llvm_ir_59B551B71BC7CCC0 dense<[108, 108, 118, 109, 45, 105, 114]> : vector<7xi8>
vm.global.i32 @_device_match_id_0 mutable : i32
vm.rodata @_utf8_vmla_EC74E8E47AC10E22 dense<[118, 109, 108, 97]> : vector<4xi8>
vm.global.i32 @_device_match_id_1 mutable : i32
vm.global.ref @_executable_simple_mul_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
vm.rodata @_simple_mul_ex_dispatch_0_binary_llvm dense<"0x100000004C4C564D08000C000400080008000000541B000004000000481B00003B204D6F64756C654944203D20274C4C564D4469616C6563744D6F64756C65270A736F757263655F66696C656E616D65203D20224C4C564D4469616C6563744D6F64756C65220A0A6465636C6172652069382A20406D616C6C6F6328693634290A0A6465636C61726520766F69642040667265652869382A290A0A646566696E6520766F6964204073696D706C655F6D756C5F65785F64697370617463685F3028666C6F61742A2025302C20666C6F61742A2025312C206936342025322C206936342025332C206936342025342C20666C6F61742A2025352C20666C6F61742A2025362C206936342025372C206936342025382C206936342025392C20666C6F61742A202531302C20666C6F61742A202531312C20693634202531322C20693634202531332C2069363420253134292021646267202133207B0A2020253136203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D20756E6465662C20666C6F61742A2025302C20300A2020253137203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20666C6F61742A2025312C20310A2020253138203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531372C206936342025322C20320A2020253139203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531382C206936342025332C20332C20300A2020253230203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531392C206936342025342C20342C20300A2020253231203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D20756E6465662C20666C6F61742A2025352C20300A2020253232203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532312C20666C6F61742A2025362C20310A2020253233203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532322C206936342025372C20320A2020253234203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532332C206936342025382C20332C20300A2020253235203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532342C206936342025392C20342C20300A2020253236203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D20756E6465662C20666C6F61742A202531302C20300A2020253237203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532362C20666C6F61742A202531312C20310A2020253238203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532372C20693634202531322C20320A2020253239203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532382C20693634202531332C20332C20300A2020253330203D20696E7365727476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532392C20693634202531342C20342C20300A20206272206C6162656C202533312C20216462672021370A0A33313A20202020202020202020202020202020202020202020202020202020202020202020202020202020202020202020203B207072656473203D202533342C202531350A2020253332203D2070686920693634205B202534372C20253334205D2C205B20302C20253135205D0A2020253333203D2069636D7020736C7420693634202533322C20342C20216462672021370A20206272206931202533332C206C6162656C202533342C206C6162656C202534382C20216462672021370A0A33343A20202020202020202020202020202020202020202020202020202020202020202020202020202020202020202020203B207072656473203D202533310A2020253335203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532302C20312C20216462672021370A2020253336203D206D756C20693634202533322C20312C20216462672021370A2020253337203D206164642069363420302C202533362C20216462672021370A2020253338203D20676574656C656D656E7470747220666C6F61742C20666C6F61742A202533352C20693634202533372C20216462672021370A2020253339203D206C6F616420666C6F61742C20666C6F61742A202533382C20216462672021370A2020253430203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202532352C20312C20216462672021370A2020253431203D20676574656C656D656E7470747220666C6F61742C20666C6F61742A202534302C20693634202533372C20216462672021370A2020253432203D206C6F616420666C6F61742C20666C6F61742A202534312C20216462672021370A2020253433203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202533302C20312C20216462672021370A2020253434203D20676574656C656D656E7470747220666C6F61742C20666C6F61742A202534332C20693634202533372C20216462672021370A2020253435203D206C6F616420666C6F61742C20666C6F61742A202534342C20216462672021370A2020253436203D20666D756C20666C6F6174202533392C202534322C20216462672021370A202073746F726520666C6F6174202534362C20666C6F61742A202534342C20216462672021370A2020253437203D2061646420693634202533322C20312C20216462672021370A20206272206C6162656C202533312C20216462672021370A0A34383A20202020202020202020202020202020202020202020202020202020202020202020202020202020202020202020203B207072656473203D202533310A202072657420766F69642C20216462672021370A7D0A0A646566696E6520766F696420405F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30287B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025302C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025312C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A202532292021646267202139207B0A20202534203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025302C2021646267202131300A20202535203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20302C2021646267202131300A20202536203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20312C2021646267202131300A20202537203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20322C2021646267202131300A20202538203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20332C20302C2021646267202131300A20202539203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2025342C20342C20302C2021646267202131300A2020253130203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025312C2021646267202131300A2020253131203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20302C2021646267202131300A2020253132203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20312C2021646267202131300A2020253133203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20322C2021646267202131300A2020253134203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20332C20302C2021646267202131300A2020253135203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531302C20342C20302C2021646267202131300A2020253136203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025322C2021646267202131300A2020253137203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20302C2021646267202131300A2020253138203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20312C2021646267202131300A2020253139203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20322C2021646267202131300A2020253230203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20332C20302C2021646267202131300A2020253231203D206578747261637476616C7565207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D202531362C20342C20302C2021646267202131300A202063616C6C20766F6964204073696D706C655F6D756C5F65785F64697370617463685F3028666C6F61742A2025352C20666C6F61742A2025362C206936342025372C206936342025382C206936342025392C20666C6F61742A202531312C20666C6F61742A202531322C20693634202531332C20693634202531342C20693634202531352C20666C6F61742A202531372C20666C6F61742A202531382C20693634202531392C20693634202532302C2069363420253231292C2021646267202131300A202072657420766F69642C2021646267202131300A7D0A0A646566696E6520766F69642040696E766F6B655F5F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F302869382A2A20253029207B0A20202532203D20676574656C656D656E747074722069382A2C2069382A2A2025302C2069363420300A20202533203D206C6F61642069382A2C2069382A2A2025320A20202534203D20626974636173742069382A20253320746F207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A0A20202535203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A2025340A20202536203D20676574656C656D656E747074722069382A2C2069382A2A2025302C2069363420310A20202537203D206C6F61642069382A2C2069382A2A2025360A20202538203D20626974636173742069382A20253720746F207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A0A20202539203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A2025380A2020253130203D20676574656C656D656E747074722069382A2C2069382A2A2025302C2069363420320A2020253131203D206C6F61642069382A2C2069382A2A202531300A2020253132203D20626974636173742069382A2025313120746F207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A0A2020253133203D206C6F6164207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2A202531320A202063616C6C20766F696420405F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30287B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025352C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A2025392C207B20666C6F61742A2C20666C6F61742A2C206936342C205B312078206936345D2C205B312078206936345D207D2A20253133290A202072657420766F69640A7D0A0A216C6C766D2E6462672E6375203D20217B21307D0A216C6C766D2E6D6F64756C652E666C616773203D20217B21327D0A0A2130203D2064697374696E637420214449436F6D70696C65556E6974286C616E67756167653A2044575F4C414E475F432C2066696C653A2021312C2070726F64756365723A20226D6C6972222C2069734F7074696D697A65643A20747275652C2072756E74696D6556657273696F6E3A20302C20656D697373696F6E4B696E643A2046756C6C4465627567290A2131203D2021444946696C652866696C656E616D653A20224C4C564D4469616C6563744D6F64756C65222C206469726563746F72793A20222F22290A2132203D20217B69333220322C202122446562756720496E666F2056657273696F6E222C2069333220337D0A2133203D2064697374696E63742021444953756270726F6772616D286E616D653A202273696D706C655F6D756C5F65785F64697370617463685F30222C206C696E6B6167654E616D653A202273696D706C655F6D756C5F65785F64697370617463685F30222C2073636F70653A206E756C6C2C2066696C653A2021342C206C696E653A20332C20747970653A2021352C2073636F70654C696E653A20332C207370466C6167733A2044495350466C6167446566696E6974696F6E207C2044495350466C61674F7074696D697A65642C20756E69743A2021302C2072657461696E65644E6F6465733A202136290A2134203D2021444946696C652866696C656E616D653A2022697265655C5C73616D706C65735C5C73696D706C655F656D62656464696E675C5C73696D706C655F656D62656464696E675F746573742E6D6C69722073706C6974206174206C696E65202331222C206469726563746F72793A2022443A5C5C4465765C5C6972656522290A2135203D20214449537562726F7574696E65547970652874797065733A202136290A2136203D20217B7D0A2137203D202144494C6F636174696F6E286C696E653A20332C20636F6C756D6E3A20382C2073636F70653A202138290A2138203D202144494C65786963616C426C6F636B46696C652873636F70653A2021332C2066696C653A2021342C206469736372696D696E61746F723A2030290A2139203D2064697374696E63742021444953756270726F6772616D286E616D653A20225F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30222C206C696E6B6167654E616D653A20225F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30222C2073636F70653A206E756C6C2C2066696C653A2021342C206C696E653A20332C20747970653A2021352C2073636F70654C696E653A20332C207370466C6167733A2044495350466C6167446566696E6974696F6E207C2044495350466C61674F7074696D697A65642C20756E69743A2021302C2072657461696E65644E6F6465733A202136290A213130203D202144494C6F636174696F6E286C696E653A20332C20636F6C756D6E3A20382C2073636F70653A20213131290A213131203D202144494C65786963616C426C6F636B46696C652873636F70653A2021392C2066696C653A2021342C206469736372696D696E61746F723A2030290A0100000004000000250000005F6D6C69725F6369666163655F73696D706C655F6D756C5F65785F64697370617463685F30000000"> : vector<7068xi8>
vm.rodata @_simple_mul_ex_dispatch_0_binary_vmla dense<"0x10000000564D4C41000006000800040006000000040000005403000020000000424D4F44180020001C001800140010000C0000000000000008000400180000008C0200005C0000005000000044000000280000001400000004000000060000006D6F64756C65000003000000500200002C0200000C02000005000000B00100006C01000020010000D80000009C0000000100000058000000010000001400000001000000000000009F0000000500050078FEFFFF080000001400000006FEFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000B4FEFFFF080000001400000042FEFFFF0400000001000000020000001800000073696D706C655F6D756C5F65785F64697370617463685F3000000000F0FEFFFF080000001C0000007EFEFFFF04000000030000000100000001000000010000000C000000766D6C612E6D756C2E6633320000000028FFFFFF0800000024000000B6FEFFFF0400000005000000010000000000000001000000000000000000000010000000766D6C612E6275666665722E636F7079000000006CFFFFFF080000002800000078FFFFFF080000000C00000001000000010000000300000001000000000000000000000010000000766D6C612E6275666665722E7669657700000000B4FFFFFF0800000020000000C0FFFFFF080000000C0000000100000001000000010000000000000011000000766D6C612E6275666665722E616C6C6F63000000F4FFFFFF100000003000000008000C000800040008000000080000000C00000001000000010000000300000002000000000000000000000016000000766D6C612E696E746572666163652E62696E64696E670000CAFFFFFF040000000F00000021766D6C612E696E7465726661636500E6FFFFFF040000000C00000021766D6C612E6275666665720000060008000400060000000400000003000000693332009F00000008000009100000000100090100000002000902000000030052000000800300008000000000010001805202000080030001C0000001000100018052000000800300008000000200010002805202000080030002C00000010001000280520100008001000100010003805204000080030001C002C0038000005200000080030000C000000300010000805203000080050003C0000000C000000100000054000000"> : vector<880xi8>
vm.func @simple_mul(%arg0: !vm.ref<!hal.buffer>, %arg1: !vm.ref<!hal.buffer>) -> !vm.ref<!hal.buffer> {
%c50 = vm.const.i32 50 : i32
%c15 = vm.const.i32 15 : i32
%c3 = vm.const.i32 3 : i32
%c50331680 = vm.const.i32 50331680 : i32
%c2 = vm.const.i32 2 : i32
%c1 = vm.const.i32 1 : i32
%zero = vm.const.i32.zero : i32
%c20 = vm.const.i32 20 : i32
%c5 = vm.const.i32 5 : i32
%c8 = vm.const.i32 8 : i32
%c4 = vm.const.i32 4 : i32
%ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
%0 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [%c4], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
%ref_1 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %0) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
%ref_2 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
vm.call @hal.command_buffer.begin(%ref_2) : (!vm.ref<!hal.command_buffer>) -> ()
%_executable_simple_mul_ex_dispatch_0 = vm.global.load.ref @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable>
%_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
%ref_3 = vm.call @hal.buffer.allocator(%arg0) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
%1 = vm.call.variadic @hal.allocator.compute_size(%ref_3, [%c4], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
%ref_4 = vm.call @hal.buffer.allocator(%arg1) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
%2 = vm.call.variadic @hal.allocator.compute_size(%ref_4, [%c4], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_2, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%arg0, %arg1, %ref_1], [%zero, %zero, %zero], [%1, %2, %0]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
%_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
vm.cond_br %_device_match_id_0, ^bb1, ^bb2
^bb1: // pred: ^bb0
vm.call @hal.command_buffer.dispatch(%ref_2, %_executable_simple_mul_ex_dispatch_0, %zero, %c4, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
vm.br ^bb5
^bb2: // pred: ^bb0
%_device_match_id_1 = vm.global.load.i32 @_device_match_id_1 : i32
vm.cond_br %_device_match_id_1, ^bb3, ^bb4
^bb3: // pred: ^bb2
vm.call @hal.command_buffer.dispatch(%ref_2, %_executable_simple_mul_ex_dispatch_0, %zero, %c4, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
vm.br ^bb5
^bb4: // pred: ^bb2
vm.br ^bb5
^bb5: // 3 preds: ^bb1, ^bb3, ^bb4
vm.call.variadic @hal.command_buffer.execution_barrier(%ref_2, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
vm.call @hal.command_buffer.end(%ref_2) : (!vm.ref<!hal.command_buffer>) -> ()
vm.call @hal.ex.submit_and_wait(%ref, %ref_2) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
vm.return %ref_1 : !vm.ref<!hal.buffer>
}
vm.export @simple_mul as("simple_mul$raw")
vm.func @simple_mul$sync(%arg0: !vm.ref<!hal.buffer_view>, %arg1: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view> attributes {iree.reflection = {f = "I11!B3!d4B3!d4R6!B3!d4", fv = "1"}} {
%c50 = vm.const.i32 50 : i32
%c15 = vm.const.i32 15 : i32
%c3 = vm.const.i32 3 : i32
%c2 = vm.const.i32 2 : i32
%c1 = vm.const.i32 1 : i32
%zero = vm.const.i32.zero : i32
%c20 = vm.const.i32 20 : i32
%c5 = vm.const.i32 5 : i32
%c8 = vm.const.i32 8 : i32
%c4 = vm.const.i32 4 : i32
%c50331680 = vm.const.i32 50331680 : i32
%ref = vm.call @hal.buffer_view.buffer(%arg0) : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer>
%ref_0 = vm.call @hal.buffer_view.buffer(%arg1) : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer>
%ref_1 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%ref_2 = vm.call @hal.device.allocator(%ref_1) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
%0 = vm.call.variadic @hal.allocator.compute_size(%ref_2, [%c4], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
%ref_3 = vm.call @hal.allocator.allocate(%ref_2, %c50, %c15, %0) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
vm.call @hal.ex.defer_release(%ref_3) : (!vm.ref<!hal.buffer>) -> ()
%ref_4 = vm.call @hal.command_buffer.create(%ref_1, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
vm.call @hal.command_buffer.begin(%ref_4) : (!vm.ref<!hal.command_buffer>) -> ()
%_executable_simple_mul_ex_dispatch_0 = vm.global.load.ref @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable>
%_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
%ref_5 = vm.call @hal.buffer.allocator(%ref) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
%1 = vm.call.variadic @hal.allocator.compute_size(%ref_5, [%c4], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
%ref_6 = vm.call @hal.buffer.allocator(%ref_0) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
%2 = vm.call.variadic @hal.allocator.compute_size(%ref_6, [%c4], %c50331680) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_4, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%ref, %ref_0, %ref_3], [%zero, %zero, %zero], [%1, %2, %0]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
%_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
vm.cond_br %_device_match_id_0, ^bb1, ^bb2
^bb1: // pred: ^bb0
vm.call @hal.command_buffer.dispatch(%ref_4, %_executable_simple_mul_ex_dispatch_0, %zero, %c4, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
vm.br ^bb5
^bb2: // pred: ^bb0
%_device_match_id_1 = vm.global.load.i32 @_device_match_id_1 : i32
vm.cond_br %_device_match_id_1, ^bb3, ^bb4
^bb3: // pred: ^bb2
vm.call @hal.command_buffer.dispatch(%ref_4, %_executable_simple_mul_ex_dispatch_0, %zero, %c4, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
vm.br ^bb5
^bb4: // pred: ^bb2
vm.br ^bb5
^bb5: // 3 preds: ^bb1, ^bb3, ^bb4
vm.call.variadic @hal.command_buffer.execution_barrier(%ref_4, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
vm.call @hal.command_buffer.end(%ref_4) : (!vm.ref<!hal.command_buffer>) -> ()
vm.call @hal.ex.submit_and_wait(%ref_1, %ref_4) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
vm.br ^bb6(%ref_3 : !vm.ref<!hal.buffer>)
^bb6(%3: !vm.ref<!hal.buffer>): // pred: ^bb5
%ref_7 = vm.call.variadic @hal.buffer_view.create(%3, [%c4], %c50331680) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
vm.return %ref_7 : !vm.ref<!hal.buffer_view>
}
vm.export @simple_mul$sync as("simple_mul")
vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
vm.func @__init() {
%c1 = vm.const.i32 1 : i32
%c2 = vm.const.i32 2 : i32
%c6 = vm.const.i32 6 : i32
%zero = vm.const.i32.zero : i32
%c1280071245 = vm.const.i32 1280071245 : i32
%c1447906369 = vm.const.i32 1447906369 : i32
%null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
%c7 = vm.const.i32 7 : i32
%ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%_utf8_llvm_ir_59B551B71BC7CCC0 = vm.const.ref.rodata @_utf8_llvm_ir_59B551B71BC7CCC0 : !vm.ref<!iree.byte_buffer>
%0 = vm.call @hal.device.match.id(%ref, %_utf8_llvm_ir_59B551B71BC7CCC0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
vm.global.store.i32 %0, @_device_match_id_0 : i32
%ref_0 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%_utf8_vmla_EC74E8E47AC10E22 = vm.const.ref.rodata @_utf8_vmla_EC74E8E47AC10E22 : !vm.ref<!iree.byte_buffer>
%1 = vm.call @hal.device.match.id(%ref_0, %_utf8_vmla_EC74E8E47AC10E22) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
vm.global.store.i32 %1, @_device_match_id_1 : i32
%ref_1 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%ref_2 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_1, %c1, [%zero, %c7, %c1, %c1, %c7, %c1, %c2, %c7, %c6]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
vm.global.store.ref %ref_2, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
%_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
%ref_3 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%ref_4 = vm.call.variadic @hal.executable_layout.create(%ref_3, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
vm.global.store.ref %ref_4, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
%ref_5 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
%_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
%ref_6 = vm.call @hal.executable_cache.create(%ref_5, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
%_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
%2 = vm.call.variadic @hal.executable_cache.select_format(%ref_6, [%c1280071245, %c1447906369]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
%_simple_mul_ex_dispatch_0_binary_llvm = vm.const.ref.rodata @_simple_mul_ex_dispatch_0_binary_llvm : !vm.ref<!iree.byte_buffer>
%_simple_mul_ex_dispatch_0_binary_vmla = vm.const.ref.rodata @_simple_mul_ex_dispatch_0_binary_vmla : !vm.ref<!iree.byte_buffer>
%ref_7 = vm.switch.ref %2[%_simple_mul_ex_dispatch_0_binary_llvm, %_simple_mul_ex_dispatch_0_binary_vmla] else %null : !vm.ref<!iree.byte_buffer>
%ref_8 = vm.call @hal.executable_cache.prepare(%ref_6, %_executable_layout_0, %c7, %ref_7) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
vm.global.store.ref %ref_8, @_executable_simple_mul_ex_dispatch_0 : !vm.ref<!hal.executable>
vm.global.store.ref %ref_6, @_executable_cache : !vm.ref<!hal.executable_cache>
vm.return
}
vm.export @__init
}
}
Raw
6-module.txt
{
name: "module",
types: [ {
full_name: "i32"
}, {
full_name: "!hal.allocator"
}, {
full_name: "!hal.buffer"
}, {
full_name: "!hal.buffer_view"
}, {
full_name: "!hal.command_buffer"
}, {
full_name: "!hal.descriptor_set_layout"
}, {
full_name: "!hal.device"
}, {
full_name: "!hal.executable"
}, {
full_name: "!hal.executable_cache"
}, {
full_name: "!hal.executable_layout"
}, {
full_name: "!iree.byte_buffer"
} ],
imported_functions: [ {
full_name: "hal.ex.shared_device",
signature: {
result_types: [ 6 ]
}
}, {
full_name: "hal.ex.defer_release",
signature: {
argument_types: [ 0 ]
}
}, {
full_name: "hal.ex.submit_and_wait",
signature: {
argument_types: [ 6, 4 ]
}
}, {
full_name: "hal.allocator.compute_size",
signature: {
argument_types: [ 1, 0, 0 ],
result_types: [ 0 ]
}
}, {
full_name: "hal.allocator.allocate",
signature: {
argument_types: [ 1, 0, 0, 0 ],
result_types: [ 2 ]
}
}, {
full_name: "hal.buffer.allocator",
signature: {
argument_types: [ 2 ],
result_types: [ 1 ]
}
}, {
full_name: "hal.buffer_view.create",
signature: {
argument_types: [ 2, 0, 0 ],
result_types: [ 3 ]
}
}, {
full_name: "hal.buffer_view.buffer",
signature: {
argument_types: [ 3 ],
result_types: [ 2 ]
}
}, {
full_name: "hal.command_buffer.create",
signature: {
argument_types: [ 6, 0, 0 ],
result_types: [ 4 ]
}
}, {
full_name: "hal.command_buffer.begin",
signature: {
argument_types: [ 4 ]
}
}, {
full_name: "hal.command_buffer.end",
signature: {
argument_types: [ 4 ]
}
}, {
full_name: "hal.command_buffer.execution_barrier",
signature: {
argument_types: [ 4, 0, 0, 0, 0 ]
}
}, {
full_name: "hal.command_buffer.push_descriptor_set",
signature: {
argument_types: [ 4, 9, 0, 0, 2, 0, 0 ]
}
}, {
full_name: "hal.command_buffer.dispatch",
signature: {
argument_types: [ 4, 7, 0, 0, 0, 0 ]
}
}, {
full_name: "hal.descriptor_set_layout.create",
signature: {
argument_types: [ 6, 0, 0 ],
result_types: [ 5 ]
}
}, {
full_name: "hal.device.allocator",
signature: {
argument_types: [ 6 ],
result_types: [ 1 ]
}
}, {
full_name: "hal.device.match.id",
signature: {
argument_types: [ 6, 10 ],
result_types: [ 0 ]
}
}, {
full_name: "hal.executable_cache.create",
signature: {
argument_types: [ 6, 10 ],
result_types: [ 8 ]
}
}, {
full_name: "hal.executable_cache.select_format",
signature: {
argument_types: [ 8, 0 ],
result_types: [ 0 ]
}
}, {
full_name: "hal.executable_cache.prepare",
signature: {
argument_types: [ 8, 9, 0, 10 ],
result_types: [ 7 ]
}
}, {
full_name: "hal.executable_layout.create",
signature: {
argument_types: [ 6, 5, 0 ],
result_types: [ 9 ]
}
} ],
exported_functions: [ {
local_name: "simple_mul$raw",
signature: {
argument_types: [ 2, 2 ],
result_types: [ 2 ]
}
}, {
local_name: "simple_mul",
signature: {
argument_types: [ 3, 3 ],
result_types: [ 3 ]
},
internal_ordinal: 1
}, {
local_name: "__init",
signature: {
},
internal_ordinal: 2
} ],
internal_functions: [ {
local_name: "simple_mul",
signature: {
argument_types: [ 2, 2 ],
result_types: [ 2 ]
}
}, {
local_name: "simple_mul$sync",
signature: {
argument_types: [ 3, 3 ],
result_types: [ 3 ],
reflection_attrs: [ {
key: "f",
value: "I11!B3!d4B3!d4R6!B3!d4"
}, {
key: "fv",
value: "1"
} ]
}
}, {
local_name: "__init",
signature: {
}
} ],
rodata_segments: [ {
data: [ 108, 108, 118, 109, 45, 105, 114 ]
}, {
data: [ 118, 109, 108, 97 ]
}, {
data: [ 100, 101, 102, 97, 117, 108, 116 ]
}, {
data: [ 16, 0, 0, 0, 76, 76, 86, 77, 8, 0, 12, 0, 4, 0, 8, 0, 8, 0, 0, 0, 84, 27, 0, 0, 4, 0, 0, 0, 72, 27, 0, 0, 59, 32, 77, 111, 100, 117, 108, 101, 73, 68, 32, 61, 32, 39, 76, 76, 86, 77, 68, 105, 97, 108, 101, 99, 116, 77, 111, 100, 117, 108, 101, 39, 10, 115, 111, 117, 114, 99, 101, 95, 102, 105, 108, 101, 110, 97, 109, 101, 32, 61, 32, 34, 76, 76, 86, 77, 68, 105, 97, 108, 101, 99, 116, 77, 111, 100, 117, 108, 101, 34, 10, 10, 100, 101, 99, 108, 97, 114, 101, 32, 105, 56, 42, 32, 64, 109, 97, 108, 108, 111, 99, 40, 105, 54, 52, 41, 10, 10, 100, 101, 99, 108, 97, 114, 101, 32, 118, 111, 105, 100, 32, 64, 102, 114, 101, 101, 40, 105, 56, 42, 41, 10, 10, 100, 101, 102, 105, 110, 101, 32, 118, 111, 105, 100, 32, 64, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 40, 102, 108, 111, 97, 116, 42, 32, 37, 48, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 44, 32, 105, 54, 52, 32, 37, 50, 44, 32, 105, 54, 52, 32, 37, 51, 44, 32, 105, 54, 52, 32, 37, 52, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 53, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 54, 44, 32, 105, 54, 52, 32, 37, 55, 44, 32, 105, 54, 52, 32, 37, 56, 44, 32, 105, 54, 52, 32, 37, 57, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 48, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 49, 44, 32, 105, 54, 52, 32, 37, 49, 50, 44, 32, 105, 54, 52, 32, 37, 49, 51, 44, 32, 105, 54, 52, 32, 37, 49, 52, 41, 32, 33, 100, 98, 103, 32, 33, 51, 32, 123, 10, 32, 32, 37, 49, 54, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 117, 110, 100, 101, 102, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 48, 44, 32, 48, 10, 32, 32, 37, 49, 55, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 54, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 44, 32, 49, 10, 32, 32, 37, 49, 56, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 55, 44, 32, 105, 54, 52, 32, 37, 50, 44, 32, 50, 10, 32, 32, 37, 49, 57, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 56, 44, 32, 105, 54, 52, 32, 37, 51, 44, 32, 51, 44, 32, 48, 10, 32, 32, 37, 50, 48, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 57, 44, 32, 105, 54, 52, 32, 37, 52, 44, 32, 52, 44, 32, 48, 10, 32, 32, 37, 50, 49, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 117, 110, 100, 101, 102, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 53, 44, 32, 48, 10, 32, 32, 37, 50, 50, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 49, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 54, 44, 32, 49, 10, 32, 32, 37, 50, 51, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 50, 44, 32, 105, 54, 52, 32, 37, 55, 44, 32, 50, 10, 32, 32, 37, 50, 52, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 51, 44, 32, 105, 54, 52, 32, 37, 56, 44, 32, 51, 44, 32, 48, 10, 32, 32, 37, 50, 53, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 52, 44, 32, 105, 54, 52, 32, 37, 57, 44, 32, 52, 44, 32, 48, 10, 32, 32, 37, 50, 54, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 117, 110, 100, 101, 102, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 48, 44, 32, 48, 10, 32, 32, 37, 50, 55, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 54, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 49, 44, 32, 49, 10, 32, 32, 37, 50, 56, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 55, 44, 32, 105, 54, 52, 32, 37, 49, 50, 44, 32, 50, 10, 32, 32, 37, 50, 57, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 56, 44, 32, 105, 54, 52, 32, 37, 49, 51, 44, 32, 51, 44, 32, 48, 10, 32, 32, 37, 51, 48, 32, 61, 32, 105, 110, 115, 101, 114, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 57, 44, 32, 105, 54, 52, 32, 37, 49, 52, 44, 32, 52, 44, 32, 48, 10, 32, 32, 98, 114, 32, 108, 97, 98, 101, 108, 32, 37, 51, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 10, 51, 49, 58, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 59, 32, 112, 114, 101, 100, 115, 32, 61, 32, 37, 51, 52, 44, 32, 37, 49, 53, 10, 32, 32, 37, 51, 50, 32, 61, 32, 112, 104, 105, 32, 105, 54, 52, 32, 91, 32, 37, 52, 55, 44, 32, 37, 51, 52, 32, 93, 44, 32, 91, 32, 48, 44, 32, 37, 49, 53, 32, 93, 10, 32, 32, 37, 51, 51, 32, 61, 32, 105, 99, 109, 112, 32, 115, 108, 116, 32, 105, 54, 52, 32, 37, 51, 50, 44, 32, 52, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 98, 114, 32, 105, 49, 32, 37, 51, 51, 44, 32, 108, 97, 98, 101, 108, 32, 37, 51, 52, 44, 32, 108, 97, 98, 101, 108, 32, 37, 52, 56, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 10, 51, 52, 58, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 59, 32, 112, 114, 101, 100, 115, 32, 61, 32, 37, 51, 49, 10, 32, 32, 37, 51, 53, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 48, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 51, 54, 32, 61, 32, 109, 117, 108, 32, 105, 54, 52, 32, 37, 51, 50, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 51, 55, 32, 61, 32, 97, 100, 100, 32, 105, 54, 52, 32, 48, 44, 32, 37, 51, 54, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 51, 56, 32, 61, 32, 103, 101, 116, 101, 108, 101, 109, 101, 110, 116, 112, 116, 114, 32, 102, 108, 111, 97, 116, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 51, 53, 44, 32, 105, 54, 52, 32, 37, 51, 55, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 51, 57, 32, 61, 32, 108, 111, 97, 100, 32, 102, 108, 111, 97, 116, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 51, 56, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 48, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 50, 53, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 49, 32, 61, 32, 103, 101, 116, 101, 108, 101, 109, 101, 110, 116, 112, 116, 114, 32, 102, 108, 111, 97, 116, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 52, 48, 44, 32, 105, 54, 52, 32, 37, 51, 55, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 50, 32, 61, 32, 108, 111, 97, 100, 32, 102, 108, 111, 97, 116, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 52, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 51, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 51, 48, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 52, 32, 61, 32, 103, 101, 116, 101, 108, 101, 109, 101, 110, 116, 112, 116, 114, 32, 102, 108, 111, 97, 116, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 52, 51, 44, 32, 105, 54, 52, 32, 37, 51, 55, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 53, 32, 61, 32, 108, 111, 97, 100, 32, 102, 108, 111, 97, 116, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 52, 52, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 54, 32, 61, 32, 102, 109, 117, 108, 32, 102, 108, 111, 97, 116, 32, 37, 51, 57, 44, 32, 37, 52, 50, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 115, 116, 111, 114, 101, 32, 102, 108, 111, 97, 116, 32, 37, 52, 54, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 52, 52, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 37, 52, 55, 32, 61, 32, 97, 100, 100, 32, 105, 54, 52, 32, 37, 51, 50, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 32, 32, 98, 114, 32, 108, 97, 98, 101, 108, 32, 37, 51, 49, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 10, 52, 56, 58, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 32, 59, 32, 112, 114, 101, 100, 115, 32, 61, 32, 37, 51, 49, 10, 32, 32, 114, 101, 116, 32, 118, 111, 105, 100, 44, 32, 33, 100, 98, 103, 32, 33, 55, 10, 125, 10, 10, 100, 101, 102, 105, 110, 101, 32, 118, 111, 105, 100, 32, 64, 95, 109, 108, 105, 114, 95, 99, 105, 102, 97, 99, 101, 95, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 40, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 48, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 49, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 50, 41, 32, 33, 100, 98, 103, 32, 33, 57, 32, 123, 10, 32, 32, 37, 52, 32, 61, 32, 108, 111, 97, 100, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 53, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 52, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 54, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 52, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 55, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 52, 44, 32, 50, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 56, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 52, 44, 32, 51, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 57, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 52, 44, 32, 52, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 48, 32, 61, 32, 108, 111, 97, 100, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 49, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 49, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 48, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 50, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 48, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 51, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 48, 44, 32, 50, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 52, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 48, 44, 32, 51, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 53, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 48, 44, 32, 52, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 54, 32, 61, 32, 108, 111, 97, 100, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 50, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 55, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 54, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 56, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 54, 44, 32, 49, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 49, 57, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 54, 44, 32, 50, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 50, 48, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 54, 44, 32, 51, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 37, 50, 49, 32, 61, 32, 101, 120, 116, 114, 97, 99, 116, 118, 97, 108, 117, 101, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 32, 37, 49, 54, 44, 32, 52, 44, 32, 48, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 99, 97, 108, 108, 32, 118, 111, 105, 100, 32, 64, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 40, 102, 108, 111, 97, 116, 42, 32, 37, 53, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 54, 44, 32, 105, 54, 52, 32, 37, 55, 44, 32, 105, 54, 52, 32, 37, 56, 44, 32, 105, 54, 52, 32, 37, 57, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 49, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 50, 44, 32, 105, 54, 52, 32, 37, 49, 51, 44, 32, 105, 54, 52, 32, 37, 49, 52, 44, 32, 105, 54, 52, 32, 37, 49, 53, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 55, 44, 32, 102, 108, 111, 97, 116, 42, 32, 37, 49, 56, 44, 32, 105, 54, 52, 32, 37, 49, 57, 44, 32, 105, 54, 52, 32, 37, 50, 48, 44, 32, 105, 54, 52, 32, 37, 50, 49, 41, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 32, 32, 114, 101, 116, 32, 118, 111, 105, 100, 44, 32, 33, 100, 98, 103, 32, 33, 49, 48, 10, 125, 10, 10, 100, 101, 102, 105, 110, 101, 32, 118, 111, 105, 100, 32, 64, 105, 110, 118, 111, 107, 101, 95, 95, 109, 108, 105, 114, 95, 99, 105, 102, 97, 99, 101, 95, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 40, 105, 56, 42, 42, 32, 37, 48, 41, 32, 123, 10, 32, 32, 37, 50, 32, 61, 32, 103, 101, 116, 101, 108, 101, 109, 101, 110, 116, 112, 116, 114, 32, 105, 56, 42, 44, 32, 105, 56, 42, 42, 32, 37, 48, 44, 32, 105, 54, 52, 32, 48, 10, 32, 32, 37, 51, 32, 61, 32, 108, 111, 97, 100, 32, 105, 56, 42, 44, 32, 105, 56, 42, 42, 32, 37, 50, 10, 32, 32, 37, 52, 32, 61, 32, 98, 105, 116, 99, 97, 115, 116, 32, 105, 56, 42, 32, 37, 51, 32, 116, 111, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 42, 10, 32, 32, 37, 53, 32, 61, 32, 108, 111, 97, 100, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 42, 32, 37, 52, 10, 32, 32, 37, 54, 32, 61, 32, 103, 101, 116, 101, 108, 101, 109, 101, 110, 116, 112, 116, 114, 32, 105, 56, 42, 44, 32, 105, 56, 42, 42, 32, 37, 48, 44, 32, 105, 54, 52, 32, 49, 10, 32, 32, 37, 55, 32, 61, 32, 108, 111, 97, 100, 32, 105, 56, 42, 44, 32, 105, 56, 42, 42, 32, 37, 54, 10, 32, 32, 37, 56, 32, 61, 32, 98, 105, 116, 99, 97, 115, 116, 32, 105, 56, 42, 32, 37, 55, 32, 116, 111, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 42, 10, 32, 32, 37, 57, 32, 61, 32, 108, 111, 97, 100, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 42, 32, 37, 56, 10, 32, 32, 37, 49, 48, 32, 61, 32, 103, 101, 116, 101, 108, 101, 109, 101, 110, 116, 112, 116, 114, 32, 105, 56, 42, 44, 32, 105, 56, 42, 42, 32, 37, 48, 44, 32, 105, 54, 52, 32, 50, 10, 32, 32, 37, 49, 49, 32, 61, 32, 108, 111, 97, 100, 32, 105, 56, 42, 44, 32, 105, 56, 42, 42, 32, 37, 49, 48, 10, 32, 32, 37, 49, 50, 32, 61, 32, 98, 105, 116, 99, 97, 115, 116, 32, 105, 56, 42, 32, 37, 49, 49, 32, 116, 111, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 42, 10, 32, 32, 37, 49, 51, 32, 61, 32, 108, 111, 97, 100, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 42, 32, 37, 49, 50, 10, 32, 32, 99, 97, 108, 108, 32, 118, 111, 105, 100, 32, 64, 95, 109, 108, 105, 114, 95, 99, 105, 102, 97, 99, 101, 95, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 40, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 53, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 57, 44, 32, 123, 32, 102, 108, 111, 97, 116, 42, 44, 32, 102, 108, 111, 97, 116, 42, 44, 32, 105, 54, 52, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 44, 32, 91, 49, 32, 120, 32, 105, 54, 52, 93, 32, 125, 42, 32, 37, 49, 51, 41, 10, 32, 32, 114, 101, 116, 32, 118, 111, 105, 100, 10, 125, 10, 10, 33, 108, 108, 118, 109, 46, 100, 98, 103, 46, 99, 117, 32, 61, 32, 33, 123, 33, 48, 125, 10, 33, 108, 108, 118, 109, 46, 109, 111, 100, 117, 108, 101, 46, 102, 108, 97, 103, 115, 32, 61, 32, 33, 123, 33, 50, 125, 10, 10, 33, 48, 32, 61, 32, 100, 105, 115, 116, 105, 110, 99, 116, 32, 33, 68, 73, 67, 111, 109, 112, 105, 108, 101, 85, 110, 105, 116, 40, 108, 97, 110, 103, 117, 97, 103, 101, 58, 32, 68, 87, 95, 76, 65, 78, 71, 95, 67, 44, 32, 102, 105, 108, 101, 58, 32, 33, 49, 44, 32, 112, 114, 111, 100, 117, 99, 101, 114, 58, 32, 34, 109, 108, 105, 114, 34, 44, 32, 105, 115, 79, 112, 116, 105, 109, 105, 122, 101, 100, 58, 32, 116, 114, 117, 101, 44, 32, 114, 117, 110, 116, 105, 109, 101, 86, 101, 114, 115, 105, 111, 110, 58, 32, 48, 44, 32, 101, 109, 105, 115, 115, 105, 111, 110, 75, 105, 110, 100, 58, 32, 70, 117, 108, 108, 68, 101, 98, 117, 103, 41, 10, 33, 49, 32, 61, 32, 33, 68, 73, 70, 105, 108, 101, 40, 102, 105, 108, 101, 110, 97, 109, 101, 58, 32, 34, 76, 76, 86, 77, 68, 105, 97, 108, 101, 99, 116, 77, 111, 100, 117, 108, 101, 34, 44, 32, 100, 105, 114, 101, 99, 116, 111, 114, 121, 58, 32, 34, 47, 34, 41, 10, 33, 50, 32, 61, 32, 33, 123, 105, 51, 50, 32, 50, 44, 32, 33, 34, 68, 101, 98, 117, 103, 32, 73, 110, 102, 111, 32, 86, 101, 114, 115, 105, 111, 110, 34, 44, 32, 105, 51, 50, 32, 51, 125, 10, 33, 51, 32, 61, 32, 100, 105, 115, 116, 105, 110, 99, 116, 32, 33, 68, 73, 83, 117, 98, 112, 114, 111, 103, 114, 97, 109, 40, 110, 97, 109, 101, 58, 32, 34, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 34, 44, 32, 108, 105, 110, 107, 97, 103, 101, 78, 97, 109, 101, 58, 32, 34, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 34, 44, 32, 115, 99, 111, 112, 101, 58, 32, 110, 117, 108, 108, 44, 32, 102, 105, 108, 101, 58, 32, 33, 52, 44, 32, 108, 105, 110, 101, 58, 32, 51, 44, 32, 116, 121, 112, 101, 58, 32, 33, 53, 44, 32, 115, 99, 111, 112, 101, 76, 105, 110, 101, 58, 32, 51, 44, 32, 115, 112, 70, 108, 97, 103, 115, 58, 32, 68, 73, 83, 80, 70, 108, 97, 103, 68, 101, 102, 105, 110, 105, 116, 105, 111, 110, 32, 124, 32, 68, 73, 83, 80, 70, 108, 97, 103, 79, 112, 116, 105, 109, 105, 122, 101, 100, 44, 32, 117, 110, 105, 116, 58, 32, 33, 48, 44, 32, 114, 101, 116, 97, 105, 110, 101, 100, 78, 111, 100, 101, 115, 58, 32, 33, 54, 41, 10, 33, 52, 32, 61, 32, 33, 68, 73, 70, 105, 108, 101, 40, 102, 105, 108, 101, 110, 97, 109, 101, 58, 32, 34, 105, 114, 101, 101, 92, 92, 115, 97, 109, 112, 108, 101, 115, 92, 92, 115, 105, 109, 112, 108, 101, 95, 101, 109, 98, 101, 100, 100, 105, 110, 103, 92, 92, 115, 105, 109, 112, 108, 101, 95, 101, 109, 98, 101, 100, 100, 105, 110, 103, 95, 116, 101, 115, 116, 46, 109, 108, 105, 114, 32, 115, 112, 108, 105, 116, 32, 97, 116, 32, 108, 105, 110, 101, 32, 35, 49, 34, 44, 32, 100, 105, 114, 101, 99, 116, 111, 114, 121, 58, 32, 34, 68, 58, 92, 92, 68, 101, 118, 92, 92, 105, 114, 101, 101, 34, 41, 10, 33, 53, 32, 61, 32, 33, 68, 73, 83, 117, 98, 114, 111, 117, 116, 105, 110, 101, 84, 121, 112, 101, 40, 116, 121, 112, 101, 115, 58, 32, 33, 54, 41, 10, 33, 54, 32, 61, 32, 33, 123, 125, 10, 33, 55, 32, 61, 32, 33, 68, 73, 76, 111, 99, 97, 116, 105, 111, 110, 40, 108, 105, 110, 101, 58, 32, 51, 44, 32, 99, 111, 108, 117, 109, 110, 58, 32, 56, 44, 32, 115, 99, 111, 112, 101, 58, 32, 33, 56, 41, 10, 33, 56, 32, 61, 32, 33, 68, 73, 76, 101, 120, 105, 99, 97, 108, 66, 108, 111, 99, 107, 70, 105, 108, 101, 40, 115, 99, 111, 112, 101, 58, 32, 33, 51, 44, 32, 102, 105, 108, 101, 58, 32, 33, 52, 44, 32, 100, 105, 115, 99, 114, 105, 109, 105, 110, 97, 116, 111, 114, 58, 32, 48, 41, 10, 33, 57, 32, 61, 32, 100, 105, 115, 116, 105, 110, 99, 116, 32, 33, 68, 73, 83, 117, 98, 112, 114, 111, 103, 114, 97, 109, 40, 110, 97, 109, 101, 58, 32, 34, 95, 109, 108, 105, 114, 95, 99, 105, 102, 97, 99, 101, 95, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 34, 44, 32, 108, 105, 110, 107, 97, 103, 101, 78, 97, 109, 101, 58, 32, 34, 95, 109, 108, 105, 114, 95, 99, 105, 102, 97, 99, 101, 95, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 34, 44, 32, 115, 99, 111, 112, 101, 58, 32, 110, 117, 108, 108, 44, 32, 102, 105, 108, 101, 58, 32, 33, 52, 44, 32, 108, 105, 110, 101, 58, 32, 51, 44, 32, 116, 121, 112, 101, 58, 32, 33, 53, 44, 32, 115, 99, 111, 112, 101, 76, 105, 110, 101, 58, 32, 51, 44, 32, 115, 112, 70, 108, 97, 103, 115, 58, 32, 68, 73, 83, 80, 70, 108, 97, 103, 68, 101, 102, 105, 110, 105, 116, 105, 111, 110, 32, 124, 32, 68, 73, 83, 80, 70, 108, 97, 103, 79, 112, 116, 105, 109, 105, 122, 101, 100, 44, 32, 117, 110, 105, 116, 58, 32, 33, 48, 44, 32, 114, 101, 116, 97, 105, 110, 101, 100, 78, 111, 100, 101, 115, 58, 32, 33, 54, 41, 10, 33, 49, 48, 32, 61, 32, 33, 68, 73, 76, 111, 99, 97, 116, 105, 111, 110, 40, 108, 105, 110, 101, 58, 32, 51, 44, 32, 99, 111, 108, 117, 109, 110, 58, 32, 56, 44, 32, 115, 99, 111, 112, 101, 58, 32, 33, 49, 49, 41, 10, 33, 49, 49, 32, 61, 32, 33, 68, 73, 76, 101, 120, 105, 99, 97, 108, 66, 108, 111, 99, 107, 70, 105, 108, 101, 40, 115, 99, 111, 112, 101, 58, 32, 33, 57, 44, 32, 102, 105, 108, 101, 58, 32, 33, 52, 44, 32, 100, 105, 115, 99, 114, 105, 109, 105, 110, 97, 116, 111, 114, 58, 32, 48, 41, 10, 1, 0, 0, 0, 4, 0, 0, 0, 37, 0, 0, 0, 95, 109, 108, 105, 114, 95, 99, 105, 102, 97, 99, 101, 95, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0 ]
}, {
data: [ 16, 0, 0, 0, 86, 77, 76, 65, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 84, 3, 0, 0, 32, 0, 0, 0, 66, 77, 79, 68, 24, 0, 32, 0, 28, 0, 24, 0, 20, 0, 16, 0, 12, 0, 0, 0, 0, 0, 0, 0, 8, 0, 4, 0, 24, 0, 0, 0, 140, 2, 0, 0, 92, 0, 0, 0, 80, 0, 0, 0, 68, 0, 0, 0, 40, 0, 0, 0, 20, 0, 0, 0, 4, 0, 0, 0, 6, 0, 0, 0, 109, 111, 100, 117, 108, 101, 0, 0, 3, 0, 0, 0, 80, 2, 0, 0, 44, 2, 0, 0, 12, 2, 0, 0, 5, 0, 0, 0, 176, 1, 0, 0, 108, 1, 0, 0, 32, 1, 0, 0, 216, 0, 0, 0, 156, 0, 0, 0, 1, 0, 0, 0, 88, 0, 0, 0, 1, 0, 0, 0, 20, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 159, 0, 0, 0, 5, 0, 5, 0, 120, 254, 255, 255, 8, 0, 0, 0, 20, 0, 0, 0, 6, 254, 255, 255, 4, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 24, 0, 0, 0, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 180, 254, 255, 255, 8, 0, 0, 0, 20, 0, 0, 0, 66, 254, 255, 255, 4, 0, 0, 0, 1, 0, 0, 0, 2, 0, 0, 0, 24, 0, 0, 0, 115, 105, 109, 112, 108, 101, 95, 109, 117, 108, 95, 101, 120, 95, 100, 105, 115, 112, 97, 116, 99, 104, 95, 48, 0, 0, 0, 0, 240, 254, 255, 255, 8, 0, 0, 0, 28, 0, 0, 0, 126, 254, 255, 255, 4, 0, 0, 0, 3, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 12, 0, 0, 0, 118, 109, 108, 97, 46, 109, 117, 108, 46, 102, 51, 50, 0, 0, 0, 0, 40, 255, 255, 255, 8, 0, 0, 0, 36, 0, 0, 0, 182, 254, 255, 255, 4, 0, 0, 0, 5, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 16, 0, 0, 0, 118, 109, 108, 97, 46, 98, 117, 102, 102, 101, 114, 46, 99, 111, 112, 121, 0, 0, 0, 0, 108, 255, 255, 255, 8, 0, 0, 0, 40, 0, 0, 0, 120, 255, 255, 255, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 3, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 16, 0, 0, 0, 118, 109, 108, 97, 46, 98, 117, 102, 102, 101, 114, 46, 118, 105, 101, 119, 0, 0, 0, 0, 180, 255, 255, 255, 8, 0, 0, 0, 32, 0, 0, 0, 192, 255, 255, 255, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 17, 0, 0, 0, 118, 109, 108, 97, 46, 98, 117, 102, 102, 101, 114, 46, 97, 108, 108, 111, 99, 0, 0, 0, 244, 255, 255, 255, 16, 0, 0, 0, 48, 0, 0, 0, 8, 0, 12, 0, 8, 0, 4, 0, 8, 0, 0, 0, 8, 0, 0, 0, 12, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 3, 0, 0, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 22, 0, 0, 0, 118, 109, 108, 97, 46, 105, 110, 116, 101, 114, 102, 97, 99, 101, 46, 98, 105, 110, 100, 105, 110, 103, 0, 0, 202, 255, 255, 255, 4, 0, 0, 0, 15, 0, 0, 0, 33, 118, 109, 108, 97, 46, 105, 110, 116, 101, 114, 102, 97, 99, 101, 0, 230, 255, 255, 255, 4, 0, 0, 0, 12, 0, 0, 0, 33, 118, 109, 108, 97, 46, 98, 117, 102, 102, 101, 114, 0, 0, 6, 0, 8, 0, 4, 0, 6, 0, 0, 0, 4, 0, 0, 0, 3, 0, 0, 0, 105, 51, 50, 0, 159, 0, 0, 0, 8, 0, 0, 9, 16, 0, 0, 0, 1, 0, 9, 1, 0, 0, 0, 2, 0, 9, 2, 0, 0, 0, 3, 0, 82, 0, 0, 0, 128, 3, 0, 0, 128, 0, 0, 0, 0, 1, 0, 1, 128, 82, 2, 0, 0, 128, 3, 0, 1, 192, 0, 0, 1, 0, 1, 0, 1, 128, 82, 0, 0, 0, 128, 3, 0, 0, 128, 0, 0, 2, 0, 1, 0, 2, 128, 82, 2, 0, 0, 128, 3, 0, 2, 192, 0, 0, 1, 0, 1, 0, 2, 128, 82, 1, 0, 0, 128, 1, 0, 1, 0, 1, 0, 3, 128, 82, 4, 0, 0, 128, 3, 0, 1, 192, 2, 192, 3, 128, 0, 0, 82, 0, 0, 0, 128, 3, 0, 0, 192, 0, 0, 3, 0, 1, 0, 0, 128, 82, 3, 0, 0, 128, 5, 0, 3, 192, 0, 0, 0, 192, 0, 0, 1, 0, 0, 0, 84, 0, 0, 0 ]
} ],
module_state: {
global_bytes_capacity: 2,
global_ref_count: 4
},
function_descriptors: [ {
bytecode_offset: 0,
bytecode_length: 500,
i32_register_count: 13,
ref_register_count: 9
}, {
bytecode_offset: 500,
bytecode_length: 562,
i32_register_count: 13,
ref_register_count: 9
}, {
bytecode_offset: 1062,
bytecode_length: 391,
i32_register_count: 9,
ref_register_count: 6
} ],
bytecode_data: [ 9, 50, 0, 0, 0, 0, 0, 9, 15, 0, 0, 0, 1, 0, 9, 3, 0, 0, 0, 2, 0, 9, 32, 0, 0, 3, 3, 0, 9, 2, 0, 0, 0, 4, 0, 9, 1, 0, 0, 0, 5, 0, 8, 6, 0, 9, 20, 0, 0, 0, 7, 0, 9, 5, 0, 0, 0, 8, 0, 9, 8, 0, 0, 0, 9, 0, 9, 4, 0, 0, 0, 10, 0, 82, 0, 0, 0, 128, 0, 0, 1, 0, 2, 128, 82, 15, 0, 0, 128, 1, 0, 2, 128, 1, 0, 3, 128, 83, 3, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 3, 128, 10, 0, 3, 0, 1, 0, 11, 0, 82, 4, 0, 0, 128, 4, 0, 3, 192, 0, 0, 1, 0, 11, 0, 1, 0, 3, 128, 82, 1, 0, 0, 128, 1, 0, 3, 128, 0, 0, 82, 8, 0, 0, 128, 3, 0, 2, 128, 5, 0, 2, 0, 1, 0, 4, 128, 82, 9, 0, 0, 128, 1, 0, 4, 128, 0, 0, 4, 0, 0, 0, 0, 7, 0, 0, 0, 5, 128, 4, 2, 0, 0, 0, 9, 0, 0, 0, 6, 128, 82, 5, 0, 0, 128, 1, 0, 0, 128, 1, 0, 7, 128, 83, 3, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 7, 192, 10, 0, 3, 0, 1, 0, 0, 0, 82, 5, 0, 0, 128, 1, 0, 1, 128, 1, 0, 7, 128, 83, 3, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 7, 192, 10, 0, 3, 0, 1, 0, 1, 0, 83, 12, 0, 0, 128, 7, 0, 255, 255, 255, 255, 255, 255, 3, 0, 3, 0, 3, 0, 3, 0, 15, 0, 4, 128, 6, 192, 6, 0, 6, 0, 5, 0, 4, 0, 0, 192, 1, 192, 3, 128, 6, 0, 6, 0, 6, 0, 0, 0, 1, 0, 11, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 81, 0, 0, 99, 1, 0, 0, 0, 0, 127, 1, 0, 0, 0, 0, 82, 13, 0, 0, 128, 6, 0, 4, 128, 5, 192, 6, 0, 10, 0, 5, 0, 5, 0, 0, 0, 80, 184, 1, 0, 0, 0, 0, 0, 4, 0, 0, 0, 0, 0, 81, 0, 0, 149, 1, 0, 0, 0, 0, 177, 1, 0, 0, 0, 0, 82, 13, 0, 0, 128, 6, 0, 4, 128, 5, 192, 6, 0, 10, 0, 5, 0, 5, 0, 0, 0, 80, 184, 1, 0, 0, 0, 0, 80, 184, 1, 0, 0, 0, 0, 83, 11, 0, 0, 128, 5, 0, 255, 255, 255, 255, 255, 255, 1, 0, 0, 0, 5, 0, 4, 128, 7, 0, 8, 0, 9, 0, 10, 0, 0, 0, 82, 10, 0, 0, 128, 1, 0, 4, 128, 0, 0, 82, 2, 0, 0, 128, 2, 0, 2, 192, 4, 192, 0, 0, 84, 1, 0, 3, 192, 9, 50, 0, 0, 0, 0, 0, 9, 15, 0, 0, 0, 1, 0, 9, 3, 0, 0, 0, 2, 0, 9, 2, 0, 0, 0, 3, 0, 9, 1, 0, 0, 0, 4, 0, 8, 5, 0, 9, 20, 0, 0, 0, 6, 0, 9, 5, 0, 0, 0, 7, 0, 9, 8, 0, 0, 0, 8, 0, 9, 4, 0, 0, 0, 9, 0, 9, 32, 0, 0, 3, 10, 0, 82, 7, 0, 0, 128, 1, 0, 0, 192, 1, 0, 0, 128, 82, 7, 0, 0, 128, 1, 0, 1, 192, 1, 0, 1, 128, 82, 0, 0, 0, 128, 0, 0, 1, 0, 2, 128, 82, 15, 0, 0, 128, 1, 0, 2, 128, 1, 0, 3, 128, 83, 3, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 3, 128, 9, 0, 10, 0, 1, 0, 11, 0, 82, 4, 0, 0, 128, 4, 0, 3, 192, 0, 0, 1, 0, 11, 0, 1, 0, 3, 128, 82, 1, 0, 0, 128, 1, 0, 3, 128, 0, 0, 82, 8, 0, 0, 128, 3, 0, 2, 128, 4, 0, 2, 0, 1, 0, 4, 128, 82, 9, 0, 0, 128, 1, 0, 4, 128, 0, 0, 4, 0, 0, 0, 0, 7, 0, 0, 0, 5, 128, 4, 2, 0, 0, 0, 9, 0, 0, 0, 6, 128, 82, 5, 0, 0, 128, 1, 0, 0, 128, 1, 0, 7, 128, 83, 3, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 7, 192, 9, 0, 10, 0, 1, 0, 0, 0, 82, 5, 0, 0, 128, 1, 0, 1, 128, 1, 0, 7, 128, 83, 3, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 7, 192, 9, 0, 10, 0, 1, 0, 1, 0, 83, 12, 0, 0, 128, 7, 0, 255, 255, 255, 255, 255, 255, 3, 0, 3, 0, 3, 0, 3, 0, 15, 0, 4, 128, 6, 192, 5, 0, 5, 0, 4, 0, 3, 0, 0, 192, 1, 192, 3, 128, 5, 0, 5, 0, 5, 0, 0, 0, 1, 0, 11, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 81, 0, 0, 125, 1, 0, 0, 0, 0, 153, 1, 0, 0, 0, 0, 82, 13, 0, 0, 128, 6, 0, 4, 128, 5, 192, 5, 0, 9, 0, 4, 0, 4, 0, 0, 0, 80, 210, 1, 0, 0, 0, 0, 0, 4, 0, 0, 0, 0, 0, 81, 0, 0, 175, 1, 0, 0, 0, 0, 203, 1, 0, 0, 0, 0, 82, 13, 0, 0, 128, 6, 0, 4, 128, 5, 192, 5, 0, 9, 0, 4, 0, 4, 0, 0, 0, 80, 210, 1, 0, 0, 0, 0, 80, 210, 1, 0, 0, 0, 0, 83, 11, 0, 0, 128, 5, 0, 255, 255, 255, 255, 255, 255, 1, 0, 0, 0, 5, 0, 4, 128, 6, 0, 7, 0, 8, 0, 9, 0, 0, 0, 82, 10, 0, 0, 128, 1, 0, 4, 128, 0, 0, 82, 2, 0, 0, 128, 2, 0, 2, 192, 4, 192, 0, 0, 80, 20, 2, 0, 0, 1, 0, 3, 128, 0, 128, 83, 6, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 0, 192, 9, 0, 10, 0, 1, 0, 0, 128, 84, 1, 0, 0, 192, 9, 1, 0, 0, 0, 0, 0, 9, 2, 0, 0, 0, 1, 0, 9, 6, 0, 0, 0, 2, 0, 8, 3, 0, 9, 77, 86, 76, 76, 4, 0, 9, 65, 76, 77, 86, 5, 0, 10, 0, 128, 9, 7, 0, 0, 0, 6, 0, 82, 0, 0, 0, 128, 0, 0, 1, 0, 1, 128, 11, 0, 0, 0, 0, 2, 128, 82, 16, 0, 0, 128, 2, 0, 1, 192, 2, 192, 1, 0, 7, 0, 1, 0, 0, 0, 0, 7, 0, 82, 0, 0, 0, 128, 0, 0, 1, 0, 1, 128, 11, 1, 0, 0, 0, 2, 128, 82, 16, 0, 0, 128, 2, 0, 1, 192, 2, 192, 1, 0, 7, 0, 1, 4, 0, 0, 0, 7, 0, 82, 0, 0, 0, 128, 0, 0, 1, 0, 1, 128, 83, 14, 0, 0, 128, 3, 0, 255, 255, 255, 255, 9, 0, 11, 0, 1, 192, 0, 0, 3, 0, 6, 0, 0, 0, 0, 0, 6, 0, 0, 0, 1, 0, 6, 0, 2, 0, 1, 0, 1, 128, 5, 1, 0, 0, 0, 5, 0, 0, 0, 1, 192, 4, 1, 0, 0, 0, 5, 0, 0, 0, 1, 128, 82, 0, 0, 0, 128, 0, 0, 1, 0, 2, 128, 83, 20, 0, 0, 128, 3, 0, 255, 255, 1, 0, 255, 255, 3, 0, 2, 192, 1, 192, 3, 0, 1, 0, 1, 128, 5, 2, 0, 0, 0, 9, 0, 0, 0, 1, 192, 82, 0, 0, 0, 128, 0, 0, 1, 0, 1, 128, 11, 2, 0, 0, 0, 2, 128, 82, 17, 0, 0, 128, 2, 0, 1, 192, 2, 192, 1, 0, 1, 128, 4, 2, 0, 0, 0, 9, 0, 0, 0, 2, 128, 83, 18, 0, 0, 128, 2, 0, 255, 255, 2, 0, 3, 0, 1, 128, 4, 0, 5, 0, 1, 0, 0, 0, 11, 3, 0, 0, 0, 3, 128, 11, 4, 0, 0, 0, 4, 128, 33, 0, 0, 10, 0, 0, 0, 0, 192, 2, 0, 3, 128, 4, 128, 0, 128, 82, 19, 0, 0, 128, 4, 0, 1, 128, 2, 192, 6, 0, 0, 192, 1, 0, 0, 128, 5, 0, 0, 0, 0, 7, 0, 0, 0, 0, 192, 5, 3, 0, 0, 0, 8, 0, 0, 0, 1, 192, 84, 0, 0 ]
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment