IT Crowd TV series - famous and funny quote when Jan says not to talk about memory or ram - and Moss laughs and jokes that ram is memory. Jan is head of IT - but has no idea how computers work.
Memory Array - Alignment Buffers and Pains
Going to talk about memory and how it isn't always what you think it is! Especially for long sequential blocks of data!
For example the below implementation copies from
array0
to
array1
! What could be simplier? Both are just an array of floats. However, in WGSL the
array1
is stored as
array< vec3<f32> >
- which is common as you might want to group every 3 values as positions or vectors (x,y,z) - but what you might not be aware of is how it's stored in memory on the GPU!!
As the
vec3<f32>
is a structure - so all structures are aligned to 16 byte boundaries!
All because of the internal alignment/padding - which can be a real pain if you're not careful.
Below gives the complete implementation - with a link to the webgpulab at the bottom so you can test out the code.
How to Fix?
You can pad the data - by adding an extra 0 every 4 floating point values in the array. Then use two
vec3
structures in the WGSL code.
However, a more compact and simplier way is just to make sure the structure in WGSL are both
f32
- that way the input/output is the same - inside the WGSL code, just convert the array values to a
vec3
structure manually.
// Convert the flat array of f32 into vec3s
let v0 = vec3<f32>(array0[ global_id.x + 0 ],
array0[ global_id.x + 1 ],
array0[ global_id.x + 2 ] );
let v1 = vec3<f32>(array0[ global_id.x + 3 ],
array0[ global_id.x + 4 ],
array0[ global_id.x + 5 ] );
// Then use them as normal
// vec1
array1[ global_id.x+0 ] = v0.x;
array1[ global_id.x+1 ] = v0.y;
array1[ global_id.x+2 ] = v0.z;
// vec2
array1[ global_id.x+3 ] = v1.x;
array1[ global_id.x+4 ] = v1.y;
array1[ global_id.x+5 ] = v1.z;
The above gives the following output:
<?php
log:["res:",[1,2,3,4,5,6]
Implementation
The full code that can be run from a `index.js` file (including all the initialization and copying to/from the GPU and the compute shader).
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();
const array0 = new Float32Array( [1, 2, 3, 4, 5, 6] );
const array1 = new Float32Array( array0.byteLength/4 );
var array0Buffer = device.createBuffer({ size: array0.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC } );
var array1Buffer = device.createBuffer({ size: array1.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC } );
device.queue.writeBuffer(array0Buffer, 0, array0);
const shaderCode = `
@group(0) @binding(0) var<storage, read> array0 : array< vec3<f32> >;
@group(0) @binding(1) var<storage, read_write> array1 : array< f32 >;
@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
// vec1
array1[ global_id.x+0 ] = array0[ global_id.x ].x;
array1[ global_id.x+1 ] = array0[ global_id.x ].y;
array1[ global_id.x+2 ] = array0[ global_id.x ].z;
// vec2
array1[ global_id.x+3 ] = array0[ global_id.x+1 ].x;
array1[ global_id.x+4 ] = array0[ global_id.x+1 ].y;
array1[ global_id.x+5 ] = array0[ global_id.x+1 ].z;
}
`;
const pipeline = device.createComputePipeline({
layout: 'auto',
compute: {
module: device.createShaderModule({
code: shaderCode
}),
entryPoint: 'main'
}
});
const bindGroup = device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [
{ binding: 0, resource: { buffer: array0Buffer } },
{ binding: 1, resource: { buffer: array1Buffer } },
]
});
const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(pipeline);
passEncoder.setBindGroup(0, bindGroup);
passEncoder.dispatchWorkgroups( 1 , 1, 1);
await passEncoder.end();
device.queue.submit([commandEncoder.finish()]);
await device.queue.onSubmittedWorkDone();
// -------------------------------
// All the compute is done - just a matter of copying the data back from the array for analysis.
// -------------------------------
// Write a small helper function
async function getGPUBuffer( buf, siz, msg )
{
// Note this buffer is not linked to the 'STORAGE' compute (used to bring the data back to the CPU)
const gbufferTmp = device.createBuffer({ size: siz, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ});
const commandEncoder = device.createCommandEncoder();
// Encode commands for copying buffer to buffer.
commandEncoder.copyBufferToBuffer(
buf, // source buffer
0, // source offset
gbufferTmp, // destination buffer
0, // destination offset
siz // size
);
// Submit GPU commands.
const gpuCommands = commandEncoder.finish();
await device.queue.submit([gpuCommands]);
// Read buffer.
await gbufferTmp.mapAsync(GPUMapMode.READ);
const arrayBuffer = gbufferTmp.getMappedRange();
const arr = Array.from( new Float32Array(arrayBuffer) );
gbufferTmp.unmap();
//log(msg + 'array contents:', arr);
return arr;
}
// Copy array1 back and print it to the output
let res = await getGPUBuffer( array1Buffer, array1.byteLength, 'array1 ' );
console.log('res:', res );
Alignment with `align` Expression in Structures
There is the
align
expresssion - which you can add into structures to force alignment. However, be aware, the alignment but be a power of 2.
We can do a simple example - create a structure and have an input data - copy 2 array items - the first array item is fine (
data1
), it's only when the second array item is copied you notice the problem.
We use an input array with a sequence of numbers from 1 to 12, and the output is given below:
<?php
log:["res:",[1,2,3,5,6,7,9,10,11,0,0,0]]
As you can see - the array isn't a sequential count, 4 is missing, 8 is missing etc. As the data is aligned - so our sequential data in memory doens't align with the WGSL structures.
<?php
struct AlignedData {
// This field is aligned to 8 bytes
@align(8)
data1: vec3<f32>,
// This field is aligned to 8 bytes
@align(8)
data2: vec2<f32>,
// This field has the default alignment
data3: f32,
};
@group(0) @binding(0) var<storage, read> array0 : array< AlignedData >;
@group(0) @binding(1) var<storage, read_write> array1 : array< f32 >;
@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
var c = global_id.x;
// copy 2 array items
for (var i=0; i<2; i++)
{
array1[ c ] = array0[i].data1.x; c++;
array1[ c ] = array0[i].data1.y; c++;
array1[ c ] = array0[i].data1.z; c++;
array1[ c ] = array0[i].data2.x; c++;
array1[ c ] = array0[i].data2.y; c++;
array1[ c ] = array0[i].data3; c++;
}
}// end main
Larger offsets
For example if we increase the size of the array test data 1..32 - then make the alignments larger, you'll get:
<?php
struct AlignedData {
// This field is aligned to x bytes
@align(16)
data1: vec3<f32>,
// This field is aligned to x bytes
@align(32)
data2: vec2<f32>,
// This field has the default alignment
data3: f32,
};
You can't just force alignment to 1 byte boundaries - as the offset for each field variable needs to be a power of 2! Hence, you'll always get gaps - but you have a bit more control over the gaps with the align expression.
