Skip to content

Instantly share code, notes, and snippets.

@benvanik

benvanik/spirv_executable.mlir

Created February 26, 2020 04:09
Show Gist options
  • Save benvanik/26790195c7142e43594cd3754af1ab11 to your computer and use it in GitHub Desktop.
Save benvanik/26790195c7142e43594cd3754af1ab11 to your computer and use it in GitHub Desktop.
Download ZIP
Raw
spirv_executable.mlir
hal.executable @simpleMath_ex_dispatch_0 {
hal.interface @legacy_io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
}
hal.executable.entry_point @simpleMath_rgn_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<4xf32>) -> tensor<4xf32>, workgroup_size = dense<[32, 1, 1]> : vector<3xi32>}
hal.executable.binary attributes {data = dense<"0x0800000053505645A4FFFFFFA00600007000000004000000EEFFFFFF04000000010000000C000000000006000800040006000000040000000200000034000000100000000C001400040008000C0010000C000000010000000700000001000000200000000C0010000000040008000C000C0000000700000001000000200000008A010000030223070000010016000000390000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000C00050000000F00000073696D706C654D6174685F72676E5F64697370617463685F300000000500000004000000100006000F0000001100000020000000010000000100000005000A00040000005F5F6275696C74696E5F7661725F4E756D576F726B67726F7570735F5F00000005000B00050000005F5F6275696C74696E5F7661725F476C6F62616C496E766F636174696F6E49645F5F000005000A000B00000073696D706C654D6174685F72676E5F64697370617463685F305F6172675F300005000A000C00000073696D706C654D6174685F72676E5F64697370617463685F305F6172675F3100050009000F00000073696D706C654D6174685F72676E5F64697370617463685F3000000047000400040000000B0000001800000047000400050000000B0000001C000000470004000800000006000000040000004800050007000000000000002300000000000000470003000700000002000000470004000B0000002100000000000000470004000B0000002200000000000000470003000700000002000000470004000C0000002100000001000000470004000C00000022000000000000001500040003000000200000000100000017000400020000000300000003000000200004000100000001000000020000003B0004000100000004000000010000003B0004000100000005000000010000001600030009000000200000002B000400030000000A000000040000001C00040008000000090000000A0000001E000300070000000800000020000400060000000C000000070000003B000400060000000B0000000C0000003B000400060000000C0000000C000000130002000E000000210003000D0000000E0000002B0004000300000011000000010000002B0004000300000012000000200000002B000400030000001300000000000000140002001D00000020000400310000000C00000009000000360005000E0000000F000000000000000D000000F8000200100000003D00040002000000140000000500000051000500030000001500000014000000020000003D0004000200000016000000040000005100050003000000170000001600000002000000F900020018000000F800020018000000F5000700030000001B0000001500000010000000380000001C000000B10005001D0000001E0000001B00000011000000F60004001A0000001900000000000000FA0004001E000000190000001A000000F8000200190000003D000400020000001F000000050000005100050003000000200000001F000000010000003D0004000200000021000000040000005100050003000000220000002100000001000000F900020023000000F800020023000000F5000700030000002500000020000000190000003700000026000000B10005001D000000270000002500000011000000F60004001C0000002400000000000000FA00040027000000240000001C000000F8000200240000003D00040002000000280000000500000051000500030000002900000028000000000000003D000400020000002A0000000400000051000500030000002B0000002A0000000000000084000500030000002C0000002B00000012000000F90002002D000000F80002002D000000F5000700030000002F0000002900000024000000360000002E000000B10005001D000000300000002F0000000A000000F6000400260000002E00000000000000FA000400300000002E00000026000000F80002002E0000004100060031000000320000000B000000130000002F0000003D00040009000000330000003200000081000500090000003400000033000000330000004100060031000000350000000C000000130000002F0000003E00030035000000340000008000050003000000360000002F0000002C000000F90002002D000000F8000200260000008000050003000000370000002500000022000000F900020023000000F80002001C0000008000050003000000380000001B00000017000000F900020018000000F80002001A000000FD0001003800010001000000040000001900000073696D706C654D6174685F72676E5F64697370617463685F30000000"> : vector<1748xi8>, format = 1397773893 : i32} {
module {
spv.module "Logical" "GLSL450" {
spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
spv.globalVariable @__builtin_var_GlobalInvocationId__ built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
spv.globalVariable @simpleMath_rgn_dispatch_0_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
spv.globalVariable @simpleMath_rgn_dispatch_0_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
spv.func @simpleMath_rgn_dispatch_0() "None" {
%0 = spv.constant 4 : i32
%1 = spv.constant 1 : i32
%2 = spv.constant 32 : i32
%3 = spv.constant 0 : i32
%4 = spv._address_of @simpleMath_rgn_dispatch_0_arg_1 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
%5 = spv._address_of @simpleMath_rgn_dispatch_0_arg_0 : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
%6 = spv._address_of @__builtin_var_GlobalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
%7 = spv.Load "Input" %6 : vector<3xi32>
%8 = spv.CompositeExtract %7[2 : i32] : vector<3xi32>
%9 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
%10 = spv.Load "Input" %9 : vector<3xi32>
%11 = spv.CompositeExtract %10[2 : i32] : vector<3xi32>
spv.loop {
spv.Branch ^bb1(%8 : i32)
^bb1(%12: i32): // 2 preds: ^bb0, ^bb2
%13 = spv.SLessThan %12, %1 : i32
spv.BranchConditional %13, ^bb2, ^bb3
^bb2: // pred: ^bb1
%14 = spv._address_of @__builtin_var_GlobalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
%15 = spv.Load "Input" %14 : vector<3xi32>
%16 = spv.CompositeExtract %15[1 : i32] : vector<3xi32>
%17 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
%18 = spv.Load "Input" %17 : vector<3xi32>
%19 = spv.CompositeExtract %18[1 : i32] : vector<3xi32>
spv.loop {
spv.Branch ^bb1(%16 : i32)
^bb1(%21: i32): // 2 preds: ^bb0, ^bb2
%22 = spv.SLessThan %21, %1 : i32
spv.BranchConditional %22, ^bb2, ^bb3
^bb2: // pred: ^bb1
%23 = spv._address_of @__builtin_var_GlobalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
%24 = spv.Load "Input" %23 : vector<3xi32>
%25 = spv.CompositeExtract %24[0 : i32] : vector<3xi32>
%26 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
%27 = spv.Load "Input" %26 : vector<3xi32>
%28 = spv.CompositeExtract %27[0 : i32] : vector<3xi32>
%29 = spv.IMul %28, %2 : i32
spv.loop {
spv.Branch ^bb1(%25 : i32)
^bb1(%31: i32): // 2 preds: ^bb0, ^bb2
%32 = spv.SLessThan %31, %0 : i32
spv.BranchConditional %32, ^bb2, ^bb3
^bb2: // pred: ^bb1
%33 = spv.AccessChain %5[%3, %31] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
%34 = spv.Load "StorageBuffer" %33 : f32
%35 = spv.FAdd %34, %34 : f32
%36 = spv.AccessChain %4[%3, %31] : !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
spv.Store "StorageBuffer" %36, %35 : f32
%37 = spv.IAdd %31, %29 : i32
spv.Branch ^bb1(%37 : i32)
^bb3: // pred: ^bb1
spv._merge
}
%30 = spv.IAdd %21, %19 : i32
spv.Branch ^bb1(%30 : i32)
^bb3: // pred: ^bb1
spv._merge
}
%20 = spv.IAdd %12, %11 : i32
spv.Branch ^bb1(%20 : i32)
^bb3: // pred: ^bb1
spv._merge
}
spv.Return
}
spv.EntryPoint "GLCompute" @simpleMath_rgn_dispatch_0, @__builtin_var_GlobalInvocationId__, @__builtin_var_NumWorkgroups__
spv.ExecutionMode @simpleMath_rgn_dispatch_0 "LocalSize", 32, 1, 1
} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
}
}
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment