Road floats to prevent deltas

Author: LegNeato

tests/difftests/lib/Cargo.toml

@@ -15,7 +15,7 @@ use-compiled-tools = [
     "spirv-builder/use-compiled-tools"
 ]
 
-[dependencies]
+[target.'cfg(not(target_arch = "spirv"))'.dependencies]
 spirv-builder.workspace = true
 serde = { version = "1.0", features = ["derive"] }
 serde_json = "1.0"

tests/difftests/lib/src/lib.rs

@@ -1,6 +1,21 @@
+#![cfg_attr(target_arch = "spirv", no_std)]
+
+#[cfg(not(target_arch = "spirv"))]
 pub mod config;
+#[cfg(not(target_arch = "spirv"))]
 pub mod scaffold;
 
+/// Macro to round a f32 value to 6 decimal places for cross-platform consistency
+/// in floating-point operations. This helps ensure difftest results are consistent
+/// across different platforms (Linux, Mac, Windows) which may have slight differences
+/// in floating-point implementations.
+#[macro_export]
+macro_rules! round6 {
+    ($v:expr) => {
+        (($v) * 1_000_000.0).round() / 1_000_000.0
+    };
+}
+
 #[cfg(test)]
 mod tests {
     use super::config::Config;

tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/Cargo.toml

@@ -10,11 +10,9 @@ crate-type = ["dylib"]
 
 # Common deps
 [dependencies]
-
-# GPU deps
 spirv-std.workspace = true
+difftest.workspace = true
 
 # CPU deps
 [target.'cfg(not(target_arch = "spirv"))'.dependencies]
-difftest.workspace = true
-bytemuck.workspace = true
+bytemuck.workspace = true

tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/lib.rs

@@ -1,5 +1,6 @@
 #![no_std]
 
+use difftest::round6;
 #[allow(unused_imports)]
 use spirv_std::num_traits::Float;
 use spirv_std::spirv;
@@ -24,28 +25,28 @@ pub fn main_cs(
     }
 
     // Basic arithmetic
-    output[base_offset + 0] = x + 1.5;
-    output[base_offset + 1] = x - 0.5;
-    output[base_offset + 2] = x * 2.0;
-    output[base_offset + 3] = x / 2.0;
-    output[base_offset + 4] = x % 3.0;
+    output[base_offset + 0] = round6!(x + 1.5);
+    output[base_offset + 1] = round6!(x - 0.5);
+    output[base_offset + 2] = round6!(x * 2.0);
+    output[base_offset + 3] = round6!(x / 2.0);
+    output[base_offset + 4] = round6!(x % 3.0);
 
     // Trigonometric functions (simplified for consistent results)
-    output[base_offset + 5] = x.sin();
-    output[base_offset + 6] = x.cos();
-    output[base_offset + 7] = x.tan().clamp(-10.0, 10.0);
+    output[base_offset + 5] = round6!(x.sin());
+    output[base_offset + 6] = round6!(x.cos());
+    output[base_offset + 7] = round6!(x.tan().clamp(-10.0, 10.0));
     output[base_offset + 8] = 0.0;
     output[base_offset + 9] = 0.0;
-    output[base_offset + 10] = x.atan();
+    output[base_offset + 10] = round6!(x.atan());
 
     // Exponential and logarithmic (simplified)
-    output[base_offset + 11] = x.exp().min(1e6);
-    output[base_offset + 12] = if x > 0.0 { x.ln() } else { -10.0 };
-    output[base_offset + 13] = x.abs().sqrt();
-    output[base_offset + 14] = x.abs().powf(2.0);
-    output[base_offset + 15] = if x > 0.0 { x.log2() } else { -10.0 };
-    output[base_offset + 16] = x.exp2().min(1e6);
-    output[base_offset + 17] = x.floor();
+    output[base_offset + 11] = round6!(x.exp().min(1e6));
+    output[base_offset + 12] = round6!(if x > 0.0 { x.ln() } else { -10.0 });
+    output[base_offset + 13] = round6!(x.abs().sqrt());
+    output[base_offset + 14] = round6!(x.abs() * x.abs()); // Use multiplication instead of powf
+    output[base_offset + 15] = round6!(if x > 0.0 { x.log2() } else { -10.0 });
+    output[base_offset + 16] = round6!(x.exp2().min(1e6));
+    output[base_offset + 17] = x.floor(); // floor/ceil/round are exact
     output[base_offset + 18] = x.ceil();
     output[base_offset + 19] = x.round();
 

tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/shader.wgsl

@@ -4,6 +4,11 @@ var<storage, read> input: array<f32>;
 @group(0) @binding(1)
 var<storage, read_write> output: array<f32>;
 
+// Helper function to round to 6 decimal places for cross-platform consistency
+fn round6(v: f32) -> f32 {
+    return round(v * 1000000.0) / 1000000.0;
+}
+
 @compute @workgroup_size(32, 1, 1)
 fn main_cs(@builtin(global_invocation_id) global_id: vec3<u32>) {
     let tid = global_id.x;
@@ -20,28 +25,28 @@ fn main_cs(@builtin(global_invocation_id) global_id: vec3<u32>) {
     }
 
     // Basic arithmetic
-    output[base_offset + 0u] = x + 1.5;
-    output[base_offset + 1u] = x - 0.5;
-    output[base_offset + 2u] = x * 2.0;
-    output[base_offset + 3u] = x / 2.0;
-    output[base_offset + 4u] = x % 3.0;
+    output[base_offset + 0u] = round6(x + 1.5);
+    output[base_offset + 1u] = round6(x - 0.5);
+    output[base_offset + 2u] = round6(x * 2.0);
+    output[base_offset + 3u] = round6(x / 2.0);
+    output[base_offset + 4u] = round6(x % 3.0);
 
     // Trigonometric functions (simplified for consistent results)
-    output[base_offset + 5u] = sin(x);
-    output[base_offset + 6u] = cos(x);
-    output[base_offset + 7u] = clamp(tan(x), -10.0, 10.0);
+    output[base_offset + 5u] = round6(sin(x));
+    output[base_offset + 6u] = round6(cos(x));
+    output[base_offset + 7u] = round6(clamp(tan(x), -10.0, 10.0));
     output[base_offset + 8u] = 0.0;
     output[base_offset + 9u] = 0.0;
-    output[base_offset + 10u] = atan(x);
+    output[base_offset + 10u] = round6(atan(x));
 
     // Exponential and logarithmic (simplified)
-    output[base_offset + 11u] = min(exp(x), 1e6);
-    output[base_offset + 12u] = select(-10.0, log(x), x > 0.0);
-    output[base_offset + 13u] = sqrt(abs(x));
-    output[base_offset + 14u] = pow(abs(x), 2.0);
-    output[base_offset + 15u] = select(-10.0, log2(x), x > 0.0);
-    output[base_offset + 16u] = min(exp2(x), 1e6);
-    output[base_offset + 17u] = floor(x);
+    output[base_offset + 11u] = round6(min(exp(x), 1e6));
+    output[base_offset + 12u] = round6(select(-10.0, log(x), x > 0.0));
+    output[base_offset + 13u] = round6(sqrt(abs(x)));
+    output[base_offset + 14u] = round6(abs(x) * abs(x)); // Use multiplication instead of pow
+    output[base_offset + 15u] = round6(select(-10.0, log2(x), x > 0.0));
+    output[base_offset + 16u] = round6(min(exp2(x), 1e6));
+    output[base_offset + 17u] = floor(x); // floor/ceil/round are exact
     output[base_offset + 18u] = ceil(x);
     output[base_offset + 19u] = round(x);
 

tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-rust/Cargo.toml

@@ -10,11 +10,9 @@ crate-type = ["dylib"]
 
 # Common deps
 [dependencies]
-
-# GPU deps
 spirv-std.workspace = true
+difftest.workspace = true
 
 # CPU deps
 [target.'cfg(not(target_arch = "spirv"))'.dependencies]
-difftest.workspace = true
-bytemuck.workspace = true
+bytemuck.workspace = true

tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-rust/src/lib.rs

@@ -1,5 +1,6 @@
 #![no_std]
 
+use difftest::round6;
 use spirv_std::glam::{Mat2, Mat3, Mat4, UVec3, Vec2, Vec3, Vec4};
 #[allow(unused_imports)]
 use spirv_std::num_traits::Float;
@@ -35,10 +36,10 @@ pub fn main_cs(
 
     // Mat2 multiplication
     let m2_mul = m2a * m2b;
-    output[base_offset + 0] = m2_mul.col(0).x;
-    output[base_offset + 1] = m2_mul.col(0).y;
-    output[base_offset + 2] = m2_mul.col(1).x;
-    output[base_offset + 3] = m2_mul.col(1).y;
+    output[base_offset + 0] = round6!(m2_mul.col(0).x);
+    output[base_offset + 1] = round6!(m2_mul.col(0).y);
+    output[base_offset + 2] = round6!(m2_mul.col(1).x);
+    output[base_offset + 3] = round6!(m2_mul.col(1).y);
 
     // Mat2 transpose
     let m2_transpose = m2a.transpose();
@@ -48,29 +49,29 @@ pub fn main_cs(
     output[base_offset + 7] = m2_transpose.col(1).y;
 
     // Mat2 determinant (with rounding for consistency)
-    output[base_offset + 8] = (m2a.determinant() * 1000.0).round() / 1000.0;
+    output[base_offset + 8] = round6!(m2a.determinant());
 
     // Mat2 * Vec2
     let v2 = Vec2::new(1.0, 2.0);
     let m2_v2 = m2a * v2;
-    output[base_offset + 9] = m2_v2.x;
-    output[base_offset + 10] = m2_v2.y;
+    output[base_offset + 9] = round6!(m2_v2.x);
+    output[base_offset + 10] = round6!(m2_v2.y);
 
     // Mat3 operations
     let m3a = Mat3::from_cols(Vec3::new(a, b, c), Vec3::new(b, c, d), Vec3::new(c, d, a));
     let m3b = Mat3::from_cols(Vec3::new(d, c, b), Vec3::new(c, b, a), Vec3::new(b, a, d));
 
     // Mat3 multiplication
     let m3_mul = m3a * m3b;
-    output[base_offset + 11] = m3_mul.col(0).x;
-    output[base_offset + 12] = m3_mul.col(0).y;
-    output[base_offset + 13] = m3_mul.col(0).z;
-    output[base_offset + 14] = m3_mul.col(1).x;
-    output[base_offset + 15] = m3_mul.col(1).y;
-    output[base_offset + 16] = m3_mul.col(1).z;
-    output[base_offset + 17] = m3_mul.col(2).x;
-    output[base_offset + 18] = m3_mul.col(2).y;
-    output[base_offset + 19] = m3_mul.col(2).z;
+    output[base_offset + 11] = round6!(m3_mul.col(0).x);
+    output[base_offset + 12] = round6!(m3_mul.col(0).y);
+    output[base_offset + 13] = round6!(m3_mul.col(0).z);
+    output[base_offset + 14] = round6!(m3_mul.col(1).x);
+    output[base_offset + 15] = round6!(m3_mul.col(1).y);
+    output[base_offset + 16] = round6!(m3_mul.col(1).z);
+    output[base_offset + 17] = round6!(m3_mul.col(2).x);
+    output[base_offset + 18] = round6!(m3_mul.col(2).y);
+    output[base_offset + 19] = round6!(m3_mul.col(2).z);
 
     // Mat3 transpose - store just diagonal elements
     let m3_transpose = m3a.transpose();
@@ -79,14 +80,14 @@ pub fn main_cs(
     output[base_offset + 22] = m3_transpose.col(2).z;
 
     // Mat3 determinant (with rounding for consistency)
-    output[base_offset + 23] = (m3a.determinant() * 1000.0).round() / 1000.0;
+    output[base_offset + 23] = round6!(m3a.determinant());
 
     // Mat3 * Vec3 (with rounding for consistency)
     let v3 = Vec3::new(1.0, 2.0, 3.0);
     let m3_v3 = m3a * v3;
-    output[base_offset + 24] = (m3_v3.x * 10000.0).round() / 10000.0;
-    output[base_offset + 25] = (m3_v3.y * 10000.0).round() / 10000.0;
-    output[base_offset + 26] = (m3_v3.z * 10000.0).round() / 10000.0;
+    output[base_offset + 24] = round6!(m3_v3.x);
+    output[base_offset + 25] = round6!(m3_v3.y);
+    output[base_offset + 26] = round6!(m3_v3.z);
 
     // Mat4 operations
     let m4a = Mat4::from_cols(
@@ -104,10 +105,10 @@ pub fn main_cs(
 
     // Mat4 multiplication (just store diagonal for brevity)
     let m4_mul = m4a * m4b;
-    output[base_offset + 27] = m4_mul.col(0).x;
-    output[base_offset + 28] = m4_mul.col(1).y;
-    output[base_offset + 29] = m4_mul.col(2).z;
-    output[base_offset + 30] = m4_mul.col(3).w;
+    output[base_offset + 27] = round6!(m4_mul.col(0).x);
+    output[base_offset + 28] = round6!(m4_mul.col(1).y);
+    output[base_offset + 29] = round6!(m4_mul.col(2).z);
+    output[base_offset + 30] = round6!(m4_mul.col(3).w);
 
     // Mat4 transpose (just store diagonal)
     let m4_transpose = m4a.transpose();
@@ -117,15 +118,15 @@ pub fn main_cs(
     output[base_offset + 34] = m4_transpose.col(3).w;
 
     // Mat4 determinant (with rounding for consistency)
-    output[base_offset + 35] = (m4a.determinant() * 1000.0).round() / 1000.0;
+    output[base_offset + 35] = round6!(m4a.determinant());
 
     // Mat4 * Vec4 (with rounding for consistency)
     let v4 = Vec4::new(1.0, 2.0, 3.0, 4.0);
     let m4_v4 = m4a * v4;
-    output[base_offset + 36] = (m4_v4.x * 10000.0).round() / 10000.0;
-    output[base_offset + 37] = (m4_v4.y * 10000.0).round() / 10000.0;
-    output[base_offset + 38] = (m4_v4.z * 10000.0).round() / 10000.0;
-    output[base_offset + 39] = (m4_v4.w * 10000.0).round() / 10000.0;
+    output[base_offset + 36] = round6!(m4_v4.x);
+    output[base_offset + 37] = round6!(m4_v4.y);
+    output[base_offset + 38] = round6!(m4_v4.z);
+    output[base_offset + 39] = round6!(m4_v4.w);
 
     // Identity matrices
     output[base_offset + 40] = Mat2::IDENTITY.col(0).x;
@@ -135,8 +136,8 @@ pub fn main_cs(
     // Matrix inverse
     if m2a.determinant().abs() > 0.0001 {
         let m2_inv = m2a.inverse();
-        output[base_offset + 43] = m2_inv.col(0).x;
-        output[base_offset + 44] = m2_inv.col(1).y;
+        output[base_offset + 43] = round6!(m2_inv.col(0).x);
+        output[base_offset + 44] = round6!(m2_inv.col(1).y);
     } else {
         output[base_offset + 43] = 0.0;
         output[base_offset + 44] = 0.0;