Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Greater-than-linear slowdown with linear increase in shader size #1598

Closed
mkeeter opened this issue Jan 19, 2021 · 2 comments
Closed

Greater-than-linear slowdown with linear increase in shader size #1598

mkeeter opened this issue Jan 19, 2021 · 2 comments
Labels
needs triage Needs to be reproduced before it can become a different issue type.

Comments

@mkeeter
Copy link

mkeeter commented Jan 19, 2021 •
edited
Loading

Good morning!

I'm attempting to use SPIRV-Cross (in the wgpu-native stack) to compile shaders for a raytracing kernel.

Doing a simple test with spheres scattered around the scene, I notice an exponential-ish slowdown with number of spheres:

Number of spheres spirv-cross time (ms)
10 13
15 30
16 42
17 72
18 122
19 225
20 443
21 907
22 1890

Here are example shaders with 10 spheres and 22 spheres (linked to the function which changes; everything above is boilerplate). You can reproduce timing with glslc shader10.frag && time spirv-cross --msl a.spv.

A stack trace shows an alarming deep set of calls into exists_unaccessed_path_to_return.

...
    frame #918: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #919: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #920: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #921: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #922: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #923: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #924: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #925: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #926: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #927: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #928: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #929: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #930: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #931: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #932: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #933: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #934: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #935: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #936: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #937: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #938: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #939: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #940: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #941: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #942: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #943: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #944: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #945: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #946: 0x0000000100efe240 libwgpu_native.dylib`exists_unaccessed_path_to_return(spirv_cross::CFG const&, unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > const&) + 640
    frame #947: 0x0000000100efdf95 libwgpu_native.dylib`spirv_cross::Compiler::analyze_parameter_preservation(spirv_cross::SPIRFunction&, spirv_cross::CFG const&, std::__1::unordered_map<unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> >, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<std::__1::pair<unsigned int const, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > > > > const&, std::__1::unordered_map<unsigned int, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> >, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<std::__1::pair<unsigned int const, std::__1::unordered_set<unsigned int, std::__1::hash<unsigned int>, std::__1::equal_to<unsigned int>, std::__1::allocator<unsigned int> > > > > const&) + 533
    frame #948: 0x0000000100f005f7 libwgpu_native.dylib`spirv_cross::Compiler::analyze_variable_scope(spirv_cross::SPIRFunction&, spirv_cross::Compiler::AnalyzeVariableScopeAccessHandler&) + 263
    frame #949: 0x0000000100f0425f libwgpu_native.dylib`spirv_cross::Compiler::build_function_control_flow_graphs_and_analyze() + 847
    frame #950: 0x0000000100fadb21 libwgpu_native.dylib`spirv_cross::CompilerMSL::compile() + 481
    frame #951: 0x0000000100ef1b07 libwgpu_native.dylib`sc_internal_compiler_msl_compile + 167
    frame #952: 0x0000000100eeddfa libwgpu_native.dylib`spirv_cross::msl::_$LT$impl$u20$spirv_cross..spirv..Compile$LT$spirv_cross..msl..Target$GT$$u20$for$u20$spirv_cross..spirv..Ast$LT$spirv_cross..msl..Target$GT$$GT$::compile::h560210ea9c97c90b + 74
    frame #953: 0x0000000100ec5abc libwgpu_native.dylib`gfx_backend_metal::device::Device::compile_shader_library::h93eb59e12fa03a92 + 1068
    frame #954: 0x0000000100ec6f37 libwgpu_native.dylib`gfx_backend_metal::device::Device::load_shader::h7551513e699ed3c7 + 2103
    frame #955: 0x0000000100ec7f50 libwgpu_native.dylib`_$LT$gfx_backend_metal..device..Device$u20$as$u20$gfx_hal..device..Device$LT$gfx_backend_metal..Backend$GT$$GT$::create_graphics_pipeline::h0506d789ace60e17 + 1264
    frame #956: 0x0000000100db2b1e libwgpu_native.dylib`wgpu_core::device::_$LT$impl$u20$wgpu_core..hub..Global$LT$G$GT$$GT$::device_create_render_pipeline::h6b9afe1e50b2e585 + 6974
    frame #957: 0x0000000100dd7140 libwgpu_native.dylib`wgpu_device_create_render_pipeline + 80
    frame #958: 0x00000001000ca57f rayray`preview.Preview.init(alloc=0x00007ffeefbff490, scene=<unavailable>, device=4611686022722355200, uniform_buf=4611686022722355200) at preview.zig:117
    frame #959: 0x00000001000c7c41 rayray`renderer.Renderer.init(alloc=0x00007ffeefbff490, scene=<unavailable>, options=<unavailable>, device=4611686022722355200) at renderer.zig:57
    frame #960: 0x00000001000be687 rayray`window.Window.init(alloc=0x00007ffeefbff490, options_=<unavailable>, name="rayray") at window.zig:108
    frame #961: 0x00000001000bc70b rayray`main at main.zig:18
    frame #962: 0x00000001000becb3 rayray`std.start.main at start.zig:345
    frame #963: 0x00000001000beca7 rayray`std.start.main [inlined] std.start.initEventLoopAndCallMain at start.zig:287
    frame #964: 0x00000001000beca7 rayray`std.start.main [inlined] std.start.callMainWithArgs(argc=1, argv=0x00007ffeefbff768, envp=[][*:0]u8 @ 0x00007ffeefbff600) at start.zig:250
    frame #965: 0x00000001000bec46 rayray`std.start.main(c_argc=1, c_argv=0x00007ffeefbff768, c_envp=0x00007ffeefbff778) at start.zig:257
    frame #966: 0x00007fff6cbfe015 libdyld.dylib`start + 1
    frame #967: 0x00007fff6cbfe015 libdyld.dylib`start + 1

Any ideas?
@mkeeter mkeeter changed the title Extreme slowdown on large scenes Exponential slowdown with linear increase in shader size Jan 19, 2021
@mkeeter mkeeter changed the title Exponential slowdown with linear increase in shader size Greater-than-linear slowdown with linear increase in shader size Jan 19, 2021
@mkeeter
Copy link
Author

mkeeter commented Jan 19, 2021

Oh, I just noticed #1594, which seems relevant 😆 – feel free to close if that's the fix!

@HansKristian-Work HansKristian-Work added the needs triage Needs to be reproduced before it can become a different issue type. label Jan 19, 2021
@HansKristian-Work
Copy link
Contributor

The 22 sphere one takes ~3ms per iteration for MSL now, so calling it fixed.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
needs triage Needs to be reproduced before it can become a different issue type.
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
2 participants
@mkeeter @HansKristian-Work