wgebra 0.2.0

Composable WGSL shaders for linear algebra.
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110

use crate::WgQuat;
use nalgebra::{Matrix4x3, Vector4};
use wgcore::Shader;
#[cfg(test)]
use {
    naga_oil::compose::NagaModuleDescriptor,
    wgpu::{ComputePipeline, Device},
};

#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
#[repr(C)]
/// A 3D SVD as represented on the gpu, with padding (every fourth rows
/// can be ignored).
// TODO: switch to encase?
pub struct GpuSvd3 {
    /// First orthogonal matrix of the SVD.
    u: Matrix4x3<f32>,
    /// Singular values.
    s: Vector4<f32>,
    /// Second orthogonal matrix of the SVD.
    vt: Matrix4x3<f32>,
}

#[derive(Shader)]
#[shader(derive(WgQuat), src = "svd3.wgsl")]
/// Shader for computing the Singular Value Decomposition of 3x3 matrices.
pub struct WgSvd3;

impl WgSvd3 {
    #[cfg(test)]
    pub fn tests(device: &Device) -> ComputePipeline {
        let test_kernel = r#"
@group(0) @binding(0)
var<storage, read_write> in: array<mat3x3<f32>>;
@group(0) @binding(1)
var<storage, read_write> out: array<Svd>;

@compute @workgroup_size(1, 1, 1)
fn test(@builtin(global_invocation_id) invocation_id: vec3<u32>) {
    let i = invocation_id.x;
    out[i] = svd(in[i]);
}
        "#;

        let src = format!("{}\n{}", Self::src(), test_kernel);
        let module = WgQuat::composer()
            .make_naga_module(NagaModuleDescriptor {
                source: &src,
                file_path: Self::FILE_PATH,
                ..Default::default()
            })
            .unwrap();
        wgcore::utils::load_module(device, "test", module)
    }
}

#[cfg(test)]
mod test {
    use super::GpuSvd3;
    use approx::assert_relative_eq;
    use nalgebra::{DVector, Matrix3, Matrix4x3};
    use wgcore::gpu::GpuInstance;
    use wgcore::kernel::{KernelInvocationBuilder, KernelInvocationQueue};
    use wgcore::tensor::GpuVector;
    use wgpu::BufferUsages;

    #[futures_test::test]
    #[serial_test::serial]
    async fn gpu_svd3() {
        let gpu = GpuInstance::new().await.unwrap();
        let svd = super::WgSvd3::tests(gpu.device());
        let mut queue = KernelInvocationQueue::new(gpu.device());
        let mut encoder = gpu.device().create_command_encoder(&Default::default());

        const LEN: usize = 345;
        let matrices: DVector<Matrix4x3<f32>> = DVector::new_random(LEN);
        let inputs = GpuVector::init(gpu.device(), &matrices, BufferUsages::STORAGE);
        let result: GpuVector<GpuSvd3> = GpuVector::uninit(
            gpu.device(),
            matrices.len() as u32,
            BufferUsages::STORAGE | BufferUsages::COPY_SRC,
        );
        let staging: GpuVector<GpuSvd3> = GpuVector::uninit(
            gpu.device(),
            matrices.len() as u32,
            BufferUsages::MAP_READ | BufferUsages::COPY_DST,
        );

        // Queue the test.
        KernelInvocationBuilder::new(&mut queue, &svd)
            .bind0([inputs.buffer(), result.buffer()])
            .queue(matrices.len() as u32);

        // Run.
        queue.encode(&mut encoder, None);
        staging.copy_from(&mut encoder, &result);
        gpu.queue().submit(Some(encoder.finish()));

        // Check the result is correct.
        let gpu_result = staging.read(gpu.device()).await.unwrap();

        for (m, svd) in matrices.iter().zip(gpu_result.iter()) {
            let m = m.fixed_rows::<3>(0).into_owned();
            let reconstructed = svd.u.fixed_rows::<3>(0).into_owned()
                * Matrix3::from_diagonal(&svd.s.fixed_rows::<3>(0))
                * svd.vt.fixed_rows::<3>(0).into_owned();
            assert_relative_eq!(m, reconstructed, epsilon = 1.0e-4);
        }
    }
}