Skip to content
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
75 changes: 36 additions & 39 deletions sample/particles/main.ts
Original file line number Diff line number Diff line change
Expand Up @@ -182,39 +182,30 @@ quadVertexBuffer.unmap();
//////////////////////////////////////////////////////////////////////////////
// Texture
//////////////////////////////////////////////////////////////////////////////
let texture: GPUTexture;
let textureWidth = 1;
let textureHeight = 1;
let numMipLevels = 1;
{
const response = await fetch('../../assets/img/webgpu.png');
const imageBitmap = await createImageBitmap(await response.blob());

// Calculate number of mip levels required to generate the probability map
while (
textureWidth < imageBitmap.width ||
textureHeight < imageBitmap.height
) {
textureWidth *= 2;
textureHeight *= 2;
numMipLevels++;
}
texture = device.createTexture({
size: [imageBitmap.width, imageBitmap.height, 1],
mipLevelCount: numMipLevels,
format: 'rgba8unorm',
usage:
GPUTextureUsage.TEXTURE_BINDING |
GPUTextureUsage.STORAGE_BINDING |
GPUTextureUsage.COPY_DST |
GPUTextureUsage.RENDER_ATTACHMENT,
});
device.queue.copyExternalImageToTexture(
{ source: imageBitmap },
{ texture: texture },
[imageBitmap.width, imageBitmap.height]
);
}
const isPowerOf2 = (v: number) => Math.log2(v) % 1 === 0;
const response = await fetch('../../assets/img/webgpu.png');
const imageBitmap = await createImageBitmap(await response.blob());
assert(imageBitmap.width === imageBitmap.height, 'image must be square');
assert(isPowerOf2(imageBitmap.width), 'image must be a power of 2');

// Calculate number of mip levels required to generate the probability map
const mipLevelCount =
(Math.log2(Math.max(imageBitmap.width, imageBitmap.height)) + 1) | 0;
const texture = device.createTexture({
size: [imageBitmap.width, imageBitmap.height, 1],
mipLevelCount,
format: 'rgba8unorm',
usage:
GPUTextureUsage.TEXTURE_BINDING |
GPUTextureUsage.STORAGE_BINDING |
GPUTextureUsage.COPY_DST |
GPUTextureUsage.RENDER_ATTACHMENT,
});
device.queue.copyExternalImageToTexture(
{ source: imageBitmap },
{ texture: texture },
[imageBitmap.width, imageBitmap.height]
);

//////////////////////////////////////////////////////////////////////////////
// Probability map generation
Expand Down Expand Up @@ -247,22 +238,22 @@ let numMipLevels = 1;
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
});
const buffer_a = device.createBuffer({
size: textureWidth * textureHeight * 4,
size: texture.width * texture.height * 4,
usage: GPUBufferUsage.STORAGE,
});
const buffer_b = device.createBuffer({
size: textureWidth * textureHeight * 4,
size: buffer_a.size,
usage: GPUBufferUsage.STORAGE,
});
device.queue.writeBuffer(
probabilityMapUBOBuffer,
0,
new Int32Array([textureWidth])
new Uint32Array([texture.width])
);
const commandEncoder = device.createCommandEncoder();
for (let level = 0; level < numMipLevels; level++) {
const levelWidth = textureWidth >> level;
const levelHeight = textureHeight >> level;
for (let level = 0; level < texture.mipLevelCount; level++) {
const levelWidth = Math.max(1, texture.width >> level);
const levelHeight = Math.max(1, texture.height >> level);
const pipeline =
level == 0
? probabilityMapImportLevelPipeline.getBindGroupLayout(0)
Expand Down Expand Up @@ -472,3 +463,9 @@ function frame() {
}
configureContext();
requestAnimationFrame(frame);

function assert(cond: boolean, msg = '') {
if (!cond) {
throw new Error(msg);
}
}
29 changes: 13 additions & 16 deletions sample/particles/probabilityMap.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -2,17 +2,12 @@ struct UBO {
width : u32,
}

struct Buffer {
weights : array<f32>,
}

@binding(0) @group(0) var<uniform> ubo : UBO;
@binding(1) @group(0) var<storage, read> buf_in : Buffer;
@binding(2) @group(0) var<storage, read_write> buf_out : Buffer;
@binding(1) @group(0) var<storage, read> buf_in : array<f32>;
@binding(2) @group(0) var<storage, read_write> buf_out : array<f32>;
@binding(3) @group(0) var tex_in : texture_2d<f32>;
@binding(3) @group(0) var tex_out : texture_storage_2d<rgba8unorm, write>;


////////////////////////////////////////////////////////////////////////////////
// import_level
//
Expand All @@ -21,9 +16,11 @@ struct Buffer {
////////////////////////////////////////////////////////////////////////////////
@compute @workgroup_size(64)
fn import_level(@builtin(global_invocation_id) coord : vec3u) {
_ = &buf_in;
let offset = coord.x + coord.y * ubo.width;
buf_out.weights[offset] = textureLoad(tex_in, vec2i(coord.xy), 0).w;
_ = &buf_in; // so the bindGroups are similar.
if (all(coord.xy < vec2u(textureDimensions(tex_in)))) {
let offset = coord.x + coord.y * ubo.width;
buf_out[offset] = textureLoad(tex_in, vec2i(coord.xy), 0).w;
}
}

////////////////////////////////////////////////////////////////////////////////
Expand All @@ -40,13 +37,13 @@ fn export_level(@builtin(global_invocation_id) coord : vec3u) {
let dst_offset = coord.x + coord.y * ubo.width;
let src_offset = coord.x*2u + coord.y*2u * ubo.width;

let a = buf_in.weights[src_offset + 0u];
let b = buf_in.weights[src_offset + 1u];
let c = buf_in.weights[src_offset + 0u + ubo.width];
let d = buf_in.weights[src_offset + 1u + ubo.width];
let sum = dot(vec4f(a, b, c, d), vec4f(1.0));
let a = buf_in[src_offset + 0u];
let b = buf_in[src_offset + 1u];
let c = buf_in[src_offset + 0u + ubo.width];
let d = buf_in[src_offset + 1u + ubo.width];
let sum = a + b + c + d;

buf_out.weights[dst_offset] = sum / 4.0;
buf_out[dst_offset] = sum / 4.0;

let probabilities = vec4f(a, a+b, a+b+c, sum) / max(sum, 0.0001);
textureStore(tex_out, vec2i(coord.xy), probabilities);
Expand Down
Loading