 | Mixing Data Types in Same Buffer - floats, integers, shorts, ... |  |
When you create buffers or arrays of buffers - there are occasions when you want to mix types. For example, floats, unsigned integers or ascii characters.
Like sweeties! Data very rarely comes along in a single format - it's usually mixed! You never know what you're going to have! So you need to be able to pack different data types together into structures.
First, let's start with the WGSL compute shader - what the data will look like and how it'll be packed together:
<?php
struct MyData {
value_f32: f32,
value_u32: u32,
value_i32: i32
};
@group(0) @binding(0) var<storage, read_write> mixedData: MyData;
@compute @workgroup_size(1)
fn main() {
// Use the uniform values here
let f: f32 = mixedData.value_f32;
let u: u32 = mixedData.value_u32;
let i: i32 = mixedData.value_i32;
// ...
// Modify the values (for testing purposes)
mixedData.value_f32 = f + 1.0;
mixedData.value_u32 = u + 1;
mixedData.value_i32 = i + 1;
}
We've got a structure called MyData which contains a mixture of floating point, signed, unsigned integers. So how do we go about creating this structure on the CPU side in JavaScript?
Let's create a JavaScript object that matches the structure defined in the shader. You also need to create a GPUBuffer to store this data and bind it to the shader.
Very very very important that the structure and alignment of the data in the JavaScript side matches the shader definition (including any hidden alignment padding). Otherwise, it'll be bad! Really really bad!
The solution is to use a ArrayBuffer and DataView so you can manually set each element in the buffer (including the type and any offset).
<?php
const arrayBuffer = new ArrayBuffer(bufferSize);
const dataView = new DataView(arrayBuffer);
dataView.setFloat32(0, data.value_f32, true); // offset 0, little endian
dataView.setUint32(4, data.value_u32, true); // offset 4, little endian
dataView.setInt32(8, data.value_i32, true); // offset 8, little endian
Just for testing - to make sure there isn't any problems - let's take the take over to the compute shader - do a bit of jiggly piggly with it (add some values) - then read it back and check they're correct!
| Complete Code | |
// Initialize WebGPU
if (!navigator.gpu) {
console.error("WebGPU not supported on this browser.");
return;
}
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();
// Define your data - use different distinct numbers
const data = {
value_f32: 10.65,
value_u32: 321,
value_i32: -987,
};
// Create a buffer with the appropriate size
const bufferSize = 12; // 4 bytes for each f32, u32, and i32
const dataBuffer = device.createBuffer({
size: bufferSize,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
});
// Create an array buffer and set the values
const arrayBuffer = new ArrayBuffer(bufferSize);
const dataView = new DataView(arrayBuffer);
dataView.setFloat32(0, data.value_f32, true); // offset 0, little endian
dataView.setUint32(4, data.value_u32, true); // offset 4, little endian
dataView.setInt32(8, data.value_i32, true); // offset 8, little endian
// Upload the data to the GPU
device.queue.writeBuffer(dataBuffer, 0, arrayBuffer);
// Create a bind group layout and bind group
const bindGroupLayout = device.createBindGroupLayout({
entries: [
{ binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
],
});
const bindGroup = device.createBindGroup({
layout: bindGroupLayout,
entries: [
{ binding: 0, resource: { buffer: dataBuffer } },
],
});
// Load and compile your compute shader
const shaderModule = device.createShaderModule({
code: `
struct MyData {
value_f32: f32,
value_u32: u32,
value_i32: i32
};
@group(0) @binding(0) var<storage, read_write> mixedData: MyData;
@compute @workgroup_size(1)
fn main() {
// Use the uniform values here
let f: f32 = mixedData.value_f32;
let u: u32 = mixedData.value_u32;
let i: i32 = mixedData.value_i32;
// ...
// Modify the values (for testing purposes)
mixedData.value_f32 = f + 1.0;
mixedData.value_u32 = u + 1;
mixedData.value_i32 = i + 1;
}
`,
});
// Create the compute pipeline
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout],
}),
compute: {
module: shaderModule,
entryPoint: 'main',
},
});
// Create a command encoder and dispatch the compute shader
const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
await passEncoder.dispatchWorkgroups(1);
await passEncoder.end();
// Submit the commands
await device.queue.submit([commandEncoder.finish()]);
{
// Read back the data
const commandEncoder2 = device.createCommandEncoder();
const readBuffer = device.createBuffer({
size: bufferSize,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
});
commandEncoder2.copyBufferToBuffer(dataBuffer, 0, readBuffer, 0, bufferSize);
await device.queue.submit([commandEncoder2.finish()]);
await readBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = readBuffer.getMappedRange();
const outputDataView = new DataView(copyArrayBuffer);
console.log("value_f32:", outputDataView.getFloat32(0, true));
console.log("value_u32:", outputDataView.getUint32(4, true));
console.log("value_i32:", outputDataView.getInt32(8, true));
readBuffer.unmap();
}
If you run the above code, you'll get this in the output:
<?php
log:["value_f32:",11.649999618530273]
log:["value_u32:",322]
log:["value_i32:",-986]
| Resources and Links | |
• WebGPU Lab Demo [LINK]
|
