wasm2spirv - Compile your WebAssembly programs into SPIR-V shaders
Warning
wasm2spirvis still in early development, and not production ready.
This repository contains the code for both, the CLI and library for wasm2spirv. wasm2spirv allows you to compile any WebAssembly program into a SPIR-V shader
Features
- Compiles your WebAssembly programs into SPIR-V
- Can transpile into other various shading languages
- Supports validation and optimization of the resulting SPIR-V
- Can be compiled to WebAssembly itself
- You won't be able to use
spirv-tools,spirv_crossortree-sitterin WebAssembly - CLI will have to be compiled to WASI
- You won't be able to use
Caveats
- Still in early development
- Unexpected bugs and crashes are to be expected
- Still working through the WebAssembly MVP
- WebAssembly programs with memory allocations will not work
- You can customize whether the
memory.growinstruction errors the compilation (hard errors) or always returns -1 (soft errors)
- You can customize whether the
- You'll have to manually provide quite some extra information
- This is because SPIR-V has a lot of constructs compared to the simplicity of WebAssembly.
- wasm2spirv can do some inferrence based on the WebAssembly program itself, but it's usually better to specify most the information on the configuration.
- The plan for the future is to be able to store the config information inside the WebAssembly program itself.
Compilation Targets
| Target | Windows | Linux | macOS | WebAssembly |
|---|---|---|---|---|
| SPIR-V | ✅ |
|
|
|
| GLSL |
|
☑️ (spvc-glsl/naga-glsl) | ☑️ (spvc-glsl/naga-glsl) | ☑️ (naga-glsl) |
| HLSL |
|
|
|
|
| Metal (MSL) |
|
|
|
☑️ (naga-msl) |
| WGSL |
|
|
|
☑️ (naga-wgsl) |
| DXIL |
|
|
❌ |
|
| OpenCL C |
|
|
|
|
| Cuda |
|
|
❌ |
|
| Validation |
|
|
|
|
-
✅ Supported - ☑️ Supported, but requires cargo feature(s)
-
❌ Unsupported
Note
The CLI programs built by the releases use the Khronos compilers/validators whenever possible, faling back to naga compilers/validators if the Khronos are not available or are not supported on that platform.
Examples
You can find a few examples on the "examples" directory, with their Zig file, translated WebAssembly Text, and compilation configuration file.
Saxpy example
Zig program
export fn main(n: usize, alpha: f32, x: [*]const f32, y: [*]f32) void {
var i = gl_GlobalInvocationID(0);
const size = gl_NumWorkGroups(0);
while (i < n) {
y[i] += alpha * x[i];
i += size;
}
}
extern "spir_global" fn gl_GlobalInvocationID(u32) usize;
extern "spir_global" fn gl_NumWorkGroups(u32) usize;
WebAssembly text
(module
(type (;0;) (func (param i32) (result i32)))
(type (;1;) (func (param i32 f32 i32 i32)))
(import "spir_global" "gl_GlobalInvocationID" (func (;0;) (type 0)))
(import "spir_global" "gl_NumWorkGroups" (func (;1;) (type 0)))
(func (;2;) (type 1) (param i32 f32 i32 i32)
(local i32 i32 i32 i32 i32)
i32.const 0
call 0
local.tee 4
i32.const 2
i32.shl
local.set 5
i32.const 0
call 1
local.tee 6
i32.const 2
i32.shl
local.set 7
block ;; label = @1
loop ;; label = @2
local.get 4
local.get 0
i32.ge_u
br_if 1 (;@1;)
local.get 3
local.get 5
i32.add
local.tee 8
local.get 8
f32.load
local.get 2
local.get 5
i32.add
f32.load
local.get 1
f32.mul
f32.add
f32.store
local.get 5
local.get 7
i32.add
local.set 5
local.get 4
local.get 6
i32.add
local.set 4
br 0 (;@2;)
end
end)
(memory (;0;) 16)
(global (;0;) (mut i32) (i32.const 1048576))
(export "memory" (memory 0))
(export "main" (func 2)))
Configuration file (in JSON)
{
"platform": {
"vulkan": "1.1"
},
"version": "1.3",
"addressing_model": "logical",
"memory_model": "GLSL450",
"capabilities": { "dynamic": ["VariablePointers"] },
"extensions": ["VH_KHR_variable_pointers"],
"functions": {
"2": {
"execution_model": "GLCompute",
"execution_modes": [{
"local_size": [1, 1, 1]
}],
"params": {
"0": {
"type": "i32",
"kind": {
"descriptor_set": {
"storage_class": "StorageBuffer",
"set": 0,
"binding": 0
}
}
},
"1": {
"type": "f32",
"kind": {
"descriptor_set": {
"storage_class": "StorageBuffer",
"set": 0,
"binding": 1
}
}
},
"2": {
"type": {
"size": "fat",
"storage_class": "StorageBuffer",
"pointee": "f32"
},
"kind": {
"descriptor_set": {
"storage_class": "StorageBuffer",
"set": 0,
"binding": 2
}
}
},
"3": {
"type": {
"size": "fat",
"storage_class": "StorageBuffer",
"pointee": "f32"
},
"kind": {
"descriptor_set": {
"storage_class": "StorageBuffer",
"set": 0,
"binding": 3
}
}
}
}
}
}
}
SPIR-V result
; SPIR-V
; Version: 1.3
; Generator: rspirv
; Bound: 73
OpCapability VariablePointers
OpCapability Shader
OpExtension "VH_KHR_variable_pointers"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %3 "main" %6 %7
OpExecutionMode %3 LocalSize 1 1 1
OpDecorate %6 BuiltIn GlobalInvocationId
OpDecorate %7 BuiltIn NumWorkgroups
OpMemberDecorate %10 0 Offset 0
OpDecorate %10 Block
OpDecorate %12 DescriptorSet 0
OpDecorate %12 Binding 0
OpMemberDecorate %14 0 Offset 0
OpDecorate %14 Block
OpDecorate %16 DescriptorSet 0
OpDecorate %16 Binding 1
OpDecorate %17 ArrayStride 4
OpMemberDecorate %18 0 Offset 0
OpDecorate %18 Block
OpDecorate %20 DescriptorSet 0
OpDecorate %20 Binding 2
OpDecorate %21 DescriptorSet 0
OpDecorate %21 Binding 3
%1 = OpTypeInt 32 0
%2 = OpConstant %1 1048576
%4 = OpTypeVector %1 3
%5 = OpTypePointer Input %4
%6 = OpVariable %5 Input
%7 = OpVariable %5 Input
%8 = OpTypeVoid
%9 = OpTypeFunction %8
%10 = OpTypeStruct %1
%11 = OpTypePointer StorageBuffer %10
%12 = OpVariable %11 StorageBuffer
%13 = OpTypeFloat 32
%14 = OpTypeStruct %13
%15 = OpTypePointer StorageBuffer %14
%16 = OpVariable %15 StorageBuffer
%17 = OpTypeRuntimeArray %13
%18 = OpTypeStruct %17
%19 = OpTypePointer StorageBuffer %18
%20 = OpVariable %19 StorageBuffer
%21 = OpVariable %19 StorageBuffer
%23 = OpTypePointer Function %1
%28 = OpConstant %1 2
%39 = OpTypeBool
%41 = OpConstant %1 0
%42 = OpTypePointer StorageBuffer %1
%46 = OpTypePointer Function %19
%50 = OpTypePointer StorageBuffer %13
%51 = OpConstant %1 4
%3 = OpFunction %8 None %9
%22 = OpLabel
%48 = OpVariable %23 Function %41
%47 = OpVariable %46 Function
%33 = OpVariable %23 Function
%30 = OpVariable %23 Function
%27 = OpVariable %23 Function
%24 = OpVariable %23 Function
%25 = OpLoad %4 %6
%26 = OpCompositeExtract %1 %25 0
OpStore %24 %26
%29 = OpShiftLeftLogical %1 %26 %28
OpStore %27 %29
%31 = OpLoad %4 %7
%32 = OpCompositeExtract %1 %31 0
OpStore %30 %32
%34 = OpShiftLeftLogical %1 %32 %28
OpStore %33 %34
OpBranch %35
%35 = OpLabel
OpBranch %36
%36 = OpLabel
%40 = OpLoad %1 %24
%43 = OpAccessChain %42 %12 %41
%44 = OpLoad %1 %43
%45 = OpUGreaterThanEqual %39 %40 %44
OpLoopMerge %37 %38 None
OpBranchConditional %45 %37 %38
%38 = OpLabel
OpStore %47 %21
%49 = OpLoad %1 %27
OpStore %48 %49
%52 = OpUDiv %1 %49 %51
%53 = OpAccessChain %50 %21 %41 %52
%54 = OpLoad %19 %47
%55 = OpLoad %1 %48
%56 = OpUDiv %1 %55 %51
%57 = OpAccessChain %50 %54 %41 %56
%58 = OpLoad %13 %57 Aligned 4
%59 = OpLoad %1 %27
%60 = OpUDiv %1 %59 %51
%61 = OpAccessChain %50 %20 %41 %60
%62 = OpLoad %13 %61 Aligned 4
%63 = OpAccessChain %50 %16 %41
%64 = OpLoad %13 %63
%65 = OpFMul %13 %62 %64
%66 = OpFAdd %13 %58 %65
OpStore %53 %66 Aligned 4
%67 = OpLoad %1 %27
%68 = OpLoad %1 %33
%69 = OpIAdd %1 %67 %68
OpStore %27 %69
%70 = OpLoad %1 %24
%71 = OpLoad %1 %30
%72 = OpIAdd %1 %70 %71
OpStore %24 %72
OpBranch %36
%37 = OpLabel
OpReturn
OpFunctionEnd
Metal translation
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct _10
{
uint _m0;
};
struct _14
{
float _m0;
};
struct _18
{
float _m0[1];
};
kernel void main0(device _10& _12 [[buffer(0)]], device _14& _16 [[buffer(1)]], device _18& _20 [[buffer(2)]], device _18& _21 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]])
{
uint _48 = 0u;
uint _29 = gl_GlobalInvocationID.x << 2u;
uint _34 = gl_NumWorkGroups.x << 2u;
device _18* _47;
for (uint _24 = gl_GlobalInvocationID.x, _27 = _29, _30 = gl_NumWorkGroups.x, _33 = _34; !(_24 >= _12._m0); )
{
_47 = &_21;
_48 = _27;
_21._m0[_27 / 4u] = _47->_m0[_48 / 4u] + (_20._m0[_27 / 4u] * _16._m0);
_27 += _33;
_24 += _30;
continue;
}
}
Installation
To add wasm2spirv as a library for your Rust project, run this command on you'r project's root directory.
cargo add wasm2spirv
To install the latest version of the wasm2spirv CLI, run this command.
cargo install wasm2spirv
Cargo features
spirv-toolsenables optimization and validation.spirv_crossenables cross-compilation to GLSL, HLSL and MSL.tree-sitterenables syntax highlighting on the CLI.nagaenables cross-compilation for GLSL, HLSL, MSL and WGSL.
Related projects
- SPIRV-LLVM is an official Khronos tool to compile LLVM IR into SPIR-V.
- Wasmer is a WebAssembly runtime that runs WebAssembly programs on the host machine.
- Bytecoder can translate JVM code into JavaScript, WebAssembly and OpenCL.
- Naga is a translator from, and to, various shading languages and IRs.
551 Dec 18, 2022
5.5k Jan 9, 2023
71 Dec 16, 2022
354 Jan 1, 2023
4 May 10, 2022
24 Feb 1, 2023
11 Aug 28, 2023
2.3k Dec 31, 2022
343 Dec 22, 2022
15 Oct 11, 2022
39 Dec 21, 2022
20 Dec 30, 2022
145 Apr 10, 2023
17 Oct 14, 2023
18 Nov 27, 2023
53 Oct 11, 2022
37 Aug 1, 2022
2 Dec 17, 2022