Wumpf
commented
4 years ago Known issues
- Can't do atomic ops on images:
- spirv-cross makes a separate binding that's a buffer pointing to the same storage, which confuses the wgpu
- could use a buffer instead of an image for the one usecase I have right now. (this could also become come a rabbit hole though and perf implications is unknown)
- https://github.com/Wumpf/blub/blob/79beda0/shader/simulation/transfer_build_linkedlist.comp#L26
- see https://github.com/gfx-rs/wgpu/issues/618 / https://github.com/gpuweb/gpuweb/issues/728
- spirv-cross makes a separate binding that's a buffer pointing to the same storage, which confuses the wgpu
- has somehow trouble with some load (??) in the remove_divergence.comp shader:
Andreas R, [23.06.20 00:26] [2020-06-22T22:25:31Z ERROR gfx_backend_metal::device] Error compiling the shader CompilationFailed("Compilation failed: \n\nprogram_source:63:38: error: no matching member function for call to \'read\'\n uint neighborType = MarkerVolume.read(uint3(coord)).x;\n ~~~~~~~~~~~~~^~~~\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.11/include/metal/metal_texture:2811:24: note: candidate function not viable: \'this\' argument has type \'volatile texture3d<uint>\' (aka \'volatile texture3d<unsigned int>\'), but method is not marked volatile\n METAL_FUNC vec<T, 4> read(uint3 coord, uint lod = 0) const thread\n ^\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.11/include/metal/metal_texture:2828:24: note: candidate function not viable: \'this\' argument (\'volatile texture3d<uint>\' (aka \'volatile texture3d<unsigned int>\')) is in address space 0, but parameter must be in address space device\n METAL_FUNC vec<T, 4> read(uint3 coord, uint lod = 0) const device\n ^\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.11/include/metal/metal_texture:2846:24: note: candidate function not viable: \'this\' argument (\'volatile texture3d<uint>\' (aka \'volatile texture3d<unsigned int>\')) is in address space 0, but parameter must be in address space constant\n METAL_FUNC vec<T, 4> read(uint3 coord, uint lod = 0) const constant\n ^\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.11/include/metal/metal_texture:2803:24: note: candidate function not viable: \'this\' argument has type \'volatile texture3d<uint>\' (aka \'volatile texture3d<unsigned int>\'), but method is not marked volatile\n METAL_FUNC vec<T, 4> read(ushort3 coord, ushort lod = 0) const thread\n ^\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.11/include/metal/metal_texture:2820:24: note: candidate function not viable: \'this\' argument (\'volatile texture3d<uint>\' (aka \'volatile texture3d<unsigned int>\')) is in address space 0, but parameter must be in address space device\n METAL_FUNC vec<T, 4> read(ushort3 coord, ushort lod = 0) const device\n ^\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.11/include/metal/metal_texture:2838:24: note: candidate function not viable: \'this\' argument (\'volatile texture3d<uint>\' (aka \'volatile texture3d<unsigned int>\')) is in address space 0, but parameter must be in address space constant\n METAL_FUNC vec<T, 4> read(ushort3 coord, ushort lod = 0) const constant\n ^\nprogram_source:86:12: warning: unused variable \'velocityGridPosition\'\n float3 velocityGridPosition = float3(gl_GlobalInvocationID);\n ^\n") thread 'main' panicked at 'called Result::unwrap() on an Err value: Other', /Users/andreasreich/.cargo/git/checkouts/wgpu-53e70f8674b08dd4/a02a566/wgpu-core/src/device/mod.rs:2412:17 note: run with RUST_BACKTRACE=1 environment variable to display a backtrace
pannous
Looks like there's some things in the shaders that have trouble being converted to MSL!