forked from 1b-agent/opencadstudio-upstream
perf(render): keep solid meshes resident
Keep all solid geometry GPU-resident during navigation. Use vertex-rate instancing, compact untextured vertices, vertex-cache ordering, and cached silhouette sources without culling geometry.
This commit is contained in:
Hakan Seven committed
1 parent
15ab027d70
commit
5e8221f8f1
9 files changed
+756
-1012
No files matched your search
+57
-7
@@ -98,9 +98,6 @@ struct MeshInstance {
|
||||
normal_row_2: vec4<f32>,
|
||||
};
|
||||
|
||||
@group(1) @binding(15)
|
||||
var<storage, read> mesh_instances: array<MeshInstance>;
|
||||
|
||||
struct VertexIn {
|
||||
@location(0) position: vec3<f32>,
|
||||
@location(1) normal: vec3<f32>,
|
||||
@@ -111,6 +108,22 @@ struct VertexIn {
|
||||
@location(12) uv_refraction_normal: vec4<f32>,
|
||||
};
|
||||
|
||||
struct PlainVertexIn {
|
||||
@location(0) position: vec3<f32>,
|
||||
@location(1) normal: vec3<f32>,
|
||||
@location(3) position_low: vec3<f32>,
|
||||
};
|
||||
|
||||
struct InstanceIn {
|
||||
@location(4) model_row_0: vec4<f32>,
|
||||
@location(5) model_row_1: vec4<f32>,
|
||||
@location(7) model_row_2: vec4<f32>,
|
||||
@location(8) translation_low: vec4<f32>,
|
||||
@location(9) normal_row_0: vec4<f32>,
|
||||
@location(13) normal_row_1: vec4<f32>,
|
||||
@location(14) normal_row_2: vec4<f32>,
|
||||
};
|
||||
|
||||
struct VertexOut {
|
||||
@builtin(position) clip_pos: vec4<f32>,
|
||||
@location(0) normal: vec3<f32>,
|
||||
@@ -153,13 +166,25 @@ fn relative_position(
|
||||
return (world_high - u.eye_high) + (world_low - u.eye_low);
|
||||
}
|
||||
|
||||
fn mesh_instance(input: InstanceIn) -> MeshInstance {
|
||||
return MeshInstance(
|
||||
input.model_row_0,
|
||||
input.model_row_1,
|
||||
input.model_row_2,
|
||||
input.translation_low,
|
||||
input.normal_row_0,
|
||||
input.normal_row_1,
|
||||
input.normal_row_2,
|
||||
);
|
||||
}
|
||||
|
||||
@vertex
|
||||
fn vs_main(
|
||||
v: VertexIn,
|
||||
@builtin(instance_index) instance_index: u32,
|
||||
i: InstanceIn,
|
||||
) -> VertexOut {
|
||||
var out: VertexOut;
|
||||
let instance = mesh_instances[instance_index];
|
||||
let instance = mesh_instance(i);
|
||||
let rel = relative_position(v.position, v.position_low, instance);
|
||||
out.clip_pos = u.view_rot * vec4<f32>(rel, 1.0);
|
||||
out.normal = normalize(vec3<f32>(
|
||||
@@ -178,13 +203,38 @@ fn vs_main(
|
||||
return out;
|
||||
}
|
||||
|
||||
@vertex
|
||||
fn vs_main_plain(
|
||||
v: PlainVertexIn,
|
||||
i: InstanceIn,
|
||||
) -> VertexOut {
|
||||
var out: VertexOut;
|
||||
let instance = mesh_instance(i);
|
||||
let rel = relative_position(v.position, v.position_low, instance);
|
||||
out.clip_pos = u.view_rot * vec4<f32>(rel, 1.0);
|
||||
out.normal = normalize(vec3<f32>(
|
||||
dot(instance.normal_row_0.xyz, v.normal),
|
||||
dot(instance.normal_row_1.xyz, v.normal),
|
||||
dot(instance.normal_row_2.xyz, v.normal),
|
||||
));
|
||||
out.world_pos = rel;
|
||||
out.uv_diffuse = vec2<f32>(0.0);
|
||||
out.uv_specular = vec2<f32>(0.0);
|
||||
out.uv_reflection = vec2<f32>(0.0);
|
||||
out.uv_opacity = vec2<f32>(0.0);
|
||||
out.uv_bump = vec2<f32>(0.0);
|
||||
out.uv_refraction = vec2<f32>(0.0);
|
||||
out.uv_normal = vec2<f32>(0.0);
|
||||
return out;
|
||||
}
|
||||
|
||||
@vertex
|
||||
fn vs_edge(
|
||||
v: EdgeVertexIn,
|
||||
@builtin(instance_index) instance_index: u32,
|
||||
i: InstanceIn,
|
||||
) -> EdgeVertexOut {
|
||||
var out: EdgeVertexOut;
|
||||
let instance = mesh_instances[instance_index];
|
||||
let instance = mesh_instance(i);
|
||||
let rel = relative_position(v.position, v.position_low, instance);
|
||||
out.clip_pos = u.view_rot * vec4<f32>(rel, 1.0);
|
||||
out.color = v.color;
|
||||
|
||||
@@ -1,78 +0,0 @@
|
||||
struct CullUniform {
|
||||
view_rot: mat4x4<f32>,
|
||||
eye: vec4<f32>,
|
||||
count: vec4<u32>,
|
||||
};
|
||||
|
||||
struct CullItem {
|
||||
min: vec4<f32>,
|
||||
max: vec4<f32>,
|
||||
counts: vec4<u32>,
|
||||
info: vec4<u32>,
|
||||
};
|
||||
|
||||
struct DrawIndexed {
|
||||
index_count: u32,
|
||||
instance_count: u32,
|
||||
first_index: u32,
|
||||
base_vertex: u32,
|
||||
first_instance: u32,
|
||||
};
|
||||
|
||||
struct Draw {
|
||||
vertex_count: u32,
|
||||
instance_count: u32,
|
||||
first_vertex: u32,
|
||||
first_instance: u32,
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<uniform> u: CullUniform;
|
||||
@group(0) @binding(1) var<storage, read> items: array<CullItem>;
|
||||
@group(0) @binding(2) var<storage, read_write> opaque: array<DrawIndexed>;
|
||||
@group(0) @binding(3) var<storage, read_write> transparent: array<DrawIndexed>;
|
||||
@group(0) @binding(4) var<storage, read_write> wire: array<DrawIndexed>;
|
||||
@group(0) @binding(5) var<storage, read_write> edge: array<Draw>;
|
||||
|
||||
fn chunk_visible(item: CullItem) -> bool {
|
||||
if (item.info.y == 0u) {
|
||||
return false;
|
||||
}
|
||||
var min_ndc = vec2<f32>(1e30, 1e30);
|
||||
var max_ndc = vec2<f32>(-1e30, -1e30);
|
||||
for (var corner = 0u; corner < 8u; corner++) {
|
||||
let point = vec3<f32>(
|
||||
select(item.min.x, item.max.x, (corner & 1u) != 0u),
|
||||
select(item.min.y, item.max.y, (corner & 2u) != 0u),
|
||||
select(item.min.z, item.max.z, (corner & 4u) != 0u),
|
||||
);
|
||||
let clip = u.view_rot * vec4<f32>(point - u.eye.xyz, 1.0);
|
||||
// A box touching or crossing the eye plane is conservatively retained.
|
||||
if (clip.w <= 1e-6) {
|
||||
return true;
|
||||
}
|
||||
let ndc = clip.xy / clip.w;
|
||||
min_ndc = min(min_ndc, ndc);
|
||||
max_ndc = max(max_ndc, ndc);
|
||||
}
|
||||
// 25% viewport margin on each side is 0.5 in NDC.
|
||||
return !(
|
||||
max_ndc.x < -1.5
|
||||
|| min_ndc.x > 1.5
|
||||
|| max_ndc.y < -1.5
|
||||
|| min_ndc.y > 1.5
|
||||
);
|
||||
}
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
|
||||
let index = id.x;
|
||||
if (index >= u.count.x) {
|
||||
return;
|
||||
}
|
||||
let item = items[index];
|
||||
let instances = select(0u, item.info.x, chunk_visible(item));
|
||||
opaque[index] = DrawIndexed(item.counts.x, instances, 0u, 0u, 0u);
|
||||
transparent[index] = DrawIndexed(item.counts.y, instances, 0u, 0u, 0u);
|
||||
wire[index] = DrawIndexed(item.counts.z, instances, 0u, 0u, 0u);
|
||||
edge[index] = Draw(item.counts.w, instances, 0u, 0u);
|
||||
}
|
||||
@@ -51,19 +51,24 @@ struct MeshInstance {
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<uniform> u: Uniforms;
|
||||
@group(1) @binding(15) var<storage, read> mesh_instances: array<MeshInstance>;
|
||||
|
||||
struct VertexIn {
|
||||
@location(0) position: vec3<f32>,
|
||||
@location(3) position_low: vec3<f32>,
|
||||
};
|
||||
|
||||
struct InstanceIn {
|
||||
@location(4) model_row_0: vec4<f32>,
|
||||
@location(5) model_row_1: vec4<f32>,
|
||||
@location(7) model_row_2: vec4<f32>,
|
||||
@location(8) translation_low: vec4<f32>,
|
||||
};
|
||||
|
||||
@vertex
|
||||
fn vs_main(
|
||||
vertex: VertexIn,
|
||||
@builtin(instance_index) instance_index: u32,
|
||||
instance: InstanceIn,
|
||||
) -> @builtin(position) vec4<f32> {
|
||||
let instance = mesh_instances[instance_index];
|
||||
let world_high = vec3<f32>(
|
||||
dot(instance.model_row_0.xyz, vertex.position) + instance.model_row_0.w,
|
||||
dot(instance.model_row_1.xyz, vertex.position) + instance.model_row_1.w,
|
||||
|
||||
Reference in new issue
Block a user