lvngd/mortonCodes.html

## mortonCodes.html
<!DOCTYPE html>
<html lang="en">
	<body>
		<canvas></canvas>
	<script>

	async function main(device) {

	    /* WEBGPU CONTEXT SETUP */
	    const canvas = document.querySelector('canvas');
	    const context = canvas.getContext('webgpu');

	    // number of coordinates sets / number of Morton Codes
	    const coordinateCount = 16;

	    //using a vec3f which requires padding, so 3 vertices + 1 = 4
	    const numCoordinateVertices = 3; // x,y,z
	    const numVerticesWithPadding = numCoordinateVertices + 1;
	    const totalVertexCount = coordinateCount * numCoordinateVertices;

	    /* GENERATE COORDINATES */
	    var testPoints = [];
	    for (let i = 0; i < numVerticesWithPadding * coordinateCount; i++) {
	        testPoints.push(Math.random());
	    }

	    const mortonCodeInput = new Float32Array(testPoints);
	    // output buffer size -> 1 morton code for every 4 numbers(x,y,z + padding)
	    const mortonOutputBufferSize = Math.floor(mortonCodeInput.byteLength / 4);

	    // buffer for coordinate points
	    let mortonCodeInputBuffer = device.createBuffer({
	        label: 'morton code points input buffer - A',
	        size: mortonCodeInput.byteLength,
	        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
	    });
	    let mortonCodeOutputBuffer = device.createBuffer({
	        label: 'morton code buffer',
	        size: mortonOutputBufferSize,
	        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
	    });

	    // write points data to buffer
	    device.queue.writeBuffer(mortonCodeInputBuffer, 0, mortonCodeInput);

	    /* SHADER MODULE */
	    const mortonCodeModule = device.createShaderModule({
	        label: 'compute morton codes from input points',
	        code: `
	            const numCoordinates = ${coordinateCount};

	            fn expandBits(v: u32) -> u32
	                {
	                    var vv = (v * 0x00010001) & 0xFF0000FF;
	                    vv = (vv * 0x00000101) & 0x0F00F00F;
	                    vv = (vv * 0x00000011) & 0xC30C30C3;
	                    vv = (vv * 0x00000005) & 0x49249249;
	                    return vv;
	                }

	            fn morton3D(x: f32, y: f32, z: f32) -> u32
	                {
	                    var xx = min(max(x * 1024.0, 0.0), 1023.0);
	                    var yy = min(max(y * 1024.0, 0.0), 1023.0);
	                    var zz = min(max(z * 1024.0, 0.0), 1023.0);

	                    var xxE = expandBits(u32(xx));
	                    var yyE = expandBits(u32(yy));
	                    var zzE = expandBits(u32(zz));
	                    return xxE * 4 + yyE * 2 + zzE;
	                }

	            @group(0) @binding(0) var<storage, read> data: array<vec3f, numCoordinates>;
	            @group(0) @binding(1) var<storage, read_write> dataOut: array<u32>;

	            // for now using a thread for each coordinate set
	            // this will be updated later because there is a max of 256
	            @compute @workgroup_size(numCoordinates, 1, 1) fn computeMortonCodes(
	                @builtin(global_invocation_id) id: vec3<u32>
	            ) {
	                let i = id.x;
	                let mc = data[i];
	                dataOut[i] = u32(morton3D(mc.x, mc.y, mc.z));
	            }
					`,
	    });

	    /* PIPELINE */
	    const mortonCodePipelineLabel = 'morton code pipeline';

	    const mortonCodePipeline = device.createComputePipeline({
	        label: mortonCodePipelineLabel,
	        layout: 'auto',
	        compute: {
	            module: mortonCodeModule,
	            entryPoint: 'computeMortonCodes',
	        },
	    });

	    /* BIND GROUP */
	    const mortonBindGroupLabel = 'bind group for morton code compute';

	    const mortonBindGroupEntries = [{
	            binding: 0,
	            resource: {
	                buffer: mortonCodeInputBuffer
	            }
	        },
	        {
	            binding: 1,
	            resource: {
	                buffer: mortonCodeOutputBuffer
	            }
	        }
	    ];

	    const mortonCodeBindGroup = device.createBindGroup({
	        label: mortonBindGroupLabel,
	        layout: mortonCodePipeline.getBindGroupLayout(0),
	        entries: mortonBindGroupEntries
	    });

	    async function completeComputePipeline(device) {

	    	/* CREATE COMMAND BUFFER */
	        const encoder = device.createCommandEncoder({
	            label: 'morton encoder'
	        });

	        const pass = encoder.beginComputePass({
	            label: 'morton compute pass'
	        });

	        /* SET PIPELINE AND BIND GROUP */
	        pass.setPipeline(mortonCodePipeline);
	        pass.setBindGroup(0, mortonCodeBindGroup);
	        pass.dispatchWorkgroups(1);

	        /* END COMPUTE PASS */
	        pass.end();

	        // create buffer for output
	        let mortonCodeBufferResult = device.createBuffer({
	            label: 'morton codes output result',
	            size: mortonCodeOutputBuffer.size,
	            usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
	        });

	        /* COPY RESULTS TO OUTPUT BUFFER */
	        encoder.copyBufferToBuffer(mortonCodeOutputBuffer, 0, mortonCodeBufferResult, 0, mortonCodeBufferResult.size);


	        /* FINISH + SUBMIT COMMAND BUFFER */
	        let commandBuffer = encoder.finish();
	        device.queue.submit([commandBuffer]);

	        // print to console
	        await mortonCodeBufferResult.mapAsync(GPUMapMode.READ);
	        const mortonCodeResult = new Int32Array(mortonCodeBufferResult.getMappedRange());
	        console.log('morton codes: ', mortonCodeResult);
	        return;
	    }

	    completeComputePipeline(device);


	} //end async main
	</script>

	<script>
	async function start() {
	    if (!navigator.gpu) {
	        fail('need a browser that supports WebGPU');
	        return;
	    }
	    const adapter = await navigator.gpu.requestAdapter();

	    if (!adapter) {
	        fail('this browser supports WebGPU but it appears disabled');
	        return;
	    }
	    const device = await adapter?.requestDevice();

	    device.lost.then((info) => {
	        console.error(`WebGPU device was lost: ${info.message}`);
	        if (info.reason !== 'destroyed') {
	            // try again
	            start();
	        }
	    });

	    function fail(msg) {
	        alert(msg);
	    }

	    main(device);
	} //end start function

	start();
	</script>

	</body>
</html>
	<!DOCTYPE html>
	<html lang="en">
	<body>
	<canvas></canvas>
	<script>

	async function main(device) {

	/* WEBGPU CONTEXT SETUP */
	const canvas = document.querySelector('canvas');
	const context = canvas.getContext('webgpu');

	// number of coordinates sets / number of Morton Codes
	const coordinateCount = 16;

	//using a vec3f which requires padding, so 3 vertices + 1 = 4
	const numCoordinateVertices = 3; // x,y,z
	const numVerticesWithPadding = numCoordinateVertices + 1;
	const totalVertexCount = coordinateCount * numCoordinateVertices;

	/* GENERATE COORDINATES */
	var testPoints = [];
	for (let i = 0; i < numVerticesWithPadding * coordinateCount; i++) {
	testPoints.push(Math.random());
	}

	const mortonCodeInput = new Float32Array(testPoints);
	// output buffer size -> 1 morton code for every 4 numbers(x,y,z + padding)
	const mortonOutputBufferSize = Math.floor(mortonCodeInput.byteLength / 4);

	// buffer for coordinate points
	let mortonCodeInputBuffer = device.createBuffer({
	label: 'morton code points input buffer - A',
	size: mortonCodeInput.byteLength,
	usage: GPUBufferUsage.STORAGE \| GPUBufferUsage.COPY_SRC \| GPUBufferUsage.COPY_DST,
	});
	let mortonCodeOutputBuffer = device.createBuffer({
	label: 'morton code buffer',
	size: mortonOutputBufferSize,
	usage: GPUBufferUsage.STORAGE \| GPUBufferUsage.COPY_SRC \| GPUBufferUsage.COPY_DST,
	});

	// write points data to buffer
	device.queue.writeBuffer(mortonCodeInputBuffer, 0, mortonCodeInput);

	/* SHADER MODULE */
	const mortonCodeModule = device.createShaderModule({
	label: 'compute morton codes from input points',
	code: `
	const numCoordinates = ${coordinateCount};

	fn expandBits(v: u32) -> u32
	{
	var vv = (v * 0x00010001) & 0xFF0000FF;
	vv = (vv * 0x00000101) & 0x0F00F00F;
	vv = (vv * 0x00000011) & 0xC30C30C3;
	vv = (vv * 0x00000005) & 0x49249249;
	return vv;
	}

	fn morton3D(x: f32, y: f32, z: f32) -> u32
	{
	var xx = min(max(x * 1024.0, 0.0), 1023.0);
	var yy = min(max(y * 1024.0, 0.0), 1023.0);
	var zz = min(max(z * 1024.0, 0.0), 1023.0);

	var xxE = expandBits(u32(xx));
	var yyE = expandBits(u32(yy));
	var zzE = expandBits(u32(zz));
	return xxE * 4 + yyE * 2 + zzE;
	}

	@group(0) @binding(0) var<storage, read> data: array<vec3f, numCoordinates>;
	@group(0) @binding(1) var<storage, read_write> dataOut: array<u32>;

	// for now using a thread for each coordinate set
	// this will be updated later because there is a max of 256
	@compute @workgroup_size(numCoordinates, 1, 1) fn computeMortonCodes(
	@builtin(global_invocation_id) id: vec3<u32>
	) {
	let i = id.x;
	let mc = data[i];
	dataOut[i] = u32(morton3D(mc.x, mc.y, mc.z));
	}
	`,
	});

	/* PIPELINE */
	const mortonCodePipelineLabel = 'morton code pipeline';

	const mortonCodePipeline = device.createComputePipeline({
	label: mortonCodePipelineLabel,
	layout: 'auto',
	compute: {
	module: mortonCodeModule,
	entryPoint: 'computeMortonCodes',
	},
	});

	/* BIND GROUP */
	const mortonBindGroupLabel = 'bind group for morton code compute';

	const mortonBindGroupEntries = [{
	binding: 0,
	resource: {
	buffer: mortonCodeInputBuffer
	}
	},
	{
	binding: 1,
	resource: {
	buffer: mortonCodeOutputBuffer
	}
	}
	];

	const mortonCodeBindGroup = device.createBindGroup({
	label: mortonBindGroupLabel,
	layout: mortonCodePipeline.getBindGroupLayout(0),
	entries: mortonBindGroupEntries
	});

	async function completeComputePipeline(device) {

	/* CREATE COMMAND BUFFER */
	const encoder = device.createCommandEncoder({
	label: 'morton encoder'
	});

	const pass = encoder.beginComputePass({
	label: 'morton compute pass'
	});

	/* SET PIPELINE AND BIND GROUP */
	pass.setPipeline(mortonCodePipeline);
	pass.setBindGroup(0, mortonCodeBindGroup);
	pass.dispatchWorkgroups(1);

	/* END COMPUTE PASS */
	pass.end();

	// create buffer for output
	let mortonCodeBufferResult = device.createBuffer({
	label: 'morton codes output result',
	size: mortonCodeOutputBuffer.size,
	usage: GPUBufferUsage.MAP_READ \| GPUBufferUsage.COPY_DST
	});

	/* COPY RESULTS TO OUTPUT BUFFER */
	encoder.copyBufferToBuffer(mortonCodeOutputBuffer, 0, mortonCodeBufferResult, 0, mortonCodeBufferResult.size);


	/* FINISH + SUBMIT COMMAND BUFFER */
	let commandBuffer = encoder.finish();
	device.queue.submit([commandBuffer]);

	// print to console
	await mortonCodeBufferResult.mapAsync(GPUMapMode.READ);
	const mortonCodeResult = new Int32Array(mortonCodeBufferResult.getMappedRange());
	console.log('morton codes: ', mortonCodeResult);
	return;
	}

	completeComputePipeline(device);


	} //end async main
	</script>

	<script>
	async function start() {
	if (!navigator.gpu) {
	fail('need a browser that supports WebGPU');
	return;
	}
	const adapter = await navigator.gpu.requestAdapter();

	if (!adapter) {
	fail('this browser supports WebGPU but it appears disabled');
	return;
	}
	const device = await adapter?.requestDevice();

	device.lost.then((info) => {
	console.error(`WebGPU device was lost: ${info.message}`);
	if (info.reason !== 'destroyed') {
	// try again
	start();
	}
	});

	function fail(msg) {
	alert(msg);
	}

	main(device);
	} //end start function

	start();
	</script>

	</body>
	</html>