Skip to content

Instantly share code, notes, and snippets.

@danteissaias
Created July 4, 2022 12:32
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save danteissaias/ba81b617ef1143a7ceb83e63f116c3c9 to your computer and use it in GitHub Desktop.
Save danteissaias/ba81b617ef1143a7ceb83e63f116c3c9 to your computer and use it in GitHub Desktop.
interface BufferInit {
label?: string;
usage: number;
contents: ArrayBuffer;
}
export function createBufferInit(
device: GPUDevice,
descriptor: BufferInit,
): GPUBuffer {
const contents = new Uint8Array(descriptor.contents);
const unpaddedSize = contents.byteLength;
const padding = 4 - unpaddedSize % 4;
const paddedSize = padding + unpaddedSize;
const buffer = device.createBuffer({
label: descriptor.label,
usage: descriptor.usage,
mappedAtCreation: true,
size: paddedSize,
});
const data = new Uint8Array(buffer.getMappedRange());
data.set(contents);
buffer.unmap();
return buffer;
}
const OVERFLOW = 0xffffffff;
// Get some numbers from the command line, or use the default 1, 4, 3, 295.
let numbers: Uint32Array;
if (Deno.args.length > 0) {
numbers = new Uint32Array(Deno.args.map(parseInt));
} else {
numbers = new Uint32Array([1, 4, 3, 295]);
}
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter?.requestDevice();
if (!device) {
console.error("no suitable adapter found");
Deno.exit(0);
}
const shaderCode = `
@group(0)
@binding(0)
var<storage, read_write> v_indices: array<u32>; // this is used as both input and output for convenience
// The Collatz Conjecture states that for any integer n:
// If n is even, n = n/2
// If n is odd, n = 3n+1
// And repeat this process for each new n, you will always eventually reach 1.
// Though the conjecture has not been proven, no counterexample has ever been found.
// This function returns how many times this recurrence needs to be applied to reach 1.
fn collatz_iterations(n_base: u32) -> u32{
var n: u32 = n_base;
var i: u32 = 0u;
loop {
if (n <= 1u) {
break;
}
if (n % 2u == 0u) {
n = n / 2u;
}
else {
// Overflow? (i.e. 3*n + 1 > 0xffffffffu?)
if (n >= 1431655765u) { // 0x55555555u
return 4294967295u; // 0xffffffffu
}
n = 3u * n + 1u;
}
i = i + 1u;
}
return i;
}
@compute
@workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]);
}
`;
const shaderModule = device.createShaderModule({
code: shaderCode,
});
const stagingBuffer = device.createBuffer({
size: numbers.byteLength,
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
});
const storageBuffer = createBufferInit(device, {
label: "Storage Buffer",
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST |
GPUBufferUsage.COPY_SRC,
contents: numbers.buffer,
});
const computePipeline = device.createComputePipeline({
compute: {
module: shaderModule,
entryPoint: "main",
},
});
const bindGroupLayout = computePipeline.getBindGroupLayout(0);
const bindGroup = device.createBindGroup({
layout: bindGroupLayout,
entries: [
{
binding: 0,
resource: {
buffer: storageBuffer,
},
},
],
});
const encoder = device.createCommandEncoder();
const computePass = encoder.beginComputePass();
computePass.setPipeline(computePipeline);
computePass.setBindGroup(0, bindGroup);
computePass.insertDebugMarker("compute collatz iterations");
computePass.dispatchWorkgroups(numbers.length);
computePass.end();
encoder.copyBufferToBuffer(
storageBuffer,
0,
stagingBuffer,
0,
numbers.byteLength,
);
device.queue.submit([encoder.finish()]);
await stagingBuffer.mapAsync(1);
const arrayBufferData = stagingBuffer.getMappedRange();
const uintData = new Uint32Array(arrayBufferData);
const checkedData = Array.from(uintData).map((n) => {
if (n === OVERFLOW) {
return "OVERFLOW";
} else {
return n.toString();
}
});
console.log(checkedData);
stagingBuffer.unmap();
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment