Upload shaders/bf16_to_f32.wgsl with huggingface_hub

shaders/bf16_to_f32.wgsl

@@ -0,0 +1,36 @@
+// bf16_to_f32.wgsl — convert BF16 weights to F32 in-place on GPU.
+// BF16 is F32 with the lower 16 bits zeroed: bf16_as_u16 << 16 = f32_bits.
+//
+// Input: array of u32 where each u32 holds TWO bf16 values (packed little-endian).
+//   u32 = (bf16_high << 16) | bf16_low
+// Output: array of f32, twice the length of the input u32 array.
+//
+// Dispatch: ceil(num_f32_elements / 2 / 64) workgroups.
+// Each thread converts one u32 (2 bf16 values) → 2 f32 values.
+
+const WG_SIZE: u32 = 64u;
+
+struct Params {
+    num_pairs: u32,  // number of u32 elements (each holds 2 bf16 values)
+}
+
+@group(0) @binding(0) var<storage, read>       bf16_packed: array<u32>;   // [num_pairs]
+@group(0) @binding(1) var<storage, read_write> f32_out: array<f32>;       // [num_pairs * 2]
+
+@group(1) @binding(0) var<uniform> params: Params;
+
+@compute @workgroup_size(WG_SIZE)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let i = gid.x;
+    if (i >= params.num_pairs) { return; }
+
+    let packed = bf16_packed[i];
+
+    // Low 16 bits = first bf16, high 16 bits = second bf16
+    let bf16_low  = packed & 0xFFFFu;
+    let bf16_high = (packed >> 16u) & 0xFFFFu;
+
+    // BF16 → F32: shift left by 16 bits
+    f32_out[i * 2u]     = bitcast<f32>(bf16_low << 16u);
+    f32_out[i * 2u + 1u] = bitcast<f32>(bf16_high << 16u);
+}