dawn-cmake/test/tint/benchmark/metaball-isosurface.wgsl

206 lines
6.1 KiB
WebGPU Shading Language

struct Tables {
edges : array<u32, 256>;
tris : array<i32, 4096>;
}
@group(0) @binding(0) var<storage> tables : Tables;
struct IsosurfaceVolume {
min : vec3<f32>;
max : vec3<f32>;
step : vec3<f32>;
size : vec3<u32>;
threshold : f32;
values : array<f32>;
}
@group(0) @binding(1) var<storage, write> volume : IsosurfaceVolume;
struct PositionBuffer {
values : array<f32>;
}
@group(0) @binding(2) var<storage, write> positionsOut : PositionBuffer;
struct NormalBuffer {
values : array<f32>;
}
@group(0) @binding(3) var<storage, write> normalsOut : NormalBuffer;
struct IndexBuffer {
tris : array<u32>;
}
@group(0) @binding(4) var<storage, write> indicesOut : IndexBuffer;
struct DrawIndirectArgs {
vc : u32;
vertexCount : atomic<u32>;
firstVertex : u32;
firstInstance : u32;
indexCount : atomic<u32>;
indexedInstanceCount : u32;
indexedFirstIndex : u32;
indexedBaseVertex : u32;
indexedFirstInstance : u32;
}
@group(0) @binding(5) var<storage, read_write> drawOut : DrawIndirectArgs;
fn valueAt(index : vec3<u32>) -> f32 {
if (any((index >= volume.size))) {
return 0.0;
}
let valueIndex = ((index.x + (index.y * volume.size.x)) + ((index.z * volume.size.x) * volume.size.y));
return volume.values[valueIndex];
}
fn positionAt(index : vec3<u32>) -> vec3<f32> {
return (volume.min + (volume.step * vec3<f32>(index.xyz)));
}
fn normalAt(index : vec3<u32>) -> vec3<f32> {
return vec3<f32>((valueAt((index - vec3<u32>(1u, 0u, 0u))) - valueAt((index + vec3<u32>(1u, 0u, 0u)))), (valueAt((index - vec3<u32>(0u, 1u, 0u))) - valueAt((index + vec3<u32>(0u, 1u, 0u)))), (valueAt((index - vec3<u32>(0u, 0u, 1u))) - valueAt((index + vec3<u32>(0u, 0u, 1u)))));
}
var<private> positions : array<vec3<f32>, 12>;
var<private> normals : array<vec3<f32>, 12>;
var<private> indices : array<u32, 12>;
var<private> cubeVerts : u32 = 0u;
fn interpX(index : u32, i : vec3<u32>, va : f32, vb : f32) {
let mu = ((volume.threshold - va) / (vb - va));
positions[cubeVerts] = (positionAt(i) + vec3<f32>((volume.step.x * mu), 0.0, 0.0));
let na = normalAt(i);
let nb = normalAt((i + vec3<u32>(1u, 0u, 0u)));
normals[cubeVerts] = mix(na, nb, vec3<f32>(mu, mu, mu));
indices[index] = cubeVerts;
cubeVerts = (cubeVerts + 1u);
}
fn interpY(index : u32, i : vec3<u32>, va : f32, vb : f32) {
let mu = ((volume.threshold - va) / (vb - va));
positions[cubeVerts] = (positionAt(i) + vec3<f32>(0.0, (volume.step.y * mu), 0.0));
let na = normalAt(i);
let nb = normalAt((i + vec3<u32>(0u, 1u, 0u)));
normals[cubeVerts] = mix(na, nb, vec3<f32>(mu, mu, mu));
indices[index] = cubeVerts;
cubeVerts = (cubeVerts + 1u);
}
fn interpZ(index : u32, i : vec3<u32>, va : f32, vb : f32) {
let mu = ((volume.threshold - va) / (vb - va));
positions[cubeVerts] = (positionAt(i) + vec3<f32>(0.0, 0.0, (volume.step.z * mu)));
let na = normalAt(i);
let nb = normalAt((i + vec3<u32>(0u, 0u, 1u)));
normals[cubeVerts] = mix(na, nb, vec3<f32>(mu, mu, mu));
indices[index] = cubeVerts;
cubeVerts = (cubeVerts + 1u);
}
@stage(compute) @workgroup_size(4, 4, 4)
fn computeMain(@builtin(global_invocation_id) global_id : vec3<u32>) {
let i0 = global_id;
let i1 = (global_id + vec3<u32>(1u, 0u, 0u));
let i2 = (global_id + vec3<u32>(1u, 1u, 0u));
let i3 = (global_id + vec3<u32>(0u, 1u, 0u));
let i4 = (global_id + vec3<u32>(0u, 0u, 1u));
let i5 = (global_id + vec3<u32>(1u, 0u, 1u));
let i6 = (global_id + vec3<u32>(1u, 1u, 1u));
let i7 = (global_id + vec3<u32>(0u, 1u, 1u));
let v0 = valueAt(i0);
let v1 = valueAt(i1);
let v2 = valueAt(i2);
let v3 = valueAt(i3);
let v4 = valueAt(i4);
let v5 = valueAt(i5);
let v6 = valueAt(i6);
let v7 = valueAt(i7);
var cubeIndex = 0u;
if ((v0 < volume.threshold)) {
cubeIndex = (cubeIndex | 1u);
}
if ((v1 < volume.threshold)) {
cubeIndex = (cubeIndex | 2u);
}
if ((v2 < volume.threshold)) {
cubeIndex = (cubeIndex | 4u);
}
if ((v3 < volume.threshold)) {
cubeIndex = (cubeIndex | 8u);
}
if ((v4 < volume.threshold)) {
cubeIndex = (cubeIndex | 16u);
}
if ((v5 < volume.threshold)) {
cubeIndex = (cubeIndex | 32u);
}
if ((v6 < volume.threshold)) {
cubeIndex = (cubeIndex | 64u);
}
if ((v7 < volume.threshold)) {
cubeIndex = (cubeIndex | 128u);
}
let edges = tables.edges[cubeIndex];
if (((edges & 1u) != 0u)) {
interpX(0u, i0, v0, v1);
}
if (((edges & 2u) != 0u)) {
interpY(1u, i1, v1, v2);
}
if (((edges & 4u) != 0u)) {
interpX(2u, i3, v3, v2);
}
if (((edges & 8u) != 0u)) {
interpY(3u, i0, v0, v3);
}
if (((edges & 16u) != 0u)) {
interpX(4u, i4, v4, v5);
}
if (((edges & 32u) != 0u)) {
interpY(5u, i5, v5, v6);
}
if (((edges & 64u) != 0u)) {
interpX(6u, i7, v7, v6);
}
if (((edges & 128u) != 0u)) {
interpY(7u, i4, v4, v7);
}
if (((edges & 256u) != 0u)) {
interpZ(8u, i0, v0, v4);
}
if (((edges & 512u) != 0u)) {
interpZ(9u, i1, v1, v5);
}
if (((edges & 1024u) != 0u)) {
interpZ(10u, i2, v2, v6);
}
if (((edges & 2048u) != 0u)) {
interpZ(11u, i3, v3, v7);
}
let triTableOffset = ((cubeIndex << 4u) + 1u);
let indexCount = u32(tables.tris[(triTableOffset - 1u)]);
var firstVertex = atomicAdd(&(drawOut.vertexCount), cubeVerts);
let bufferOffset = ((global_id.x + (global_id.y * volume.size.x)) + ((global_id.z * volume.size.x) * volume.size.y));
let firstIndex = (bufferOffset * 15u);
for(var i = 0u; (i < cubeVerts); i = (i + 1u)) {
positionsOut.values[((firstVertex * 3u) + (i * 3u))] = positions[i].x;
positionsOut.values[(((firstVertex * 3u) + (i * 3u)) + 1u)] = positions[i].y;
positionsOut.values[(((firstVertex * 3u) + (i * 3u)) + 2u)] = positions[i].z;
normalsOut.values[((firstVertex * 3u) + (i * 3u))] = normals[i].x;
normalsOut.values[(((firstVertex * 3u) + (i * 3u)) + 1u)] = normals[i].y;
normalsOut.values[(((firstVertex * 3u) + (i * 3u)) + 2u)] = normals[i].z;
}
for(var i = 0u; (i < indexCount); i = (i + 1u)) {
let index = tables.tris[(triTableOffset + i)];
indicesOut.tris[(firstIndex + i)] = (firstVertex + indices[index]);
}
for(var i = indexCount; (i < 15u); i = (i + 1u)) {
indicesOut.tris[(firstIndex + i)] = firstVertex;
}
}