test_test.wgsl

//2,0

struct RWBuffer {
    MAX_THREADS: i32,
    NUM_TESTING_THREADS: u32,
    counter: atomic<u32>,
    mem_0: array<atomic<i32>, 16>,
};

@group(0)
@binding(0)
var<storage,read_write> rwBuffer: RWBuffer;

@compute
@workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    var gid_x:u32 = global_id.x;
    var pc:u32 = 0u;
    var terminate:u32;

    var total_num_threads:i32 = rwBuffer.MAX_THREADS;
    var num_testing_threads:u32 = rwBuffer.NUM_TESTING_THREADS;
    var index = gid_x / num_testing_threads;

	if(gid_x % num_testing_threads == 0){
        terminate = 0u;
        while (true) {
            if(terminate == 1u) {
                break;
            }
            switch pc {
    			case 0u {
                        if(atomicExchange(&(rwBuffer.mem_0[index]), 1) == 0){
                            pc = 0u;
                        }
                        else {
                            pc = pc + 1u;
                        }
                        break;
                    }
        			case 1u {
                    terminate = 1u;
                    break;
                }
                default {
                    //shouldn't happen 
                }
    		}
		}
	}
	if(gid_x % num_testing_threads == 1){
        terminate = 0u;
        while (true) {
            if(terminate == 1u) {
                break;
            }
            switch pc {
    			case 0u {
                        if(atomicExchange(&(rwBuffer.mem_0[index]), 0) == 0){
                            pc = 0u;
                        }
                        else {
                            pc = pc + 1u;
                        }
                        break;
                    }
        			case 1u {
                    terminate = 1u;
                    break;
                }
                default {
                    //shouldn't happen 
                }
    		}
		}
	}
	atomicAdd(&rwBuffer.counter,1u);
}