Vulkan backend is much slower than dx12 (consumed time is 6x!)

[bczhc]

Hi! I made a sha256 brute force stuff (or a miner simulator) to find an input producing digests starting with zero bits. The code is shown at the bottom. I'm quite new to webgpu/wgpu and am trying to make some things. The code may be a little dumb somehow. However anyway, to my very surprise, DX12 backend of this program on Windows 10, runs much more faster than Vulkan on Windows/Linux!

Use this cmdline to test it:

❯ WGPU_BACKEND=vulkan ./sha256-miner -i64 -d32
# or change `WGPU_BACKEND` to dx12

My hardware is NVIDIA GeForce RTX 3060 Mobile. From my tests, on both Linux/Windows, the Vulkan backend finished the work in about 16s. However the DX12 backend did it only in ~2.5s (startup time not counted). I've played around with it and seems DX12 is trying to unfold for-iterations in the compute shader, so when you set a higher -i parameter, the startup time will become larger. But still, I'm not expecting such a big difference.

Another thing that is weird: when the dx12 version runs, Windows Task Manager only keeps showing a very low CPU and GPU usage (can be < 10% e.g.) (however GPU fans start spinning plus nvidia-smi reports 100%). Also the desktop is smooth (cursor moving, etc.). As a comparison, when the Vulkan version runs, Task Manager reports 100% as expected, and the desktop becomes a bit laggy.

I don't know if dx12 does some magical optimizations for my naive code. Or, does that mean, if I write the code in a more optimized way, the Vulkan version could also run in ~2.5 seconds?

You can run the program below; observe the dispatch: xx, start: xx lines also get printed faster than Vulkan's. It's not a fake illusion - DX12 runs really much faster than Vulkan here. I want an explanation.. Appreciate it!

---

Code

#!/usr/bin/env rust-script
//! ```cargo
//! [dependencies]
//! bytemuck = "1.24.0"
//! tokio = { version = "1.48.0", features = ["full"] }
//! sha2 = "0.11.0-rc.5"
//! clap = { version = "4.5.57", features = ["derive"] }
//! wgpu = "28.0.0"
//! anyhow = "1.0.100"
//! hex = "0.4.3"
//! ```
use bytemuck::{cast_slice, cast_slice_mut};
use std::process::exit;
use std::time::Instant;
use tokio::sync::oneshot;
use wgpu::wgt::PollType;
use wgpu::{
    Backends, BindGroup, BindGroupDescriptor, BindGroupEntry, BindingResource, Buffer,
    BufferBinding, BufferDescriptor, BufferUsages, ComputePipeline, ComputePipelineDescriptor,
    Device, Instance, InstanceDescriptor, MapMode, PipelineCompilationOptions, Queue,
    ShaderModuleDescriptor, ShaderSource,
};


/// Sha256 buffer type the shader uses.
type FatSha256Buf = [u32; SHA256_BYTES];

const SHA256_BYTES: usize = 32;
const INPUT_SIZE: usize = 32;
/// The shader treats `u32`s as `u8`s.
const BLOCK_BUFFER_IN_SHADER: u64 = size_of::<FatSha256Buf>() as _;

use clap::Parser;
use sha2::Digest;

macro_rules! default {
    () => {Default::default()};
}

#[derive(Parser, Debug)]
#[command(about = "GPU Sha256 Miner Simulator")]
struct Args {
    /// Number of threads per workgroup (WORKGROUP_SIZE)
    #[arg(long, default_value_t = 256)]
    workgroup_size: u32,

    /// Number of workgroups to dispatch in the X dimension (DISPATCH_X)
    #[arg(long, default_value_t = 2048)]
    dispatch_x: u32,

    /// Number of hash iterations performed by each individual thread
    #[arg(short, long, default_value_t = 256)]
    iterations: u32,

    /// Target difficulty in bits
    #[arg(short, long, default_value_t = 32)]
    difficulty: u32,
}

struct State {
    device: Device,
    queue: Queue,
    pipeline: ComputePipeline,
    input_buffer: Buffer,
    result_buffer: Buffer,
    map_read_buffer: Buffer,
    bind_group: BindGroup,
}

impl State {
    async fn new(args: &Args) -> anyhow::Result<Self> {
        let instance = Instance::new(&InstanceDescriptor {
            backends: Backends::from_env().unwrap_or(default!()),
            ..default!()
        });
        let adapter = instance.request_adapter(&default!()).await?;
        let (device, queue) = adapter.request_device(&default!()).await?;

        let shader_module = device.create_shader_module(ShaderModuleDescriptor {
            label: None,
            source: ShaderSource::Wgsl(merged_wgsl_source(args.difficulty).into()),
        });
        let pipeline = device.create_compute_pipeline(&ComputePipelineDescriptor {
            label: None,
            layout: None,
            module: &shader_module,
            entry_point: None,
            compilation_options: PipelineCompilationOptions {
                constants: &[
                    ("WORKGROUP_SIZE", args.workgroup_size as f64),
                    ("ITERATIONS_PER_THREAD", args.iterations as f64),
                    (
                        "RUNS_PER_DISPATCH",
                        (args.dispatch_x * args.workgroup_size) as f64,
                    ),
                    ("DIFFICULTY_BITS", args.difficulty as f64),
                ],
                zero_initialize_workgroup_memory: false,
            },
            cache: None,
        });

        let input_buffer = device.create_buffer(&BufferDescriptor {
            label: None,
            size: INPUT_SIZE as u64 * 4,
            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });
        let result_buffer = device.create_buffer(&BufferDescriptor {
            label: None,
            size: BLOCK_BUFFER_IN_SHADER,
            usage: BufferUsages::STORAGE | BufferUsages::COPY_DST | BufferUsages::COPY_SRC,
            mapped_at_creation: false,
        });
        let map_read_buffer = device.create_buffer(&BufferDescriptor {
            label: None,
            size: BLOCK_BUFFER_IN_SHADER,
            usage: BufferUsages::MAP_READ | BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let bind_group = device.create_bind_group(&BindGroupDescriptor {
            label: None,
            layout: &pipeline.get_bind_group_layout(0),
            entries: &[
                BindGroupEntry {
                    binding: 0,
                    resource: BindingResource::Buffer(BufferBinding {
                        buffer: &input_buffer,
                        offset: 0,
                        size: None,
                    }),
                },
                BindGroupEntry {
                    binding: 1,
                    resource: BindingResource::Buffer(BufferBinding {
                        buffer: &result_buffer,
                        offset: 0,
                        size: None,
                    }),
                },
            ],
        });

        Ok(Self {
            queue,
            device,
            pipeline,
            input_buffer,
            bind_group,
            result_buffer,
            map_read_buffer,
        })
    }

    fn write_input_data(&self, buf: &[u8]) {
        let mut input_data = [0_u32; INPUT_SIZE];
        for (i, &b) in buf.iter().enumerate() {
            input_data[i] = b as _;
        }
        self.queue
            .write_buffer(&self.input_buffer, 0, cast_slice(&input_data));
    }

    fn compute_dispatch(&self, workgroups_x: u32) {
        let mut encoder = self.device.create_command_encoder(&default!());

        let mut pass = encoder.begin_compute_pass(&default!());
        pass.set_pipeline(&self.pipeline);
        pass.set_bind_group(0, &self.bind_group, default!());
        pass.dispatch_workgroups(workgroups_x, 1, 1);
        drop(pass);

        encoder.copy_buffer_to_buffer(&self.result_buffer, 0, &self.map_read_buffer, 0, None);

        let command_buffer = encoder.finish();
        self.queue.submit([command_buffer]);
    }

    async fn read_result(&self, to: &mut [u8]) -> anyhow::Result<()> {
        let (tx, rx) = oneshot::channel();
        self.map_read_buffer.map_async(MapMode::Read, .., |e| {
            tx.send(e).unwrap();
        });
        self.device.poll(PollType::Wait {
            submission_index: None,
            timeout: None,
        })?;
        rx.await??;

        to[..(self.map_read_buffer.size() as usize)]
            .copy_from_slice(cast_slice(&*self.map_read_buffer.get_mapped_range(..)));
        self.map_read_buffer.unmap();
        Ok(())
    }
}

fn add_big_int(data: &mut [u8; 32], mut n: u32) {
    let mut carry = n;

    for byte in data.iter_mut() {
        if carry == 0 {
            break;
        }

        let sum = *byte as u32 + carry;
        *byte = (sum & 0xFF) as u8;
        carry = sum >> 8;
    }
}

#[inline(always)]
fn convert_fat_buf(buf: &FatSha256Buf) -> [u8; SHA256_BYTES] {
    buf.map(|x| x as u8)
}

fn generate_check_difficulty_wgsl(difficulty_bits: u32) -> String {
    let mut conditions = Vec::new();

    let full_bytes = difficulty_bits / 8;
    for i in 0..full_bytes {
        conditions.push(format!("buf[{}] == 0u", i));
    }

    let remaining_bits = difficulty_bits % 8;
    if remaining_bits > 0 {
        let shift = 8 - remaining_bits;
        conditions.push(format!("(buf[{}] >> {}u) == 0u", full_bytes, shift));
    }

    let final_condition = if conditions.is_empty() {
        "true".to_string()
    } else {
        conditions.join(" && ")
    };

    format!(
        r#"
fn check_difficulty(buf: ptr<function, array<u32, SHA256_BLOCK_SIZE>>) -> bool {{
    return {};
}}
"#,
        final_condition
    )
}

fn merged_wgsl_source(difficulty_bits: u32) -> String {
    let mut source = wgsl_source::WGSL_SOURCE.lines().collect::<Vec<_>>();
    source.remove(0);
    source.remove(0);
    let generated = generate_check_difficulty_wgsl(difficulty_bits);
    for x in generated.lines().into_iter().rev() {
        source.insert(0, x);
    }
    source.join("\n")
}

fn print_result_and_exit(buf: FatSha256Buf, start: Instant) {
    use sha2::Digest;
    let mut hasher = sha2::Sha256::new();
    hasher.update(convert_fat_buf(&buf));
    let hash = hex::encode(hasher.finalize());

    println!("Result:");
    println!("  input: {}", hex::encode(convert_fat_buf(&buf)));
    println!("  sha256: {}", hash);
    println!("  elapsed: {:?}", start.elapsed());
    exit(0);
}

#[tokio::main]
async fn main() -> anyhow::Result<()> {
    let args = Args::parse();
    let runs_per_dispatch = args.dispatch_x * args.workgroup_size;

    println!("Args: {:?}", args);

    let state = State::new(&args).await?;
    let mut input_data = [0_u8; INPUT_SIZE];
    let mut result = [0_u32; SHA256_BYTES];
    let mut counter = 0_usize;
    let start = Instant::now();
    loop {
        println!("dispatch: {}, start: {}", counter, hex::encode(input_data));
        state.write_input_data(&input_data);
        state.compute_dispatch(args.dispatch_x);
        add_big_int(&mut input_data, runs_per_dispatch * args.iterations);
        state.read_result(cast_slice_mut(&mut result)).await?;
        if result.iter().any(|x| *x != 0) {
            print_result_and_exit(result, start);
        }
        counter += 1;
    }
}

mod wgsl_source {
   /// Thanks to https://github.com/dezmou/SHA256-WebGPU.
    pub const WGSL_SOURCE: &str = r#"
fn check_difficulty(buf: ptr<function, array<u32, SHA256_BLOCK_SIZE>>) -> bool { /* stub */; return true; }

override WORKGROUP_SIZE = 0u;
override ITERATIONS_PER_THREAD = 0u;
override RUNS_PER_DISPATCH = 0u;
override DIFFICULTY_BITS = 0u;

struct SHA256_CTX {
    data : array<u32, 64>,
    datalen : u32,
    bitlen : array<u32, 2>,
    state : array<u32, 8>,
    info : u32,
  };

  @group(0) @binding(0) var<storage, read> start : array<u32>;
  @group(0) @binding(1) var<storage, read_write> result : array<u32>;

  const SHA256_BLOCK_SIZE = 32;
  const INPUT_SIZE = 32;

  const k = array<u32, 64> (
    0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
    0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174,
    0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da,
    0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967,
    0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85,
    0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070,
    0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3,
    0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2
  );

  fn ROTLEFT(a : u32, b : u32) -> u32{return (((a) << (b)) | ((a) >> (32-(b))));}
  fn ROTRIGHT(a : u32, b : u32) -> u32{return (((a) >> (b)) | ((a) << (32-(b))));}

  fn CH(x : u32, y : u32, z : u32) -> u32{return (((x) & (y)) ^ (~(x) & (z)));}
  fn MAJ(x : u32, y : u32, z : u32) -> u32{return (((x) & (y)) ^ ((x) & (z)) ^ ((y) & (z)));}
  fn EP0(x : u32) -> u32{return (ROTRIGHT(x,2) ^ ROTRIGHT(x,13) ^ ROTRIGHT(x,22));}
  fn EP1(x : u32) -> u32{return (ROTRIGHT(x,6) ^ ROTRIGHT(x,11) ^ ROTRIGHT(x,25));}
  fn SIG0(x : u32) -> u32{return (ROTRIGHT(x,7) ^ ROTRIGHT(x,18) ^ ((x) >> 3));}
  fn SIG1(x : u32) -> u32{return (ROTRIGHT(x,17) ^ ROTRIGHT(x,19) ^ ((x) >> 10));}

  fn sha256_transform(ctx : ptr<function, SHA256_CTX>)
  {
    var a : u32;
    var b : u32;
    var c : u32;
    var d : u32;
    var e : u32;
    var f : u32;
    var g : u32;
    var h : u32;
    var i : u32 = 0;
    var j : u32 = 0;
    var t1 : u32;
    var t2 : u32;
    var m : array<u32, 64> ;


    while(i < 16) {
      m[i] = ((*ctx).data[j] << 24) | ((*ctx).data[j + 1] << 16) | ((*ctx).data[j + 2] << 8) | ((*ctx).data[j + 3]);
      i++;
      j += 4;
    }

    while(i < 64) {
          m[i] = SIG1(m[i - 2]) + m[i - 7] + SIG0(m[i - 15]) + m[i - 16];
      i++;
    }

    a = (*ctx).state[0];
    b = (*ctx).state[1];
    c = (*ctx).state[2];
    d = (*ctx).state[3];
    e = (*ctx).state[4];
    f = (*ctx).state[5];
    g = (*ctx).state[6];
    h = (*ctx).state[7];

    i = 0;
    for (; i < 64; i++) {
      t1 = h + EP1(e) + CH(e,f,g) + k[i] + m[i];
      t2 = EP0(a) + MAJ(a,b,c);
      h = g;
      g = f;
      f = e;
      e = d + t1;
      d = c;
      c = b;
      b = a;
      a = t1 + t2;
    }


    (*ctx).state[0] += a;
    (*ctx).state[1] += b;
    (*ctx).state[2] += c;
    (*ctx).state[3] += d;
    (*ctx).state[4] += e;
    (*ctx).state[5] += f;
    (*ctx).state[6] += g;
    (*ctx).state[7] += h;
  }


  fn sha256_update(ctx : ptr<function, SHA256_CTX>, input_data: ptr<function, array<u32, INPUT_SIZE>>, len : u32)
  {
    for (var i :u32 = 0; i < len; i++) {
      (*ctx).data[(*ctx).datalen] = input_data[i];
      (*ctx).datalen++;
      if ((*ctx).datalen == 64) {
        sha256_transform(ctx);

        if ((*ctx).bitlen[0] > 0xffffffff - (512)){
          (*ctx).bitlen[1]++;
        }
        (*ctx).bitlen[0] += 512;


        (*ctx).datalen = 0;
      }
    }
  }

  fn sha256_final(ctx : ptr<function, SHA256_CTX>, hash:  ptr<function, array<u32, SHA256_BLOCK_SIZE>>  )
  {
    var i : u32 = (*ctx).datalen;

    if ((*ctx).datalen < 56) {
      (*ctx).data[i] = 0x80;
        i++;
      while (i < 56){
        (*ctx).data[i] = 0x00;
        i++;
      }
    }
    else {
      (*ctx).data[i] = 0x80;
      i++;
      while (i < 64){
        (*ctx).data[i] = 0x00;
        i++;
      }
      sha256_transform(ctx);
      for (var i = 0; i < 56 ; i++) {
        (*ctx).data[i] = 0;
      }
    }

    if ((*ctx).bitlen[0] > 0xffffffff - (*ctx).datalen * 8) {
      (*ctx).bitlen[1]++;
    }
    (*ctx).bitlen[0] += (*ctx).datalen * 8;


    (*ctx).data[63] = (*ctx).bitlen[0];
    (*ctx).data[62] = (*ctx).bitlen[0] >> 8;
    (*ctx).data[61] = (*ctx).bitlen[0] >> 16;
    (*ctx).data[60] = (*ctx).bitlen[0] >> 24;
    (*ctx).data[59] = (*ctx).bitlen[1];
    (*ctx).data[58] = (*ctx).bitlen[1] >> 8;
    (*ctx).data[57] = (*ctx).bitlen[1] >> 16;
    (*ctx).data[56] = (*ctx).bitlen[1] >> 24;
    sha256_transform(ctx);

    for (i = 0; i < 4; i++) {
      (*hash)[i] = ((*ctx).state[0] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 4] = ((*ctx).state[1] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 8] = ((*ctx).state[2] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 12] = ((*ctx).state[3] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 16] = ((*ctx).state[4] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 20] = ((*ctx).state[5] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 24] = ((*ctx).state[6] >> (24 - i * 8)) & 0x000000ff;
      (*hash)[i + 28] = ((*ctx).state[7] >> (24 - i * 8)) & 0x000000ff;
    }
  }

  fn sha256_init(ctx : ptr<function, SHA256_CTX>) {
    (*ctx).datalen = 0;
    (*ctx).bitlen[0] = 0;
    (*ctx).bitlen[1] = 0;
    (*ctx).state[0] = 0x6a09e667;
    (*ctx).state[1] = 0xbb67ae85;
    (*ctx).state[2] = 0x3c6ef372;
    (*ctx).state[3] = 0xa54ff53a;
    (*ctx).state[4] = 0x510e527f;
    (*ctx).state[5] = 0x9b05688c;
    (*ctx).state[6] = 0x1f83d9ab;
    (*ctx).state[7] = 0x5be0cd19;
  }

fn add_big_int_local(p_data: ptr<function, array<u32, 32>>, n: u32) {
    var carry = n;

    for (var i = 0u; i < 32u; i++) {
        if (carry == 0u) { break; }

        let sum = (*p_data)[i] + carry;
        (*p_data)[i] = sum & 255u;
        carry = sum >> 8u;
    }
}

@compute @workgroup_size(WORKGROUP_SIZE)
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
    let _r = result[0];

    for (var i = 0u; i < ITERATIONS_PER_THREAD; i += 1) {
      let addition = i * RUNS_PER_DISPATCH + global_id.x;
      var this_input: array<u32, SHA256_BLOCK_SIZE>;
      for (var j = 0u; j < INPUT_SIZE; j += 1) {
        this_input[j] = start[j];
      }
      add_big_int_local(&this_input, addition);
      var ctx : SHA256_CTX;
      sha256_init(&ctx);
      var buf : array<u32, SHA256_BLOCK_SIZE>;
      sha256_update(&ctx, &this_input, INPUT_SIZE);
      sha256_final(&ctx, &buf);

      if check_difficulty(&buf) {
        for (var j = 0u; j < 32; j += 1) {
          // result[j] = buf[j];
          result[j] = this_input[j];
        }
      }
    }
}
    "#;
}

[inner-daemons]

Interesting stuff. I assume this replicated across many runs? Since this would be inherently random to my knowledge.

This might be a thing with DXC optimizing very well. If you try to manually use naga to compile, and then invoke DXC with a low optimization level, and pass this through using passthrough shaders, does that give a slower DX12 run?

[inner-daemons]

Ok, its probably an issue with the pipeline layout, not really sure. This is probably a me bug, if you could get a minimal reproducible example and file an issue that'd be great.

[bczhc]

I've created a reduced reproduction code. The error occurs when the storage buffer is read-write. If changing ty: BufferBindingType::Storage { read_only: false } to ty: BufferBindingType::Storage { read_only: true } and read_write to read in the Rust and WGSL code respectively, the error won't happen. Since I have zero experience with dx12 so not sure if this is necessarily a wgpu bug, or if I missed something.

use bytemuck::bytes_of;
use std::borrow::Cow;
use wgpu::{
    Backends, BindGroupLayoutDescriptor, BindGroupLayoutEntry,
    BindingType, BufferBindingType, BufferDescriptor, BufferUsages, ComputePipelineDescriptor,
    DeviceDescriptor, Features, Instance, InstanceDescriptor, PipelineLayoutDescriptor,
    ShaderModuleDescriptorPassthrough, ShaderStages,
};

async fn async_main() -> anyhow::Result<()> {
    let instance = Instance::new(InstanceDescriptor {
        backends: Backends::from_env().unwrap_or_default(),
        ..Default::default()
    });

    let adapter = instance.request_adapter(&Default::default()).await?;
    let (device, queue) = adapter
        .request_device(&DeviceDescriptor {
            required_features: Features::PASSTHROUGH_SHADERS,
            ..Default::default()
        })
        .await?;

    let data_buffer = device.create_buffer(&BufferDescriptor {
        label: None,
        size: 4,
        usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
        mapped_at_creation: false,
    });
    queue.write_buffer(&data_buffer, 0, &bytes_of(&[0_u32]));

    let bind_group_layout = device.create_bind_group_layout(&BindGroupLayoutDescriptor {
        label: None,
        entries: &[BindGroupLayoutEntry {
            binding: 0,
            visibility: ShaderStages::COMPUTE,
            ty: BindingType::Buffer {
                ty: BufferBindingType::Storage { read_only: false },
                has_dynamic_offset: false,
                min_binding_size: None,
            },
            count: None,
        }],
    });
    let pipeline_layout = device.create_pipeline_layout(&PipelineLayoutDescriptor {
        label: None,
        bind_group_layouts: &[Some(&bind_group_layout)],
        immediate_size: 0,
    });

    #[cfg(false)]
    let shader_module = device.create_shader_module(wgpu::include_wgsl!("../compute.wgsl"));

    #[cfg(true)]
    let shader_module = {
        let desc = ShaderModuleDescriptorPassthrough {
            label: None,
            num_workgroups: (256, 0, 0),
            spirv: None,
            dxil: Some(Cow::Borrowed(include_bytes!("../compute.dxil"))),
            hlsl: None,
            metallib: None,
            msl: None,
            glsl: None,
            wgsl: None,
        };
        unsafe { device.create_shader_module_passthrough(desc) }
    };

    // This will fail.
    let pipeline = device.create_compute_pipeline(&ComputePipelineDescriptor {
        label: None,
        layout: Some(&pipeline_layout),
        module: &shader_module,
        entry_point: Some("main"),
        compilation_options: Default::default(),
        cache: None,
    });

    Ok(())
}

fn main() -> anyhow::Result<()> {
    env_logger::init();
    pollster::block_on(async_main())?;
    Ok(())
}

compute.wgsl:

@group(0) @binding(0) var<storage, read_write> data: array<u32>;

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) gid: vec3u) {
    let _i = gid.x;
    let _a = data[0];
}

The used wgpu is the latest trunk commit. Reproduction steps are similar to what it writes in my previous comments - naga, then dxc, then WGPU_BACKEND=dx12 cargo run.

[inner-daemons]

@bczhc Ok I'll check that out at some point, thanks.

[inner-daemons]

Yeah ok this is just a glaring issue with passthrough shaders, I haven't yet found a great way to expose bindings and stuff other than our own precompiled shader system with lots of additional info. So probably using any bindings at all is straightup undefined behavior on some backends (SPIRV is probably fine).

For the original issue, I am still willing to bet that DXC is responsible here, since most of the time is presumably taken on the GPU. I wanted to do investigations with spirv-opt but I doubt that would help too much.

[bczhc]

So probably using any bindings at all is straightup undefined behavior on some backends (SPIRV is probably fine).

Then I tried to use some alternative ways - for example push_constants. Snippet:

...
struct Constants {
    start: array<u32, 32>,
}
var<immediate> constants: Constants;
...

btw, encountered an oops - naga complains and just panics:

❯ naga sha256-miner-d32.wgsl
Validation successful
❯ naga sha256-miner-d32.wgsl a.hlsl

thread 'main' (2826842) panicked at /mnt/nvme/open-source/wgpu/naga/src/back/hlsl/writer.rs:1056:18:
No bind target was defined for the immediates block
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace