Slightly simpler version of atomic test

Takes the loop out, each thread does either load or store, not both.

This still fails, but doesn't test as many opportunities.

Author: raphlinus

compute-shader-hello/src/main.rs

@@ -22,10 +22,6 @@ use wgpu::util::DeviceExt;
 
 use bytemuck;
 
-// A strategy of 0 is just atomic loads.
-// A strategy of 1 replaces the flag load with an atomicOr.
-const STRATEGY: u32 = 0;
-
 const USE_SPIRV: bool = false;
 
 async fn run() {
@@ -83,7 +79,7 @@ async fn run() {
     });
     let config_buf =  device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
         label: None,
-        contents: bytemuck::bytes_of(&[STRATEGY, 0]),
+        contents: bytemuck::bytes_of(&[0u32]),
         usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::MAP_READ,
     });
     // This works if the buffer is initialized, otherwise reads all 0, for some reason.
@@ -152,12 +148,11 @@ async fn run() {
         let mut cpass = encoder.begin_compute_pass(&Default::default());
         cpass.set_pipeline(&pipeline);
         cpass.set_bind_group(0, &bind_group, &[]);
-        cpass.dispatch(256, 1, 1);
+        cpass.dispatch(512, 1, 1);
     }
     if let Some(query_set) = &query_set {
         encoder.write_timestamp(query_set, 1);
     }
-    //encoder.copy_buffer_to_buffer(&input_buf, 0, &output_buf, 0, input.len() as u64);
     if let Some(query_set) = &query_set {
         encoder.resolve_query_set(query_set, 0..2, &query_buf, 0);
     }
@@ -171,7 +166,7 @@ async fn run() {
     if buf_future.await.is_ok() {
         let data_raw = &*buf_slice.get_mapped_range();
         let data: &[u32] = bytemuck::cast_slice(data_raw);
-        println!("failures with strategy {}: {}", data[0], data[1]);
+        println!("failures: {}", data[0]);
     }
     if features.contains(wgpu::Features::TIMESTAMP_QUERY) {
         let ts_period = queue.get_timestamp_period();

compute-shader-hello/src/shader.wgsl

@@ -26,7 +26,6 @@ struct DataBuf {
 
 [[block]]
 struct ControlBuf {
-    strategy: u32;
     failures: atomic<u32>;
 };
 
@@ -46,19 +45,16 @@ fn permute_flag_ix(data_ix: u32) -> u32 {
 [[stage(compute), workgroup_size(256)]]
 fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
     let ix = global_id.x;
-    // Originally this was passed in, but is now hardcoded, as D3DCompiler
-    // thinks control flow becomes nonuniform if it's read from input.
-    let n_iter = 1024u;
-    let strategy = control_buf.strategy;
     var failures = 0u;
-    for (var i: u32 = 0u; i < n_iter; i = i + 1u) {
-        let wr_flag_ix = permute_flag_ix(ix);
-        data_buf.data[ix].data = i + 1u;
+    let a = ix >> 1u;
+    if ((ix & 1u) == 0u) {
+        let wr_flag_ix = permute_flag_ix(a);
+        data_buf.data[a].data = 1u;
         storageBarrier(); // release semantics for writing flag
-        atomicStore(&data_buf.data[wr_flag_ix].flag, i + 1u);
-
+        atomicStore(&data_buf.data[wr_flag_ix].flag, 1u);
+    } else {
         // Read from a different workgroup
-        let read_ix = ((ix & 0xffu) << 8u) | (ix >> 8u);
+        let read_ix = ((a & 0xffu) << 8u) | (a >> 8u);
         let read_flag_ix = permute_flag_ix(read_ix);
 
         let flag = atomicLoad(&data_buf.data[read_flag_ix].flag);