mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-06-03 21:21:23 +00:00
Kindly donated by Brandon Jones and Austin Eng. Tint has been used to convert all the deprecated attributes to the new style. In doing so, comments have been stripped. These are not massively important for the benchmarking. Bindings have also been adjusted to be sequential and unique so that the MSL backend doesn't have to deal with binding remapping. Existing benchmark files that used an underscore '_' have been renamed to use a dash '-' instead, to match the new files. Change-Id: If5fb507b981f107ed570f6eedb55b232448f67aa Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/77443 Reviewed-by: Brandon Jones <bajones@google.com> Reviewed-by: Austin Eng <enga@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com>
206 lines
6.1 KiB
WebGPU Shading Language
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;
|
|
}
|
|
}
|