gfx-rs / wgpu

A cross-platform, safe, pure-Rust graphics API.
https://wgpu.rs
Apache License 2.0
12.14k stars 887 forks source link

Restrict usage of `f64` with some built-in functions #6087

Open Jianqoq opened 1 month ago

Jianqoq commented 1 month ago

Description shader f64 can't compute sin, cos, tan correctly, f32 can. sinh, cosh, tanh can calculate correctly

Repro steps

@group(0) @binding(0) var<storage, read> a : array<f64>;
@group(0) @binding(1) var<storage, read_write> c : array<f64>;

@compute
@workgroup_size(16, 16, 1)
fn main(
@builtin(workgroup_id) workgroup_id : vec3 <u32>,
@builtin(local_invocation_id) local_id : vec3 <u32>
)
{
    let local_row_offset: u32 = local_id.x * 16 * 1024 + workgroup_id.x * 16 * 1024;
    let local_col_offset: u32 = workgroup_id.y * 16 + local_id.y;
    let t_id : u32 = local_row_offset + local_col_offset;

    let num_threads : u32 = u32(1024) * u32(16) * u32(1024) * u32(16);
    let tmp : u32 = 5 % num_threads;
    let start_idx : u32 = t_id * (5 / num_threads) + min(t_id, tmp);
    let end_idx : u32 = start_idx + (5 / num_threads) + u32(t_id < tmp);

    if end_idx - start_idx == 0 {
        return;
    }

    let range : u32 = end_idx - start_idx;
    for (var i: u32 = 0; i < range; i++) {
        c[t_id + i] = f64(sin(f64(a[t_id + i])));
    }
}

Expected vs observed behavior for [0, 1, 2, 3, 4] f32 I can get [0. 0.84 0.90 0.14 -0.7], but for f64 I get [0.0, 2.121995791e-314, 4.243991582e-314, 6.365987373e-314, 8.487983164e-314]

Platform windows11, vulkan | DX12, lastest wgpu

teoxoy commented 1 month ago

windows11, vulkan | DX12, lastest wgpu

The SHADER_F64 feature is only implemented on Vulkan right now.

teoxoy commented 1 month ago

Looking at the spec https://registry.khronos.org/SPIR-V/specs/unified1/GLSL.std.450.html, there are a bunch of functions that only support "16-bit or 32-bit floating-point" numbers.

We should validate that those are only used with 16-bit or 32-bit floating-point numbers.