raphlinus/webgpu-test.html

## webgpu-test.html
<html>
    <body>
        <p>WebGPU test:</p>

    </body><div id="out"></div>
    <div>Expected: [1, 1, 1]</div>
    <script>
function display(text) {
    document.getElementById("out").textContent = text;
}
async function runTest() {
    if (!navigator.gpu) {
        display("WebGPU doesn't seem to be enabled.");
        return;

    }
    const adapter = await navigator.gpu.requestAdapter();
    if (!adapter) {
        display('error initializing adapter');
        return;
    }
    const device = await adapter.requestDevice();
    const buf0 = device.createBuffer({
        mappedAtCreation: false,
        size: 1024,
        usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE,
    });
    const buf1 = device.createBuffer({
        mappedAtCreation: false,
        size: 1024,
        usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE,
    });
    const bindGroupLayout = device.createBindGroupLayout({
        entries: [
        {
                binding: 0,
                visibility: GPUShaderStage.COMPUTE,
                buffer: { type: "storage" },
            },
            {
                binding: 1,
                visibility: GPUShaderStage.COMPUTE,
                buffer: { type: "storage" },
            }
        ]
    });
    const bindGroup = device.createBindGroup({
        layout: bindGroupLayout,
        entries: [
            {
                binding: 0,
                resource: { buffer: buf0 }
            },
            {
                binding: 1,
                resource: { buffer: buf1 }
            },
        ]
    });
    const shaderModule = device.createShaderModule({
        code: `
@group(0) @binding(0)
var<storage, read_write> bump: atomic<u32>;

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

const WG_SIZE = 256u;

var<workgroup> sh_tile_offset: u32;

@compute @workgroup_size(256)
fn main(
    @builtin(global_invocation_id) global_id: vec3<u32>,
    @builtin(local_invocation_id) local_id: vec3<u32>,
) {

    let drawobj_ix = global_id.x;
    let tile_count_in = local_id.x + 1u;
    if local_id.x == WG_SIZE - 1u {
        sh_tile_offset = 1u + atomicAdd(&bump, tile_count_in);
    }
    workgroupBarrier();
    let tile_offset = sh_tile_offset;
    if drawobj_ix < 3u {
        paths[drawobj_ix] = tile_offset;
    }

}
        `
    });
    const computePipeline = device.createComputePipeline({
        layout: device.createPipelineLayout({
            bindGroupLayouts: [bindGroupLayout]
        }),
        compute: {
            module: shaderModule,
            entryPoint: "main"
        }
    });
    const commandEncoder = device.createCommandEncoder();
    const passEncoder = commandEncoder.beginComputePass();
    passEncoder.setPipeline(computePipeline);
    passEncoder.setBindGroup(0, bindGroup);
    passEncoder.dispatchWorkgroups(1);
    passEncoder.end();
    const buf2 = device.createBuffer({
        mappedAtCreation: false,
        size: 1024,
        usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
    });
    commandEncoder.copyBufferToBuffer(buf1, 0, buf2, 0, 1024);
    const gpuCommands = commandEncoder.finish();
    device.queue.submit([gpuCommands]);
    await buf2.mapAsync(GPUMapMode.READ);
    const arrayBuffer = buf2.getMappedRange();
    let a = new Uint32Array(arrayBuffer);
    display(`Actual: [${a[0]}, ${a[1]}, ${a[2]}]`);
}
runTest();
    </script>

</html>
	<html>
	<body>
	<p>WebGPU test:</p>

	</body><div id="out"></div>
	<div>Expected: [1, 1, 1]</div>
	<script>
	function display(text) {
	document.getElementById("out").textContent = text;
	}
	async function runTest() {
	if (!navigator.gpu) {
	display("WebGPU doesn't seem to be enabled.");
	return;

	}
	const adapter = await navigator.gpu.requestAdapter();
	if (!adapter) {
	display('error initializing adapter');
	return;
	}
	const device = await adapter.requestDevice();
	const buf0 = device.createBuffer({
	mappedAtCreation: false,
	size: 1024,
	usage: GPUBufferUsage.COPY_SRC \| GPUBufferUsage.STORAGE,
	});
	const buf1 = device.createBuffer({
	mappedAtCreation: false,
	size: 1024,
	usage: GPUBufferUsage.COPY_SRC \| GPUBufferUsage.STORAGE,
	});
	const bindGroupLayout = device.createBindGroupLayout({
	entries: [
	{
	binding: 0,
	visibility: GPUShaderStage.COMPUTE,
	buffer: { type: "storage" },
	},
	{
	binding: 1,
	visibility: GPUShaderStage.COMPUTE,
	buffer: { type: "storage" },
	}
	]
	});
	const bindGroup = device.createBindGroup({
	layout: bindGroupLayout,
	entries: [
	{
	binding: 0,
	resource: { buffer: buf0 }
	},
	{
	binding: 1,
	resource: { buffer: buf1 }
	},
	]
	});
	const shaderModule = device.createShaderModule({
	code: `
	@group(0) @binding(0)
	var<storage, read_write> bump: atomic<u32>;

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

	const WG_SIZE = 256u;

	var<workgroup> sh_tile_offset: u32;

	@compute @workgroup_size(256)
	fn main(
	@builtin(global_invocation_id) global_id: vec3<u32>,
	@builtin(local_invocation_id) local_id: vec3<u32>,
	) {

	let drawobj_ix = global_id.x;
	let tile_count_in = local_id.x + 1u;
	if local_id.x == WG_SIZE - 1u {
	sh_tile_offset = 1u + atomicAdd(&bump, tile_count_in);
	}
	workgroupBarrier();
	let tile_offset = sh_tile_offset;
	if drawobj_ix < 3u {
	paths[drawobj_ix] = tile_offset;
	}

	}
	`
	});
	const computePipeline = device.createComputePipeline({
	layout: device.createPipelineLayout({
	bindGroupLayouts: [bindGroupLayout]
	}),
	compute: {
	module: shaderModule,
	entryPoint: "main"
	}
	});
	const commandEncoder = device.createCommandEncoder();
	const passEncoder = commandEncoder.beginComputePass();
	passEncoder.setPipeline(computePipeline);
	passEncoder.setBindGroup(0, bindGroup);
	passEncoder.dispatchWorkgroups(1);
	passEncoder.end();
	const buf2 = device.createBuffer({
	mappedAtCreation: false,
	size: 1024,
	usage: GPUBufferUsage.COPY_DST \| GPUBufferUsage.MAP_READ,
	});
	commandEncoder.copyBufferToBuffer(buf1, 0, buf2, 0, 1024);
	const gpuCommands = commandEncoder.finish();
	device.queue.submit([gpuCommands]);
	await buf2.mapAsync(GPUMapMode.READ);
	const arrayBuffer = buf2.getMappedRange();
	let a = new Uint32Array(arrayBuffer);
	display(`Actual: [${a[0]}, ${a[1]}, ${a[2]}]`);
	}
	runTest();
	</script>

	</html>