diff --git a/bin/data/config.json b/bin/data/config.json index aa261711..e93dd336 100644 --- a/bin/data/config.json +++ b/bin/data/config.json @@ -31,7 +31,7 @@ "vxgi": { // "limiter": 0, "limiter": 0.0125, - "size": 192, + "size": 256, "dispatch": 16, "cascades": 3, "cascadePower": 2, @@ -42,8 +42,8 @@ "filtering": "LINEAR", "shadows": 0, "extents": { - "min": [ -8, -8, -8 ], - "max": [ 8, 8, 8 ] + "min": [ -16, -4, -16 ], + "max": [ 16, 4, 16 ] } }, "rt": { @@ -60,7 +60,7 @@ }, "graph": { "initial buffer elements": 128, - "global storage": false + "global storage": true }, "ext": { "vulkan": { diff --git a/bin/data/entities/burger.json b/bin/data/entities/burger.json index f7b52d49..8f1c1ae8 100644 --- a/bin/data/entities/burger.json +++ b/bin/data/entities/burger.json @@ -31,9 +31,7 @@ "gravity": [ 0, -9.81, 0 ], "mass": 10, - "type": "sphere", - "radius": 1, - "static": false, + "type": "bounding box", "recenter": false }, "graph": { diff --git a/bin/data/entities/physics_prop.json b/bin/data/entities/physics_prop.json new file mode 100644 index 00000000..061b363c --- /dev/null +++ b/bin/data/entities/physics_prop.json @@ -0,0 +1,17 @@ +{ + "assets": [], + "behaviors": [ + "SoundEmitterBehavior" + ], + "metadata": { + "holdable": true, + "physics": { + "gravity": [ 0, 0, 0 ], + "inertia": false, + + "mass": 1, + "type": "aabb", + "recenter": false + } + } +} \ No newline at end of file diff --git a/bin/data/entities/prop.json b/bin/data/entities/prop.json index 130962ac..0fb2fcdc 100644 --- a/bin/data/entities/prop.json +++ b/bin/data/entities/prop.json @@ -8,8 +8,8 @@ "physics": { "mass": 0, "inertia": false, - "type": "bounding box" - // "type": "mesh" + // "type": "bounding box" + "type": "mesh" } } } \ No newline at end of file diff --git a/bin/data/entities/scripts/player.lua b/bin/data/entities/scripts/player.lua index 4f30a3f8..51cda1f5 100644 --- a/bin/data/entities/scripts/player.lua +++ b/bin/data/entities/scripts/player.lua @@ -244,6 +244,8 @@ ent:addHook( "entity:Use.%UID%", function( payload ) heldObject.distance = offset:norm() prop:getComponent("PhysicsBody"):enableGravity(false) + + print( prop:name() ) else validUse = not string.matched( prop:name(), "/^worldspawn/" ) end diff --git a/bin/data/scenes/sourceengine/mds_mcdonalds.json b/bin/data/scenes/sourceengine/mds_mcdonalds.json index 4b3d2cfb..af7f4cfd 100644 --- a/bin/data/scenes/sourceengine/mds_mcdonalds.json +++ b/bin/data/scenes/sourceengine/mds_mcdonalds.json @@ -3,7 +3,7 @@ "assets": [ // { "filename": "./models/mds_mcdonalds.glb" } { "filename": "./models/mds_mcdonalds/graph.json" }, - { "filename": "/burger.json", "delay": 4 } + { "filename": "/burger.json", "delay": 1 } ], "metadata": { "graph": { @@ -15,8 +15,9 @@ "func_door_rotating_5568": { "action": "load", "payload": { "import": "/door.json", "metadata": { "angle":-1.570795, "normal": [1,0,0] } } }, "func_door_rotating_5584": { "action": "load", "payload": { "import": "/door.json", "metadata": { "angle":-1.570795, "normal": [1,0,0] } } }, + "prop_physics_override_5813": { "action": "load", "payload": { "import": "/physics_prop.json" } }, + // regex matches - // "/^prop_physics_override/": { "action": "load", "payload": { "import": "/prop.json", "metadata": { "physics": { "gravity": [ 0, 0, 0 ] } } } }, // "/^prop_physics_[^o]/": { "action": "load", "payload": { "import": "/prop.json" } }, "/^tools\\/toolsnodraw/": { "material": { diff --git a/bin/data/scenes/sourceengine/scene.json b/bin/data/scenes/sourceengine/scene.json index 1e49a4ec..a27dd52d 100644 --- a/bin/data/scenes/sourceengine/scene.json +++ b/bin/data/scenes/sourceengine/scene.json @@ -6,7 +6,7 @@ ], "metadata": { "light": { - "fog-0": { + "fog": { // "color": [ 0.1, 0.1, 0.1 ], // "color": [ 0.2, 0.2, 0.2 ], "color": [ 0.3, 0.3, 0.3 ], diff --git a/bin/data/shaders/common/fog.h b/bin/data/shaders/common/fog.h index 7c62eff2..84c76dfc 100644 --- a/bin/data/shaders/common/fog.h +++ b/bin/data/shaders/common/fog.h @@ -2,47 +2,48 @@ void fog( in Ray ray, inout vec3 i, float scale ) { if ( ubo.settings.fog.stepScale <= 0 || ubo.settings.fog.range.x == 0 || ubo.settings.fog.range.y == 0 ) return; #if FOG_RAY_MARCH + if ( ubo.settings.fog.stepScale <= 0 || ubo.settings.fog.range.y <= 0 ) return; + const float range = ubo.settings.fog.range.y; - const vec3 boundsMin = vec3(-range,-range,-range) + ray.origin; - const vec3 boundsMax = vec3(range,range,range) + ray.origin; - const int numSteps = int(length(boundsMax - boundsMin) * ubo.settings.fog.stepScale ); + const vec3 boundsMin = ray.origin - vec3(range); + const vec3 boundsMax = ray.origin + vec3(range); const vec2 rayBoxInfo = rayBoxDst( boundsMin, boundsMax, ray ); const float dstToBox = rayBoxInfo.x; const float dstInsideBox = rayBoxInfo.y; const float depth = surface.position.eye.z; - const float aperture = 0; // PI * 0.5; - const float coneCoefficient = 2.0 * tan(aperture * 0.5); + if ( dstInsideBox < 0 || dstToBox > depth ) return; - // march - if ( 0 <= dstInsideBox && dstToBox <= depth ) { - float stepSize = dstInsideBox / numSteps; - float dstLimit = min( depth - dstToBox, dstInsideBox ); - float totalDensity = 0; - float transmittance = 1; - float lightFactor = scale; - float coneDiameter = coneCoefficient * ray.distance; - float level = aperture > 0 ? log2( coneDiameter ) : 0; - float density = 0; - vec3 uvw; - ray.distance = dstToBox; - while ( ray.distance < dstLimit ) { - ray.distance += stepSize; - ray.position = ray.origin + ray.direction * ray.distance; - coneDiameter = coneCoefficient * ray.distance; - level = aperture > 0 ? log2( coneDiameter ) : 0; - uvw = ray.position * ubo.settings.fog.densityScale * 0.001 + ubo.settings.fog.offset * 0.01; - density = max(0, textureLod(samplerNoise, uvw, level).r - ubo.settings.fog.densityThreshold) * ubo.settings.fog.densityMultiplier; - if ( density > 0 ) { - density = exp(-density * stepSize * ubo.settings.fog.absorbtion); - transmittance *= density; - lightFactor *= density; - if ( transmittance < 0.1 ) break; - } + const float marchLimit = min( depth - dstToBox, dstInsideBox ); + const int MAX_STEPS = 32; + + const float stepSize = marchLimit / float(MAX_STEPS); + const vec3 stepVec = ray.direction * stepSize; + const float jitter = rand2(gl_GlobalInvocationID.xy); + + const float densityScale = ubo.settings.fog.densityScale * 0.001; + const vec3 densityOffset = ubo.settings.fog.offset * 0.01; + const float densityThreshold = ubo.settings.fog.densityThreshold; + const float densityMult = ubo.settings.fog.densityMultiplier; + const float absorption = ubo.settings.fog.absorbtion; + + vec3 currentPos = ray.origin + (ray.direction * (dstToBox + stepSize * jitter)); + float transmittance = 1.0; + + for ( uint j = 0; j < MAX_STEPS; ++j ) { + vec3 uvw = currentPos * densityScale + densityOffset; + float noiseVal = textureLod(samplerNoise, uvw, 0.0).r; + float density = max(0.0, noiseVal - densityThreshold) * densityMult; + if ( density > 0.0 ) { + float beer = exp(-density * stepSize * absorption); + transmittance *= beer; + if ( transmittance < 0.01 ) break; } - i.rgb = mix(ubo.settings.fog.color.rgb, i.rgb, transmittance ); + currentPos += stepVec; } + + i.rgb = mix(ubo.settings.fog.color.rgb, i.rgb, transmittance ); #endif #if FOG_BASIC const vec3 color = ubo.settings.fog.color.rgb; diff --git a/bin/data/shaders/common/pbr.h b/bin/data/shaders/common/pbr.h index 377d7897..47e02cdf 100644 --- a/bin/data/shaders/common/pbr.h +++ b/bin/data/shaders/common/pbr.h @@ -2,7 +2,7 @@ void pbr() { // per-surface, not per-light, compute once - // Freslen reflectance for a dieletric + // Fresnel reflectance for a dieletric const vec3 F0 = mix(vec3(0.04), surface.material.albedo.rgb, surface.material.metallic); // outcoming light from surface to eye const vec3 Lo = normalize( -surface.position.eye ); diff --git a/bin/data/shaders/common/vxgi.h b/bin/data/shaders/common/vxgi.h index c2c2e01a..09e6c7a3 100644 --- a/bin/data/shaders/common/vxgi.h +++ b/bin/data/shaders/common/vxgi.h @@ -1,35 +1,51 @@ // GI -uint cascadeIndex( vec3 v ) { - float x = max3( abs( v ) ); - for ( uint cascade = 0; cascade < CASCADES; ++cascade ) - if ( x / cascadePower(cascade) < 1 - voxelInfo.radianceSizeRecip ) return cascade; - return CASCADES - 1; +float cascadeScales[CASCADES]; +float cascadeScalesInv[CASCADES]; +void precomputeCascades() { + for ( int i = 0; i < CASCADES; ++i ) { + cascadeScales[i] = cascadePower(i); + cascadeScalesInv[i] = 1.0 / cascadeScales[i]; + } } - vec4 voxelTrace( inout Ray ray, float aperture, float maxDistance ) { - ray.origin += ray.direction * voxelInfo.radianceSizeRecip * 2 * SQRT2; + ray.origin += ray.direction * voxelInfo.radianceSizeRecip * 1.5; + #if VXGI_NDC ray.origin = vec3( ubo.settings.vxgi.matrix * vec4( ray.origin, 1.0 ) ); ray.direction = vec3( ubo.settings.vxgi.matrix * vec4( ray.direction, 0.0 ) ); - uint cascade = cascadeIndex(ray.origin); + float absMax = max3(abs(ray.origin)); #else - uint cascade = cascadeIndex( vec3( ubo.settings.vxgi.matrix * vec4( ray.origin, 1.0 ) ) ); + const float voxelOrigin = vec3( ubo.settings.vxgi.matrix * vec4( ray.origin, 1.0 ) ); + float absMax = max3(abs(voxelOrigin)); #endif - const float granularityRecip = ubo.settings.vxgi.granularity; //2.0; // 0.25f * (CASCADES - cascade); - const float granularity = 1.0f / granularityRecip; - const float occlusionFalloff = ubo.settings.vxgi.occlusionFalloff; //128.0f; - const vec3 voxelBounds = voxelInfo.max - voxelInfo.min; - const vec3 voxelBoundsRecip = 1.0f / voxelBounds; + + uint cascade = 0; + for ( uint c = 0; c < CASCADES - 1; ++c ) { + if ( absMax * cascadeScalesInv[c] < (1.0 - voxelInfo.radianceSizeRecip)) { + cascade = c; + break; + } + cascade = CASCADES - 1; + } + + const float maxCascadeScale = cascadeScales[CASCADES-1]; + float currentCascadeScale = cascadeScales[cascade]; + float currentCascadeScaleInv = cascadeScalesInv[cascade]; + + const float granularity = ubo.settings.vxgi.granularity; + const float occlusionFalloff = ubo.settings.vxgi.occlusionFalloff; const float coneCoefficient = 2.0 * tan(aperture * 0.5); - const uint maxSteps = uint(voxelInfo.radianceSize * cascadePower(CASCADES-1) * granularityRecip); + + const uint maxSteps = uint(voxelInfo.radianceSize * maxCascadeScale * granularity); + const float maxRadiance = 0.90; // box - const vec2 rayBoxInfoA = rayBoxDst( voxelInfo.min * cascadePower(cascade), voxelInfo.max * cascadePower(cascade), ray ); - const vec2 rayBoxInfoB = rayBoxDst( voxelInfo.min * cascadePower(CASCADES-1), voxelInfo.max * cascadePower(CASCADES-1), ray ); + const vec2 rayBoxInfoA = rayBoxDst( voxelInfo.min * currentCascadeScale, voxelInfo.max * currentCascadeScale, ray ); + const vec2 rayBoxInfoB = rayBoxDst( voxelInfo.min * maxCascadeScale, voxelInfo.max * maxCascadeScale, ray ); const float tStart = rayBoxInfoA.x; const float tEnd = maxDistance > 0 ? min(maxDistance, rayBoxInfoB.y) : rayBoxInfoB.y; - const float tDelta = voxelInfo.radianceSizeRecip * granularityRecip; - const uint MIN_VOXEL_MIP_LEVEL = 1; + const float tDelta = voxelInfo.radianceSizeRecip * granularity; + // marcher ray.distance = tStart + tDelta * ubo.settings.vxgi.traceStartOffsetFactor; ray.position = vec3(0); @@ -39,32 +55,39 @@ vec4 voxelTrace( inout Ray ray, float aperture, float maxDistance ) { float coneDiameter = coneCoefficient * ray.distance; float level = aperture > 0 ? log2( coneDiameter ) : 0; vec4 color = vec4(0); - float occlusion = 0.0; + float occlusion = 0; uint stepCounter = 0; - while ( color.a < 1.0 && occlusion < 1.0 && ray.distance < tEnd && stepCounter++ < maxSteps ) { - float stepScale = max(1.0, coneDiameter * cascadePower(cascade)); + + ray.distance += tDelta * rand2(gl_GlobalInvocationID.xy); + + while ( color.a < maxRadiance && occlusion < 1.0 && ray.distance < tEnd && stepCounter++ < maxSteps ) { + float stepScale = max(1.0, (coneDiameter * currentCascadeScale) * 1.5); ray.distance += tDelta * stepScale; - // ray.distance += tDelta * cascadePower(cascade) * max(1, coneDiameter); ray.position = ray.origin + ray.direction * ray.distance; - #if VXGI_NDC - uvw = ray.position; - #else - uvw = vec3( ubo.settings.vxgi.matrix * vec4( ray.position, 1.0 ) ); - #endif - cascade = cascadeIndex( uvw ); - uvw = (uvw / cascadePower(cascade)) * 0.5 + 0.5; - if ( cascade >= CASCADES || uvw.x < 0.0 || uvw.y < 0.0 || uvw.z < 0.0 || uvw.x >= 1.0 || uvw.y >= 1.0 || uvw.z >= 1.0 ) break; + absMax = max3(abs(ray.position)); + if ( absMax * currentCascadeScaleInv > 0.99 ) { + if ( ++cascade >= CASCADES ) break; + + currentCascadeScale = cascadeScales[cascade]; + currentCascadeScaleInv = cascadeScalesInv[cascade]; + + ray.distance += tDelta * currentCascadeScale; + continue; + } + + uvw = ray.position * currentCascadeScaleInv * 0.5 + 0.5; coneDiameter = coneCoefficient * ray.distance; level = aperture > 0 ? log2( coneDiameter ) : 0; - // level = clamp(level, MIN_VOXEL_MIP_LEVEL, voxelInfo.mipmapLevels - 1); - if ( level >= voxelInfo.mipmapLevels ) break; + radiance = textureLod(voxelOutput[nonuniformEXT(cascade)], uvw.xzy, level); - color += (1.0 - color.a) * radiance; + + color.rgb += (1.0 - color.a) * radiance.rgb * radiance.a; + color.a += (1.0 - color.a) * radiance.a; + occlusion += ((1.0f - occlusion) * radiance.a) / (1.0f + occlusionFalloff * coneDiameter); } return maxDistance > 0 ? color : vec4(color.rgb, occlusion); -// return vec4(color.rgb, occlusion); } vec4 voxelConeTrace( inout Ray ray, float aperture ) { return voxelTrace( ray, aperture, 0 ); @@ -86,9 +109,12 @@ float shadowFactorVXGI( const Light light, float def ) { return 1.0 - voxelTrace( ray, SHADOW_APERTURE, z ).a; } void indirectLightingVXGI() { + precomputeCascades(); + voxelInfo.radianceSize = textureSize( voxelOutput[0], 0 ).x; voxelInfo.radianceSizeRecip = 1.0 / voxelInfo.radianceSize; voxelInfo.mipmapLevels = log2(voxelInfo.radianceSize) + 1; + #if VXGI_NDC voxelInfo.min = vec3( -1 ); voxelInfo.max = vec3( 1 ); @@ -98,53 +124,36 @@ void indirectLightingVXGI() { voxelInfo.max = vec3( inverseOrtho * vec4( 1, 1, 1, 1 ) ); #endif - const vec3 P = surface.position.world; - const vec3 N = surface.normal.world; - -#if 1 - const vec3 right = normalize(orthogonal(N)); - const vec3 up = normalize(cross(right, N)); - - const uint CONES_COUNT = 6; - const vec3 CONES[] = { - N, - normalize(N + 0.0f * right + 0.866025f * up), - normalize(N + 0.823639f * right + 0.267617f * up), - normalize(N + 0.509037f * right + -0.7006629f * up), - normalize(N + -0.50937f * right + -0.7006629f * up), - normalize(N + -0.823639f * right + 0.267617f * up), - }; -#else - const vec3 ortho = normalize(orthogonal(N)); - const vec3 ortho2 = normalize(cross(ortho, N)); - - const vec3 corner = 0.5f * (ortho + ortho2); - const vec3 corner2 = 0.5f * (ortho - ortho2); - - const uint CONES_COUNT = 9; - const vec3 CONES[] = { - N, - normalize(mix(N, ortho, 0.5)), - normalize(mix(N, -ortho, 0.5)), - normalize(mix(N, ortho2, 0.5)), - normalize(mix(N, -ortho2, 0.5)), - normalize(mix(N, corner, 0.5)), - normalize(mix(N, -corner, 0.5)), - normalize(mix(N, corner2, 0.5)), - normalize(mix(N, -corner2, 0.5)), - }; -#endif - - const float DIFFUSE_CONE_APERTURE = 2.0f * 0.57735f; - const float DIFFUSE_INDIRECT_FACTOR = 1.0f / float(CONES_COUNT) * 0.125f; - - const float SPECULAR_CONE_APERTURE = clamp(tan(PI * 0.5f * surface.material.roughness), 0.0174533f, PI); // tan( R * PI * 0.5f * 0.1f ); - const float SPECULAR_INDIRECT_FACTOR = (1.0f - surface.material.metallic) * (1.0f - surface.material.roughness); // * 0.25; // 1.0f; - vec4 indirectDiffuse = vec4(0); vec4 indirectSpecular = vec4(0); -// outFragColor.rgb = voxelConeTrace( surface.ray, 0 ).rgb; return; + const vec3 P = surface.position.world; + const vec3 N = surface.normal.world; + + const vec3 right = normalize(orthogonal(N)); + const vec3 up = normalize(cross(right, N)); + +#if 0 + { + Ray ray; + ray.direction = N; + ray.origin = P + N * (voxelInfo.radianceSizeRecip * 4.0); + + indirectDiffuse = voxelConeTrace(ray, 1.0f); + surface.material.occlusion += 1.0 - clamp(indirectDiffuse.a, 0.0, 1.0); + } +#else + const uint CONES_COUNT = 4; + const vec3 CONES[] = { + normalize(N + right + up), + normalize(N - right + up), + normalize(N + right - up), + normalize(N - right - up) + }; + + const float DIFFUSE_CONE_APERTURE = 1.1547; + const float DIFFUSE_INDIRECT_FACTOR = 1.0f / float(CONES_COUNT) * 0.125f; + if ( DIFFUSE_INDIRECT_FACTOR > 0.0f ) { float weight = PI * 0.25f; for ( uint i = 0; i < CONES_COUNT; ++i ) { @@ -154,32 +163,35 @@ void indirectLightingVXGI() { indirectDiffuse += voxelConeTrace( ray, DIFFUSE_CONE_APERTURE ) * weight; weight = PI * 0.15f; } - // indirectDiffuse.rgb *= surface.material.albedo.rgb; surface.material.occlusion += 1.0 - clamp(indirectDiffuse.a, 0.0, 1.0); - // outFragColor.rgb = indirectDiffuse.rgb; return; - // outFragColor.rgb = vec3(surface.material.occlusion); return; } + indirectDiffuse *= DIFFUSE_INDIRECT_FACTOR; +#endif + + const float SPECULAR_CONE_APERTURE = clamp(tan(PI * 0.5f * surface.material.roughness), 0.0174533f, PI); // tan( R * PI * 0.5f * 0.1f ); + const float SPECULAR_INDIRECT_FACTOR = (1.0f - surface.material.metallic) * (1.0f - surface.material.roughness); // * 0.25; // 1.0f; if ( SPECULAR_INDIRECT_FACTOR > 0.0f ) { const vec3 R = reflect( normalize(P - surface.ray.origin), N ); Ray ray; ray.direction = R; ray.origin = P; // + ray.direction; indirectSpecular = voxelConeTrace( ray, SPECULAR_CONE_APERTURE ); - // indirectSpecular.rgb *= surface.material.albedo.rgb; - // outFragColor.rgb = indirectSpecular.rgb; return; } - -/* - if ( true ) { - gammaCorrect(indirectDiffuse.rgb, 1.0 / ubo.settings.bloom.gamma); - } -*/ - indirectDiffuse *= DIFFUSE_INDIRECT_FACTOR; indirectSpecular *= SPECULAR_INDIRECT_FACTOR; + // Calculate Fresnel + { + const vec3 V = normalize(surface.ray.origin - surface.position.world); + const vec3 N = surface.normal.world; + const float NdotV = max(dot(N, V), 0.0); + + const vec3 F0 = mix(vec3(0.04), surface.material.albedo.rgb, surface.material.metallic); + const vec3 F = fresnelSchlick(F0, NdotV); + + indirectDiffuse.rgb *= (1.0 - F) * (1.0 - surface.material.metallic) * surface.material.occlusion; + } surface.material.indirect += indirectDiffuse + indirectSpecular; -// outFragColor.rgb = surface.material.indirect.rgb; return; // deferred sampling doesn't have a blended albedo buffer // in place we'll just cone trace behind the window diff --git a/bin/data/shaders/display/bloom/comp.glsl b/bin/data/shaders/display/bloom/comp.glsl index d5b709bc..058764f3 100644 --- a/bin/data/shaders/display/bloom/comp.glsl +++ b/bin/data/shaders/display/bloom/comp.glsl @@ -17,12 +17,14 @@ layout (binding = 0) uniform UBO { float threshold; float smoothness; uint size; - uint padding1; + float padding1; + + float weights[32]; } ubo; -layout (binding = 1, rgba16f) uniform volatile coherent image2D imageColor; -layout (binding = 2, rgba16f) uniform volatile coherent image2D imageBloom; -layout (binding = 3, rgba16f) uniform volatile coherent image2D imagePingPong; +layout (binding = 1, rgba16f) uniform image2D imageColor; +layout (binding = 2, rgba16f) uniform image2D imageBloom; +layout (binding = 3, rgba16f) uniform image2D imagePingPong; #include "../../common/macros.h" #include "../../common/structs.h" @@ -34,35 +36,30 @@ void main() { const ivec2 size = imageSize( imageColor ); if ( texel.x >= size.x || texel.y >= size.y ) return; - uint kernelSize = ubo.size * 2 + 1; - float sigma = kernelSize / (6 * max(0.001f, ubo.smoothness)); - float strength = 1.0f; - float scale = 1.0f; - if ( mode == 0 ) { // fill bloom vec3 result = imageLoad( imageColor, texel ).rgb; float brightness = dot(result, vec3(0.2126, 0.7152, 0.0722)); - if ( brightness < ubo.threshold ) result = vec3(0, 0, 0); - imageStore( imageBloom, texel, vec4( result, 1.0 ) ); + if( brightness < ubo.threshold ) result = vec3(0.0); + imageStore(imageBloom, texel, vec4(result, 1.0)); } else if ( mode == 1 ) { // bloom horizontal - - vec3 result = imageLoad( imageBloom, texel ).rgb * gauss( 0, sigma ) * strength; - for( int i = 1; i < ubo.size; ++i ) { - result += imageLoad( imageBloom, texel + ivec2(i * scale, 0.0)).rgb * gauss( i * scale / ubo.size, sigma ) * strength; - result += imageLoad( imageBloom, texel - ivec2(i * scale, 0.0)).rgb * gauss( i * scale / ubo.size, sigma ) * strength; - } - // write to PingPong - imageStore( imagePingPong, texel, vec4( result, 1.0 ) ); + vec3 result = imageLoad( imageBloom, texel ).rgb * ubo.weights[0]; + for ( int i = 1; i < int(ubo.size); ++i ) { + vec3 c1 = imageLoad( imageBloom, texel + ivec2(i, 0) ).rgb; + vec3 c2 = imageLoad( imageBloom, texel - ivec2(i, 0) ).rgb; + result += (c1 + c2) * ubo.weights[i]; + } + imageStore( imagePingPong, texel, vec4(result, 1.0) ); } else if ( mode == 2 ) { // bloom vertical - vec3 result = imageLoad( imagePingPong, texel ).rgb * gauss( 0, sigma ) * strength; - for( int i = 1; i < ubo.size; ++i ) { - result += imageLoad( imagePingPong, texel + ivec2(0.0, i * scale)).rgb * gauss( i * scale / ubo.size, sigma ) * strength; - result += imageLoad( imagePingPong, texel - ivec2(0.0, i * scale)).rgb * gauss( i * scale / ubo.size, sigma ) * strength; + vec3 result = imageLoad( imagePingPong, texel ).rgb * ubo.weights[0]; + for( int i = 1; i < int(ubo.size); ++i ) { + vec3 c1 = imageLoad( imagePingPong, texel + ivec2(0, i) ).rgb; + vec3 c2 = imageLoad( imagePingPong, texel - ivec2(0, i) ).rgb; + result += (c1 + c2) * ubo.weights[i]; } - // write to Bloom - imageStore( imageBloom, texel, vec4( result, 1.0 ) ); + imageStore(imageBloom, texel, vec4(result, 1.0)); } else if ( mode == 3 ) { // combine - vec3 result = imageLoad( imageColor, texel ).rgb + imageLoad( imageBloom, texel ).rgb; - imageStore( imageColor, texel, vec4( result, 1.0 ) ); + vec3 base = imageLoad( imageColor, texel ).rgb; + vec3 bloom = imageLoad( imageBloom, texel ).rgb; + imageStore( imageColor, texel, vec4(base + bloom, 1.0) ); } } \ No newline at end of file diff --git a/bin/data/shaders/display/deferred/comp/comp.h b/bin/data/shaders/display/deferred/comp/comp.h index 4a608c1a..9e37b8db 100644 --- a/bin/data/shaders/display/deferred/comp/comp.h +++ b/bin/data/shaders/display/deferred/comp/comp.h @@ -50,9 +50,9 @@ layout (constant_id = 1) const uint CUBEMAPS = 128; #endif -layout(binding = 7, rgba16f) uniform volatile coherent image2D imageColor; -layout(binding = 8, rgba16f) uniform volatile coherent image2D imageBright; -layout(binding = 9, rg16f) uniform volatile coherent image2D imageMotion; +layout(binding = 7, rgba16f) uniform writeonly image2D imageColor; +layout(binding = 8, rgba16f) uniform writeonly image2D imageBright; +layout(binding = 9, rg16f) uniform writeonly image2D imageMotion; layout( push_constant ) uniform PushBlock { uint pass; diff --git a/bin/data/shaders/display/depth-pyramid/comp.glsl b/bin/data/shaders/display/depth-pyramid/comp.glsl index d58feb89..708815b9 100644 --- a/bin/data/shaders/display/depth-pyramid/comp.glsl +++ b/bin/data/shaders/display/depth-pyramid/comp.glsl @@ -17,7 +17,7 @@ layout( push_constant ) uniform PushBlock { } PushConstant; layout (binding = 3) uniform sampler2D inImage[MIPS]; -layout (binding = 4, r32f) uniform volatile coherent image2D outImage[MIPS]; +layout (binding = 4, r32f) uniform writeonly image2D outImage[MIPS]; void main() { vec2 imageSize = imageSize(outImage[PushConstant.pass]); diff --git a/bin/data/shaders/display/vxgi/comp.glsl b/bin/data/shaders/display/vxgi/comp.glsl index 2d49e312..853acd20 100644 --- a/bin/data/shaders/display/vxgi/comp.glsl +++ b/bin/data/shaders/display/vxgi/comp.glsl @@ -10,7 +10,7 @@ layout (local_size_x = 8, local_size_y = 8, local_size_z = 8) in; #define VXGI 1 #define MAX_CUBEMAPS CUBEMAPS -#define GAMMA_CORRECT 0 +#define GAMMA_CORRECT 1 #define PBR 0 #define LAMBERT 1 @@ -52,16 +52,16 @@ layout (binding = 8) uniform sampler2D samplerTextures[TEXTURES]; layout (binding = 9) uniform samplerCube samplerCubemaps[CUBEMAPS]; layout (binding = 10) uniform sampler3D samplerNoise; -layout (binding = 11, r32ui) uniform volatile coherent uimage3D voxelDrawId[CASCADES]; -layout (binding = 12, r32ui) uniform volatile coherent uimage3D voxelInstanceId[CASCADES]; -layout (binding = 13, r32ui) uniform volatile coherent uimage3D voxelNormalX[CASCADES]; -layout (binding = 14, r32ui) uniform volatile coherent uimage3D voxelNormalY[CASCADES]; -layout (binding = 15, r32ui) uniform volatile coherent uimage3D voxelRadianceR[CASCADES]; -layout (binding = 16, r32ui) uniform volatile coherent uimage3D voxelRadianceG[CASCADES]; -layout (binding = 17, r32ui) uniform volatile coherent uimage3D voxelRadianceB[CASCADES]; -layout (binding = 18, r32ui) uniform volatile coherent uimage3D voxelRadianceA[CASCADES]; -layout (binding = 19, r32ui) uniform volatile coherent uimage3D voxelCount[CASCADES]; -layout (binding = 20, rgba8) uniform volatile coherent image3D voxelOutput[CASCADES]; +layout (binding = 11, r32ui) uniform readonly uimage3D voxelDrawId[CASCADES]; +layout (binding = 12, r32ui) uniform readonly uimage3D voxelInstanceId[CASCADES]; +layout (binding = 13, r32ui) uniform readonly uimage3D voxelNormalX[CASCADES]; +layout (binding = 14, r32ui) uniform readonly uimage3D voxelNormalY[CASCADES]; +layout (binding = 15, r32ui) uniform readonly uimage3D voxelRadianceR[CASCADES]; +layout (binding = 16, r32ui) uniform readonly uimage3D voxelRadianceG[CASCADES]; +layout (binding = 17, r32ui) uniform readonly uimage3D voxelRadianceB[CASCADES]; +layout (binding = 18, r32ui) uniform readonly uimage3D voxelRadianceA[CASCADES]; +layout (binding = 19, r32ui) uniform readonly uimage3D voxelCount[CASCADES]; +layout (binding = 20, rgba8) uniform writeonly image3D voxelOutput[CASCADES]; #include "../../common/functions.h" #include "../../common/light.h" diff --git a/bin/data/shaders/graph/baking/frag.h b/bin/data/shaders/graph/baking/frag.h index a4c8443b..14ff8eb5 100644 --- a/bin/data/shaders/graph/baking/frag.h +++ b/bin/data/shaders/graph/baking/frag.h @@ -52,7 +52,7 @@ layout (std140, binding = 15) readonly buffer Lights { Light lights[]; }; -layout (binding = 16, rgba8) uniform volatile coherent image3D outAlbedos; +layout (binding = 16, rgba8) uniform writeonly image3D outAlbedos; #if RT layout (binding = 17) uniform accelerationStructureEXT tlas; diff --git a/bin/data/shaders/graph/voxelize/frag.glsl b/bin/data/shaders/graph/voxelize/frag.glsl index e6a33b98..529642b3 100644 --- a/bin/data/shaders/graph/voxelize/frag.glsl +++ b/bin/data/shaders/graph/voxelize/frag.glsl @@ -37,16 +37,16 @@ layout (std140, binding = 12) readonly buffer Lights { Light lights[]; }; -layout (binding = 13, r32ui) uniform volatile coherent uimage3D voxelDrawId[CASCADES]; -layout (binding = 14, r32ui) uniform volatile coherent uimage3D voxelInstanceId[CASCADES]; -layout (binding = 15, r32ui) uniform volatile coherent uimage3D voxelNormalX[CASCADES]; -layout (binding = 16, r32ui) uniform volatile coherent uimage3D voxelNormalY[CASCADES]; -layout (binding = 17, r32ui) uniform volatile coherent uimage3D voxelRadianceR[CASCADES]; -layout (binding = 18, r32ui) uniform volatile coherent uimage3D voxelRadianceG[CASCADES]; -layout (binding = 19, r32ui) uniform volatile coherent uimage3D voxelRadianceB[CASCADES]; -layout (binding = 20, r32ui) uniform volatile coherent uimage3D voxelRadianceA[CASCADES]; -layout (binding = 21, r32ui) uniform volatile coherent uimage3D voxelCount[CASCADES]; -layout (binding = 22, rgba8) uniform volatile coherent image3D voxelOutput[CASCADES]; +layout (binding = 13, r32ui) uniform volatile uimage3D voxelDrawId[CASCADES]; +layout (binding = 14, r32ui) uniform volatile uimage3D voxelInstanceId[CASCADES]; +layout (binding = 15, r32ui) uniform volatile uimage3D voxelNormalX[CASCADES]; +layout (binding = 16, r32ui) uniform volatile uimage3D voxelNormalY[CASCADES]; +layout (binding = 17, r32ui) uniform volatile uimage3D voxelRadianceR[CASCADES]; +layout (binding = 18, r32ui) uniform volatile uimage3D voxelRadianceG[CASCADES]; +layout (binding = 19, r32ui) uniform volatile uimage3D voxelRadianceB[CASCADES]; +layout (binding = 20, r32ui) uniform volatile uimage3D voxelRadianceA[CASCADES]; +layout (binding = 21, r32ui) uniform volatile uimage3D voxelCount[CASCADES]; +layout (binding = 22, rgba8) uniform writeonly image3D voxelOutput[CASCADES]; layout (location = 0) flat in uvec4 inId; layout (location = 1) flat in vec4 inPOS0; diff --git a/bin/data/shaders/graph/voxelize/geom.glsl b/bin/data/shaders/graph/voxelize/geom.glsl index 286a482e..284268b8 100644 --- a/bin/data/shaders/graph/voxelize/geom.glsl +++ b/bin/data/shaders/graph/voxelize/geom.glsl @@ -57,10 +57,11 @@ void main(){ #endif const uint CASCADE = inId[0].w; + const float power = cascadePower(CASCADE); vec3 P[3] = { - vec3( ubo.voxel * vec4( inPosition[0], 1 ) ) / cascadePower(CASCADE), - vec3( ubo.voxel * vec4( inPosition[1], 1 ) ) / cascadePower(CASCADE), - vec3( ubo.voxel * vec4( inPosition[2], 1 ) ) / cascadePower(CASCADE), + vec3( ubo.voxel * vec4( inPosition[0], 1 ) ) / power, + vec3( ubo.voxel * vec4( inPosition[1], 1 ) ) / power, + vec3( ubo.voxel * vec4( inPosition[2], 1 ) ) / power, }; #pragma unroll 3 diff --git a/bin/data/shaders/raytrace/shader.ray-gen.glsl b/bin/data/shaders/raytrace/shader.ray-gen.glsl index 988398a0..b9b9285b 100644 --- a/bin/data/shaders/raytrace/shader.ray-gen.glsl +++ b/bin/data/shaders/raytrace/shader.ray-gen.glsl @@ -32,7 +32,7 @@ layout( push_constant ) uniform PushBlock { layout (binding = 0) uniform accelerationStructureEXT tlas; -layout (binding = 1, rgba32f) uniform volatile coherent image2D outImage; +layout (binding = 1, rgba32f) uniform image2D outImage; layout (binding = 2) uniform UBO { EyeMatrices eyes[2]; diff --git a/engine/inc/uf/engine/graph/graph.h b/engine/inc/uf/engine/graph/graph.h index b3f86aa5..8a6253ba 100644 --- a/engine/inc/uf/engine/graph/graph.h +++ b/engine/inc/uf/engine/graph/graph.h @@ -29,8 +29,6 @@ namespace pod { // Render information uf::stl::vector primitives; // - uf::stl::vector instances; // ugh - uf::stl::vector drawCommands; // ugh uf::stl::vector meshes; // uf::stl::vector images; // @@ -81,10 +79,8 @@ namespace pod { // Local storage, used for save/load struct Storage { - uf::stl::KeyMap instances; // should be unused but is a counter for global instance IDs because I cannot for the life of me figure out a working way to remap otherwise - uf::stl::KeyMap instanceAddresses; + uf::stl::KeyMap> instanceAddresses; uf::stl::KeyMap> primitives; - uf::stl::KeyMap> drawCommands; // unused uf::stl::KeyMap meshes; uf::stl::KeyMap images; @@ -137,7 +133,7 @@ namespace uf { pod::Matrix4f UF_API matrix( pod::Graph&, int32_t ); // void UF_API process( uf::Object& entity ); - void UF_API initializeGraphics( pod::Graph& graph, uf::Object& entity, uf::Mesh& mesh ); + void UF_API initializeGraphics( pod::Graph& graph, uf::Object& entity, uf::Mesh& mesh, uf::stl::vector& ); void UF_API process( pod::Graph& graph ); void UF_API process( pod::Graph& graph, int32_t, uf::Object& parent ); void UF_API reload( pod::Graph& ); diff --git a/engine/inc/uf/utils/math/vector/pod.inl b/engine/inc/uf/utils/math/vector/pod.inl index db446454..3e30b0a2 100644 --- a/engine/inc/uf/utils/math/vector/pod.inl +++ b/engine/inc/uf/utils/math/vector/pod.inl @@ -483,7 +483,7 @@ T& uf::vector::negate_( T& vector ) { template T& uf::vector::normalize_( T& vector ) { typename T::type_t norm = uf::vector::norm(vector); - return ( norm == 0 ) ? T{} : ( vector = uf::vector::divide((const T&) vector, norm) ); + return ( norm < 1.0e-6 ) ? T{} : ( vector = uf::vector::divide((const T&) vector, norm) ); } template T uf::vector::min( const T& left, const T& right ) { @@ -679,7 +679,7 @@ T uf::vector::normalize( const T& vector ) { } #endif typename T::type_t norm = uf::vector::norm(vector); - if ( norm == 0 ) return vector; + if ( norm < 1.0e-6 ) return vector; #if UF_ENV_DREAMCAST && UF_ENV_DREAMCAST_SIMD if constexpr ( std::is_same_v ) { return uf::vector::multiply(vector, MATH_fsrra(norm)); diff --git a/engine/inc/uf/utils/memory/key_map.h b/engine/inc/uf/utils/memory/key_map.h index cfefa373..970f1958 100644 --- a/engine/inc/uf/utils/memory/key_map.h +++ b/engine/inc/uf/utils/memory/key_map.h @@ -17,7 +17,9 @@ namespace uf { T& operator[]( const Key& key ); void reserve( size_t i ); + size_t id( const Key& ); uf::stl::vector flatten() const; + uf::stl::vector flattenByIndex() const; void clear(); }; } @@ -40,6 +42,19 @@ void uf::stl::KeyMap::reserve( size_t i ) { map.reserve(i); } +template +size_t uf::stl::KeyMap::id( const Key& key ) { + if ( indices.count( key ) > 0 ) return indices[key]; + + size_t newIndex = keys.size(); + + indices[key] = newIndex; + keys.emplace_back( key ); + map[key]; + + return newIndex; +} + template void uf::stl::KeyMap::clear() { keys.clear(); @@ -58,5 +73,16 @@ uf::stl::vector uf::stl::KeyMap::flatten() const { res.emplace_back(map.at(key)); } } + return res; +} + +template +uf::stl::vector uf::stl::KeyMap::flattenByIndex() const { + vector res(indices.size()); + + for (auto& [key, index] : indices) { + res[index] = map.at(key); + } + return res; } \ No newline at end of file diff --git a/engine/src/engine/ext/scene/behavior.cpp b/engine/src/engine/ext/scene/behavior.cpp index 57db8d2e..e86d412c 100644 --- a/engine/src/engine/ext/scene/behavior.cpp +++ b/engine/src/engine/ext/scene/behavior.cpp @@ -947,6 +947,8 @@ void ext::ExtSceneBehavior::bindBuffers( uf::Object& self, uf::renderer::Graphic float smoothness; uint32_t size; uint32_t padding; + + float weights[32]; }; UniformDescriptor uniforms = { @@ -954,6 +956,24 @@ void ext::ExtSceneBehavior::bindBuffers( uf::Object& self, uf::renderer::Graphic .smoothness = metadata.bloom.smoothness, .size = metadata.bloom.size, }; + + uint kernelSize = uniforms.size * 2 + 1; + float sigma = kernelSize / (6 * std::max(0.001f, uniforms.smoothness)); + + float sum = 0; + uf::stl::vector tempWeights(uniforms.size); + tempWeights[0] = (1.0f / (sqrtf(2.0f * M_PI) * sigma)); + sum = tempWeights[0]; + + for ( auto i = 1; i < uniforms.size; ++i ) { + float x = (float) i; + tempWeights[i] = (expf(-(x * x) / (2.0f * sigma * sigma)) / (sqrtf(2.0f * M_PI) * sigma)); + sum += 2.0f * tempWeights[i]; + } + + // normalize + for ( auto i = 0; i < uniforms.size; ++i ) uniforms.weights[i] = tempWeights[i] / sum; + metadata.bloom.outOfDate = false; if ( shader.hasUniform("UBO") ) shader.updateBuffer( (const void*) &uniforms, sizeof(uniforms), shader.getUniformBuffer("UBO") ); } diff --git a/engine/src/engine/graph/decode.cpp b/engine/src/engine/graph/decode.cpp index 21f5a977..924e259b 100644 --- a/engine/src/engine/graph/decode.cpp +++ b/engine/src/engine/graph/decode.cpp @@ -440,13 +440,9 @@ void uf::graph::load( pod::Graph& graph, const uf::stl::string& filename, const auto name = key + value["name"].as(); // UF_MSG_DEBUG("{}", name); storage.primitives[name] = decodePrimitives( value, graph ); - graph.primitives.emplace_back(name); + storage.instanceAddresses[name] = {}; - uf::stl::vector drawCommands; - for ( auto& primitive : storage.primitives[name] ) { - drawCommands.emplace_back( primitive.drawCommand ); - } - storage.drawCommands[name] = std::move(drawCommands); + graph.primitives.emplace_back(name); }); UF_DEBUG_TIMER_MULTITRACE("Read primitives."); #if UF_ENV_DREAMCAST @@ -506,6 +502,9 @@ void uf::graph::load( pod::Graph& graph, const uf::stl::string& filename, const // UF_MSG_DEBUG("{}", name); storage.textures[name] = decodeTexture( value, graph ); graph.textures.emplace_back(name); + + // pre-allocate + storage.texture2Ds[name]; }); UF_DEBUG_TIMER_MULTITRACE("Read texture information"); #if UF_ENV_DREAMCAST diff --git a/engine/src/engine/graph/graph.cpp b/engine/src/engine/graph/graph.cpp index 2c4578a1..b0f23702 100644 --- a/engine/src/engine/graph/graph.cpp +++ b/engine/src/engine/graph/graph.cpp @@ -44,11 +44,16 @@ namespace { return hash; } - uf::stl::string objectKey( pod::Node& node ){ - return std::to_string(node.object); + size_t allocateObjectID( pod::Graph::Storage& storage ) { + return storage.entities.keys.size(); } - uf::stl::string instanceKey( uint32_t instanceID ) { - return std::to_string( instanceID ); + size_t allocateInstanceID( pod::Graph::Storage& storage, const uf::stl::string& name ) { + size_t instanceID = 0; + for ( auto& key : storage.primitives.keys ) { + if ( key == name ) break; + instanceID += storage.primitives.map[key].size(); + } + return instanceID; } pod::Graph::Storage& getGraphStorage( uf::Object& object ) { @@ -523,52 +528,49 @@ namespace { shader.aliasBuffer( "light", storage.buffers.light ); } } + #endif + } - // rt pipeline - // to-do: segregate out buffer updating code + void bindInstanceAddresses( uf::renderer::Graphic& graphic, uf::Mesh& mesh, uf::stl::vector& addresses ) { + if ( !uf::renderer::settings::invariant::deviceAddressing || !mesh.indirect.count ) return; + addresses.resize( mesh.indirect.count ); - // grab addresses - if ( uf::renderer::settings::invariant::deviceAddressing ) { - pod::DrawCommand* drawCommands = (pod::DrawCommand*) mesh.getBuffer( mesh.indirect ).data(); - for ( size_t drawID = 0; drawID < mesh.indirect.count; ++drawID ) { - auto& drawCommand = drawCommands[drawID]; - auto instanceID = drawCommand.instanceID; - auto instanceKeyName = ::instanceKey( instanceID ); // it *should* be fine as the instance IDs are already objective to the scene graph + pod::DrawCommand* drawCommands = (pod::DrawCommand*) mesh.getBuffer( mesh.indirect ).data(); + for ( size_t drawID = 0; drawID < mesh.indirect.count; ++drawID ) { + auto& drawCommand = drawCommands[drawID]; + auto instanceID = drawCommand.instanceID; - auto& instanceAddresses = storage.instanceAddresses[instanceKeyName]; - if ( mesh.vertex.count ) { - if ( mesh.isInterleaved( mesh.vertex ) ) { - instanceAddresses.vertex = graphic.buffers.at(graphic.descriptor.inputs.vertex.interleaved).getAddress(); - } else { - for ( auto& attribute : graphic.descriptor.inputs.vertex.attributes ) { - if ( attribute.buffer < 0 ) continue; - if ( attribute.descriptor.name == "position" ) instanceAddresses.position = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "uv" ) instanceAddresses.uv = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "color" ) instanceAddresses.color = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "st" ) instanceAddresses.st = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "normal" ) instanceAddresses.normal = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "tangent" ) instanceAddresses.tangent = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "joints" ) instanceAddresses.joints = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "weights" ) instanceAddresses.weights = graphic.buffers.at(attribute.buffer).getAddress(); - else if ( attribute.descriptor.name == "id" ) instanceAddresses.id = graphic.buffers.at(attribute.buffer).getAddress(); - } + auto& instanceAddresses = addresses[drawID]; // THIS IS WRONG (to-do: actually use instanceIDs) + if ( mesh.vertex.count ) { + if ( mesh.isInterleaved( mesh.vertex ) ) { + instanceAddresses.vertex = graphic.buffers.at(graphic.descriptor.inputs.vertex.interleaved).getAddress(); + } else { + for ( auto& attribute : graphic.descriptor.inputs.vertex.attributes ) { + if ( attribute.buffer < 0 ) continue; + if ( attribute.descriptor.name == "position" ) instanceAddresses.position = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "uv" ) instanceAddresses.uv = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "color" ) instanceAddresses.color = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "st" ) instanceAddresses.st = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "normal" ) instanceAddresses.normal = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "tangent" ) instanceAddresses.tangent = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "joints" ) instanceAddresses.joints = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "weights" ) instanceAddresses.weights = graphic.buffers.at(attribute.buffer).getAddress(); + else if ( attribute.descriptor.name == "id" ) instanceAddresses.id = graphic.buffers.at(attribute.buffer).getAddress(); } } - if ( mesh.index.count ) { - if ( mesh.isInterleaved( mesh.index ) ) instanceAddresses.index = graphic.buffers.at(graphic.descriptor.inputs.index.interleaved).getAddress(); - else instanceAddresses.index = graphic.buffers.at(graphic.descriptor.inputs.index.attributes.front().buffer).getAddress(); - } + } + if ( mesh.index.count ) { + if ( mesh.isInterleaved( mesh.index ) ) instanceAddresses.index = graphic.buffers.at(graphic.descriptor.inputs.index.interleaved).getAddress(); + else instanceAddresses.index = graphic.buffers.at(graphic.descriptor.inputs.index.attributes.front().buffer).getAddress(); + } - if ( mesh.indirect.count ) { - if ( mesh.isInterleaved( mesh.indirect ) ) instanceAddresses.indirect = graphic.buffers.at(graphic.descriptor.inputs.indirect.interleaved).getAddress(); - else instanceAddresses.indirect = graphic.buffers.at(graphic.descriptor.inputs.indirect.attributes.front().buffer).getAddress(); - - instanceAddresses.drawID = drawID; - } + if ( mesh.indirect.count ) { + if ( mesh.isInterleaved( mesh.indirect ) ) instanceAddresses.indirect = graphic.buffers.at(graphic.descriptor.inputs.indirect.interleaved).getAddress(); + else instanceAddresses.indirect = graphic.buffers.at(graphic.descriptor.inputs.indirect.attributes.front().buffer).getAddress(); + instanceAddresses.drawID = drawID; } } - #endif } } @@ -683,7 +685,7 @@ UF_VERTEX_INTERPOLATE(uf::graph::mesh::Skinned_u16q, { return t < 0.5 ? p1 : p2; }) -void uf::graph::initializeGraphics( pod::Graph& graph, uf::Object& entity, uf::Mesh& mesh ) { +void uf::graph::initializeGraphics( pod::Graph& graph, uf::Object& entity, uf::Mesh& mesh, uf::stl::vector& addresses ) { auto& scene = uf::scene::getCurrentScene(); auto& sceneTextures = scene.getComponent(); auto& sceneMetadataJson = scene.getComponent(); @@ -730,6 +732,7 @@ void uf::graph::initializeGraphics( pod::Graph& graph, uf::Object& entity, uf::M ::bindTextures( graphic ); ::bindShaders( graph, entity, mesh ); ::bindBuffers( graphic, mesh ); + ::bindInstanceAddresses( graphic, mesh, addresses ); graphic.process = true; } @@ -790,12 +793,6 @@ void uf::graph::process( pod::Graph& graph ) { uf::stl::unordered_map lightmapIDs; uint32_t lightmapCount = 0; - for ( auto& name : graph.instances ) { - auto& instance = storage.instances[name]; - filenames[instance.auxID] = uf::string::replace(UF_GRAPH_DEFAULT_LIGHTMAP, "%i", std::to_string(instance.auxID)); - - lightmapCount = std::max( lightmapCount, instance.auxID + 1 ); - } for ( auto& name : graph.primitives ) { auto& primitives = storage.primitives[name]; for ( auto& primitive : primitives ) { @@ -869,11 +866,6 @@ void uf::graph::process( pod::Graph& graph ) { } } - for ( auto& name : graph.instances ) { - auto& instance = storage.instances[name]; - if ( lightmapIDs.count( instance.auxID ) == 0 ) continue; - instance.lightmapID = lightmapIDs[instance.auxID]; - } for ( auto& name : graph.primitives ) { auto& primitives = storage.primitives[name]; for ( auto& primitive : primitives ) { @@ -886,12 +878,6 @@ void uf::graph::process( pod::Graph& graph ) { // setup textures storage.texture2Ds.reserve( storage.images.map.size() ); - // not having this block will cause our images and texture2Ds to be ordered out of sync - for ( auto& key : graph.images ) { - storage.images[key]; - storage.texture2Ds[key]; - } - // figure out what texture is what exactly UF_DEBUG_TIMER_MULTITRACE("Determining format of textures"); for ( auto& key : graph.materials ) { @@ -951,7 +937,7 @@ void uf::graph::process( pod::Graph& graph ) { } // patch materials/textures - UF_DEBUG_TIMER_MULTITRACE("Remapping patching/textures"); + UF_DEBUG_TIMER_MULTITRACE("Patching textures/materials"); for ( auto& name : graph.materials ) { auto& material = storage.materials[name]; auto tag = ext::json::find( name, graphMetadataJson["tags"] ); @@ -1021,7 +1007,7 @@ void uf::graph::process( pod::Graph& graph ) { } // remap textures->images IDs - UF_DEBUG_TIMER_MULTITRACE("Remapping textures"); + UF_DEBUG_TIMER_MULTITRACE("Remapping texture -> image IDs"); for ( auto& name : graph.textures ) { auto& texture = storage.textures[name]; auto& keys = storage.images.keys; @@ -1034,7 +1020,7 @@ void uf::graph::process( pod::Graph& graph ) { } // remap materials->texture IDs - UF_DEBUG_TIMER_MULTITRACE("Remapping materials"); + UF_DEBUG_TIMER_MULTITRACE("Remapping material -> texture IDs"); for ( auto& name : graph.materials ) { auto& material = storage.materials[name]; auto& keys = storage.textures.keys; @@ -1047,82 +1033,45 @@ void uf::graph::process( pod::Graph& graph ) { ID = indices[needle]; } } + // remap instance variables - UF_DEBUG_TIMER_MULTITRACE("Remapping instances"); - for ( auto& name : graph.instances ) { - auto& instance = storage.instances[name]; - - if ( 0 <= instance.materialID && instance.materialID < graph.materials.size() ) { - auto& keys = storage.materials.keys; - auto& indices = storage.materials.indices; + UF_DEBUG_TIMER_MULTITRACE("Remapping instance -> material IDs"); + for ( auto& name : graph.primitives ) { + for ( auto& primitive : storage.primitives[name] ) { + auto& instance = primitive.instance; - if ( !(0 <= instance.materialID && instance.materialID < graph.materials.size()) ) continue; - - auto& needle = graph.materials[instance.materialID]; - instance.materialID = indices[needle]; - } - if ( 0 <= instance.lightmapID && instance.lightmapID < graph.textures.size() ) { - auto& keys = storage.textures.keys; - auto& indices = storage.textures.indices; - - if ( !(0 <= instance.lightmapID && instance.lightmapID < graph.textures.size()) ) continue; - - auto& needle = graph.textures[instance.lightmapID]; - instance.lightmapID = indices[needle]; - } - #if 0 - // i genuinely dont remember what this is used for - - if ( 0 <= instance.imageID && instance.imageID < graph.images.size() ) { - auto& keys = storage.images.keys; - auto it = std::find( keys.begin(), keys.end(), graph.images[instance.imageID] ); - UF_ASSERT( it != keys.end() ); - instance.imageID = it - keys.begin(); - } - #endif - // remap a skinID as an actual jointID - if ( 0 <= instance.jointID && instance.jointID < graph.skins.size() ) { - auto& name = graph.skins[instance.jointID]; - instance.jointID = 0; - for ( auto key : storage.joints.keys ) { - if ( key == name ) break; - auto& joints = storage.joints[key]; - instance.jointID += joints.size(); - } - } - } - - // remap draw commands -#if 0 - UF_DEBUG_TIMER_MULTITRACE("Remapping drawCommands"); - for ( auto& name : graph.drawCommands ) { - auto& drawCommands = storage.drawCommands[name]; - for ( auto& drawCommand : drawCommands ) { - if ( 0 <= drawCommand.instanceID && drawCommand.instanceID < graph.instances.size() ) { - auto& keys = storage.instances.keys; - auto& indices = storage.instances.indices; + if ( 0 <= instance.materialID && instance.materialID < graph.materials.size() ) { + auto& keys = storage.materials.keys; + auto& indices = storage.materials.indices; - if ( !(0 <= drawCommand.instanceID && drawCommand.instanceID < graph.instances.size()) ) continue; + if ( !(0 <= instance.materialID && instance.materialID < graph.materials.size()) ) continue; - auto& needle = graph.instances[drawCommand.instanceID]; - #if 1 - drawCommand.instanceID = indices[needle]; - #elif 1 - for ( size_t i = 0; i < keys.size(); ++i ) { - if ( keys[i] != needle ) continue; - drawCommand.instanceID = i; - break; + auto& needle = graph.materials[instance.materialID]; + instance.materialID = indices[needle]; + } + if ( 0 <= instance.lightmapID && instance.lightmapID < graph.textures.size() ) { + auto& keys = storage.textures.keys; + auto& indices = storage.textures.indices; + + if ( !(0 <= instance.lightmapID && instance.lightmapID < graph.textures.size()) ) continue; + + auto& needle = graph.textures[instance.lightmapID]; + instance.lightmapID = indices[needle]; + } + + // remap a skinID as an actual jointID + if ( 0 <= instance.jointID && instance.jointID < graph.skins.size() ) { + auto& name = graph.skins[instance.jointID]; + instance.jointID = 0; + for ( auto key : storage.joints.keys ) { + if ( key == name ) break; + auto& joints = storage.joints[key]; + instance.jointID += joints.size(); } - #else - auto it = std::find( keys.begin(), keys.end(), needle ); - UF_ASSERT( it != keys.end() ); - drawCommand.instanceID = it - keys.begin(); - #endif } } } -#endif - +/* if ( graphMetadataJson["debug"]["print"]["lights"].as() ) { UF_MSG_DEBUG("Lights: {}", graph.lights.size()); for ( auto& pair : graph.lights ) { @@ -1135,17 +1084,6 @@ void uf::graph::process( pod::Graph& graph ) { UF_MSG_DEBUG("\tMesh: {}", name); } } - - if ( graphMetadataJson["debug"]["print"]["instances"].as() ) { - UF_MSG_DEBUG("Instances: {}", graph.instances.size()); - for ( auto& name : graph.instances ) { - auto& instance = storage.instances[name]; - UF_MSG_DEBUG("\tInstance: {} | {} | {}", name, - instance.materialID, - instance.lightmapID - ); - } - } if ( graphMetadataJson["debug"]["print"]["materials"].as() ) { UF_MSG_DEBUG("Materials: {}", graph.materials.size()); for ( auto& name : graph.materials ) { @@ -1166,6 +1104,7 @@ void uf::graph::process( pod::Graph& graph ) { UF_MSG_DEBUG("\tImage: {}", name); } } +*/ UF_DEBUG_TIMER_MULTITRACE("Updating master graph"); #if UF_GRAPH_EXTENDED @@ -1173,7 +1112,6 @@ void uf::graph::process( pod::Graph& graph ) { #endif uf::graph::reload(); - storage.instanceAddresses.keys = storage.instances.keys; UF_DEBUG_TIMER_MULTITRACE_END("Processed graph."); } void uf::graph::process( pod::Graph& graph, int32_t index, uf::Object& parent ) { @@ -1319,30 +1257,32 @@ void uf::graph::process( pod::Graph& graph, int32_t index, uf::Object& parent ) } // what a mess - node.object = storage.entities.keys.size(); - auto objectKeyName = ::objectKey( node ); - storage.entities[objectKeyName] = &entity; - auto& objectData = storage.objects[objectKeyName]; auto model = uf::transform::model( transform ); - objectData.model = model; - objectData.previous = model; // if ( 0 <= node.mesh && node.mesh < graph.meshes.size() ) { - auto model = uf::transform::model( transform ); + { + node.object = ::allocateObjectID( storage ); + auto objectKeyName = std::to_string( node.object ); + + storage.entities[objectKeyName] = &entity; + storage.objects[objectKeyName] = pod::Instance::Object{ + .model = model, + .previous = model, + }; + } + auto& mesh = storage.meshes.map[graph.meshes[node.mesh]]; auto& primitives = storage.primitives.map[graph.primitives[node.mesh]]; + auto& instanceAddresses = storage.instanceAddresses.map[graph.primitives[node.mesh]]; pod::Instance::Bounds bounds = {}; + size_t baseInstanceID = ::allocateInstanceID( storage, graph.primitives[node.mesh] ); // setup instances for ( auto i = 0; i < primitives.size(); ++i ) { auto& primitive = primitives[i]; - - size_t instanceID = storage.instances.keys.size(); - auto instanceKeyName = graph.instances.emplace_back(std::to_string(instanceID)); - - auto& instance = storage.instances[instanceKeyName]; - instance = primitive.instance; + auto& instance = primitive.instance; + size_t instanceID = baseInstanceID + i; instance.objectID = node.object; instance.jointID = graphMetadataJson["renderer"]["skinned"].as() ? 0 : -1; @@ -1360,7 +1300,7 @@ void uf::graph::process( pod::Graph& graph, int32_t index, uf::Object& parent ) } #if !UF_GRAPH_EXTENDED if ( graphMetadataJson["renderer"]["render"].as() ) { - uf::graph::initializeGraphics( graph, entity, mesh ); + uf::graph::initializeGraphics( graph, entity, mesh, addresses ); } #endif @@ -1452,29 +1392,67 @@ void uf::graph::tick( uf::Object& object ) { } bool uf::graph::tick( pod::Graph::Storage& storage ) { bool rebuild = false; - uf::stl::vector instances = storage.instances.flatten(); - uf::stl::vector instanceAddresses = storage.instanceAddresses.flatten(); - uf::stl::vector joints; joints.reserve(storage.joints.map.size()); - uf::stl::vector objects = storage.objects.flatten(); + static thread_local uf::stl::vector instances; + static thread_local uf::stl::vector instanceAddresses; + static thread_local uf::stl::vector joints; + static thread_local uf::stl::vector objects; + static thread_local uf::stl::vector materials; + static thread_local uf::stl::vector textures; + static thread_local uf::stl::vector drawCommands; + + joints.clear(); for ( auto& key : storage.joints.keys ) { - auto& matrices = storage.joints.map[key]; - joints.reserve( joints.size() + matrices.size() ); - for ( auto& mat : matrices ) joints.emplace_back( mat ); + joints.insert( joints.end(), storage.joints.map[key].begin(), storage.joints.map[key].end() ); } - - rebuild = storage.buffers.instance.update( (const void*) instances.data(), instances.size() * sizeof(pod::Instance) ) || rebuild; - rebuild = storage.buffers.instanceAddresses.update( (const void*) instanceAddresses.data(), instanceAddresses.size() * sizeof(pod::Instance::Addresses) ) || rebuild; - rebuild = storage.buffers.joint.update( (const void*) joints.data(), joints.size() * sizeof(pod::Matrix4f) ) || rebuild; + + objects.clear(); + for ( auto& key : storage.objects.keys ) { + auto& entity = *storage.entities.map[key]; + auto& object = storage.objects.map[key]; + + if ( entity.isValid() ) { + auto& metadata = entity.getComponent(); + auto& transform = entity.getComponent>(); + + if ( !metadata.system.ignoreGraph ) { + object.previous = object.model; + object.model = uf::transform::model( transform ); + } + } + + objects.emplace_back( object ); + } + + if ( !joints.empty() ) rebuild = storage.buffers.joint.update( (const void*) joints.data(), joints.size() * sizeof(pod::Matrix4f) ) || rebuild; rebuild = storage.buffers.object.update( (const void*) objects.data(), objects.size() * sizeof(pod::Instance::Object) ) || rebuild; if ( ::newGraphAdded ) { - uf::stl::vector materials = storage.materials.flatten(); - uf::stl::vector textures = storage.textures.flatten(); - uf::stl::vector drawCommands; drawCommands.reserve(storage.drawCommands.map.size()); + instances.clear(); + for ( auto& key : storage.primitives.keys ) { + for ( auto& primitive : storage.primitives[key] ) { + instances.emplace_back( primitive.instance ); + } + } - for ( auto& key : storage.drawCommands.keys ) drawCommands.insert( drawCommands.end(), storage.drawCommands.map[key].begin(), storage.drawCommands.map[key].end() ); + instanceAddresses.clear(); + for ( auto& key : storage.instanceAddresses.keys ) { + instanceAddresses.insert( instanceAddresses.end(), storage.instanceAddresses.map[key].begin(), storage.instanceAddresses.map[key].end() ); + } + textures.clear(); + for ( auto& key : storage.textures.keys ) textures.emplace_back( storage.textures.map[key] ); + + materials.clear(); + for ( auto& key : storage.materials.keys ) materials.emplace_back( storage.materials.map[key] ); + + drawCommands.clear(); + for ( auto& key : storage.primitives.keys ) { + for ( auto& primitive : storage.primitives[key] ) drawCommands.emplace_back( primitive.drawCommand ); + } + + rebuild = storage.buffers.instance.update( (const void*) instances.data(), instances.size() * sizeof(pod::Instance) ) || rebuild; + rebuild = storage.buffers.instanceAddresses.update( (const void*) instanceAddresses.data(), instanceAddresses.size() * sizeof(pod::Instance::Addresses) ) || rebuild; rebuild = storage.buffers.drawCommands.update( (const void*) drawCommands.data(), drawCommands.size() * sizeof(pod::DrawCommand) ) || rebuild; rebuild = storage.buffers.material.update( (const void*) materials.data(), materials.size() * sizeof(pod::Material) ) || rebuild; rebuild = storage.buffers.texture.update( (const void*) textures.data(), textures.size() * sizeof(pod::Texture) ) || rebuild; @@ -1482,6 +1460,7 @@ bool uf::graph::tick( pod::Graph::Storage& storage ) { ::newGraphAdded = false; } + if ( rebuild ) { UF_MSG_DEBUG("Graph buffers requesting renderer update"); uf::renderer::states::rebuild = true; @@ -1633,6 +1612,7 @@ void uf::graph::reload( pod::Graph& graph, pod::Node& node ) { auto model = uf::transform::model( transform ); auto& mesh = storage.meshes.map[graph.meshes[node.mesh]]; auto& primitives = storage.primitives.map[graph.primitives[node.mesh]]; + auto& instanceAddresses = storage.instanceAddresses.map[graph.primitives[node.mesh]]; float radius = graph.settings.stream.radius; float radiusSquared = radius * radius; @@ -1932,10 +1912,12 @@ void uf::graph::reload( pod::Graph& graph, pod::Node& node ) { // update buffers if any of them were resized (because my aliasing system is weak) if ( rebuild ) { ::bindBuffers( graphic, mesh ); + + ::bindInstanceAddresses( graphic, mesh, instanceAddresses ); uf::renderer::states::rebuild = true; } } else { - uf::graph::initializeGraphics( graph, entity, mesh ); + uf::graph::initializeGraphics( graph, entity, mesh, instanceAddresses ); } } // bind mesh to physics state @@ -1996,8 +1978,10 @@ void uf::graph::update( pod::Graph& graph, float delta ) { auto& graphic = entity.getComponent(); auto& mesh = storage.meshes.map[graph.meshes[node.mesh]]; + auto& instanceAddresses = storage.instanceAddresses.map[graph.primitives[node.mesh]]; ::bindBuffers( graphic, mesh ); + ::bindInstanceAddresses( graphic, mesh, instanceAddresses ); } } @@ -2009,50 +1993,5 @@ void uf::graph::update( pod::Graph& graph, float delta ) { } #endif -#if !UF_ENV_DREAMCAST - // update instance model - - uf::stl::unordered_map matrixCache; - for ( auto& node : graph.nodes ) { - if ( !( 0 <= node.object && node.object < storage.objects.keys.size() ) ) continue; - auto objectKeyName = ::objectKey( node ); - - auto& entity = *storage.entities[objectKeyName]; - auto& object = storage.objects[objectKeyName]; - - if ( !entity.isValid() ) continue; - auto& metadata = entity.getComponent(); - if ( metadata.system.ignoreGraph ) continue; - - auto& transform = entity.getComponent>(); - - object.previous = object.model; - object.model = uf::transform::model( transform ); - - } - /* - for ( auto& name : graph.primitives ) { - auto& primitives = storage.primitives[name]; - for ( auto& primitive : primitives ) { - // auto& drawCommand = primitive.drawCommand; - auto& instance = primitive.instance; - instance.previous = instance.model; - if ( instanceCache.count( instance.objectID ) > 0 ) { - instance.model = instanceCache[instance.objectID]; - continue; - } - - auto& entity = *storage.entities[std::to_string(instance.objectID)]; - if ( !entity.isValid() ) continue; - auto& metadata = entity.getComponent(); - if ( metadata.system.ignoreGraph ) continue; - - auto& transform = entity.getComponent>(); - instance.model = (instanceCache[instance.objectID] = uf::transform::model( transform )); - } - } - */ -#endif - uf::graph::updateAnimation( graph, delta ); } diff --git a/engine/src/engine/object/behavior.cpp b/engine/src/engine/object/behavior.cpp index da177cae..c26da4d0 100644 --- a/engine/src/engine/object/behavior.cpp +++ b/engine/src/engine/object/behavior.cpp @@ -139,6 +139,8 @@ void uf::ObjectBehavior::initialize( uf::Object& self ) { if ( type == "bounding box" || type == "aabb" ) { pod::Vector3f min = uf::vector::decode( metadataJsonPhysics["min"], pod::Vector3f{-0.5f, -0.5f, -0.5f} ); pod::Vector3f max = uf::vector::decode( metadataJsonPhysics["max"], pod::Vector3f{0.5f, 0.5f, 0.5f} ); + + UF_MSG_DEBUG("entity={}, min={}, max={}", uf::string::toString( *this ), uf::vector::toString( min ), uf::vector::toString( max )); #if UF_USE_REACTPHYSICS auto center = ( max + min ) * 0.5f; if ( metadataJsonPhysics["recenter"].as(true) ) offset = (center - transform.position); diff --git a/engine/src/ext/lua/usertypes/physics.cpp b/engine/src/ext/lua/usertypes/physics.cpp index 6b3602cc..cb257a82 100644 --- a/engine/src/ext/lua/usertypes/physics.cpp +++ b/engine/src/ext/lua/usertypes/physics.cpp @@ -7,6 +7,9 @@ namespace binds { pod::Vector3f& velocity( pod::PhysicsBody& self ) { return self.velocity; } void setVelocity( pod::PhysicsBody& self, const pod::Vector3f& v ) { self.velocity = v; } void applyVelocity( pod::PhysicsBody& self, const pod::Vector3f& v ) { self.velocity += v; } + + float getMass( const pod::PhysicsBody& self ) { return self.mass; } + void setMass( pod::PhysicsBody& self, float mass ) { self.mass = mass; } void enableGravity( pod::PhysicsBody& state, bool s ) { if ( !state.object ) return; @@ -44,13 +47,10 @@ UF_LUA_REGISTER_USERTYPE_AND_COMPONENT(pod::PhysicsBody, UF_LUA_REGISTER_USERTYPE_DEFINE( applyImpulse, UF_LUA_C_FUN(uf::physics::impl::applyImpulse) ), UF_LUA_REGISTER_USERTYPE_DEFINE( applyRotation, UF_LUA_C_FUN(::binds::applyRotation) ), UF_LUA_REGISTER_USERTYPE_DEFINE( enableGravity, UF_LUA_C_FUN(::binds::enableGravity) ), - UF_LUA_REGISTER_USERTYPE_DEFINE( rayCast, UF_LUA_C_FUN(::binds::rayCast) ) - -// UF_LUA_REGISTER_USERTYPE_DEFINE( setImpulse, UF_LUA_C_FUN(uf::physics::impl::setImpulse) ), -// UF_LUA_REGISTER_USERTYPE_DEFINE( activateCollision, UF_LUA_C_FUN(uf::physics::impl::activateCollision) ), + UF_LUA_REGISTER_USERTYPE_DEFINE( rayCast, UF_LUA_C_FUN(::binds::rayCast) ), -// UF_LUA_REGISTER_USERTYPE_DEFINE( getMass, UF_LUA_C_FUN(uf::physics::impl::getMass) ), -// UF_LUA_REGISTER_USERTYPE_DEFINE( setMass, UF_LUA_C_FUN(uf::physics::impl::setMass) ), + UF_LUA_REGISTER_USERTYPE_DEFINE( getMass, UF_LUA_C_FUN(::binds::getMass) ), + UF_LUA_REGISTER_USERTYPE_DEFINE( setMass, UF_LUA_C_FUN(::binds::setMass) ) ) #endif \ No newline at end of file diff --git a/engine/src/ext/vulkan/rendermodes/deferred.cpp b/engine/src/ext/vulkan/rendermodes/deferred.cpp index 0593e4bc..fc99b7b7 100644 --- a/engine/src/ext/vulkan/rendermodes/deferred.cpp +++ b/engine/src/ext/vulkan/rendermodes/deferred.cpp @@ -28,6 +28,20 @@ namespace { const uf::stl::string DEFERRED_MODE = "compute"; ext::vulkan::Texture depthPyramid; uf::stl::vector depthPyramidViews; + + void cmdImageBarrier(VkCommandBuffer commandBuffer, VkImage image, VkAccessFlags srcAccess, VkAccessFlags dstAccess, VkImageLayout oldLayout, VkImageLayout newLayout) { + VkImageMemoryBarrier barrier{VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER}; + barrier.srcAccessMask = srcAccess; + barrier.dstAccessMask = dstAccess; + barrier.oldLayout = oldLayout; + barrier.newLayout = newLayout; + barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + barrier.image = image; + barrier.subresourceRange = { VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1 }; + + vkCmdPipelineBarrier(commandBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 0, NULL, 0, NULL, 1, &barrier); + }; } #include "./transition.inl" @@ -858,17 +872,18 @@ void ext::vulkan::DeferredRenderMode::createCommandBuffers( const uf::stl::vecto ::transitionAttachmentsTo( this, shader, commandBuffer ); // dispatch compute shader - VkMemoryBarrier memoryBarrier{VK_STRUCTURE_TYPE_MEMORY_BARRIER}; - memoryBarrier.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT; - memoryBarrier.dstAccessMask = VK_ACCESS_SHADER_READ_BIT; + auto& attachmentColor = this->getAttachment("color"); // color + auto& attachmentBright = this->getAttachment("bright"); // bloom + auto& attachmentScratch = this->getAttachment("scratch"); // pingpong - vkCmdPipelineBarrier( commandBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_FLAGS_NONE, 1, &memoryBarrier, 0, NULL, 0, NULL ); device->UF_CHECKPOINT_MARK( commandBuffer, pod::Checkpoint::GENERIC, "bloom[1]" ); blitter.record(commandBuffer, descriptor, 0, 1); - vkCmdPipelineBarrier( commandBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_FLAGS_NONE, 1, &memoryBarrier, 0, NULL, 0, NULL ); + cmdImageBarrier( commandBuffer, attachmentScratch.image, VK_ACCESS_SHADER_WRITE_BIT, VK_ACCESS_SHADER_READ_BIT, VK_IMAGE_LAYOUT_GENERAL, VK_IMAGE_LAYOUT_GENERAL ); + device->UF_CHECKPOINT_MARK( commandBuffer, pod::Checkpoint::GENERIC, "bloom[2]" ); blitter.record(commandBuffer, descriptor, 0, 2); - vkCmdPipelineBarrier( commandBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_FLAGS_NONE, 1, &memoryBarrier, 0, NULL, 0, NULL ); + cmdImageBarrier( commandBuffer, attachmentBright.image, VK_ACCESS_SHADER_WRITE_BIT, VK_ACCESS_SHADER_READ_BIT, VK_IMAGE_LAYOUT_GENERAL, VK_IMAGE_LAYOUT_GENERAL ); + device->UF_CHECKPOINT_MARK( commandBuffer, pod::Checkpoint::GENERIC, "bloom[3]" ); blitter.record(commandBuffer, descriptor, 0, 3); diff --git a/engine/src/utils/math/physics/bvh.inl b/engine/src/utils/math/physics/bvh.inl index 201f2254..d96e71d8 100644 --- a/engine/src/utils/math/physics/bvh.inl +++ b/engine/src/utils/math/physics/bvh.inl @@ -453,7 +453,7 @@ namespace { const auto& nodeA = bvh.nodes[nodeAID]; const auto& nodeB = bvh.nodes[nodeBID]; - if ( nodeA.isAsleep() || nodeB.isAsleep() || !::aabbOverlap( bvh.bounds[nodeAID], bvh.bounds[nodeBID] ) ) return; + if ( (nodeA.isAsleep() && nodeB.isAsleep()) || !::aabbOverlap( bvh.bounds[nodeAID], bvh.bounds[nodeBID] ) ) return; if ( nodeA.getCount() > 0 && nodeB.getCount() > 0 ) { for ( auto i = 0; i < nodeA.getCount(); ++i ) { @@ -483,7 +483,7 @@ namespace { const auto& nodeA = bvhA.nodes[nodeAID]; const auto& nodeB = bvhB.nodes[nodeBID]; - if ( nodeA.isAsleep() || nodeB.isAsleep() || !::aabbOverlap( bvhA.bounds[nodeAID], bvhB.bounds[nodeBID] ) ) return; + if ( (nodeA.isAsleep() && nodeB.isAsleep()) || !::aabbOverlap( bvhA.bounds[nodeAID], bvhB.bounds[nodeBID] ) ) return; if ( nodeA.getCount() > 0 && nodeB.getCount() > 0 ) { for ( auto i = 0; i < nodeA.getCount(); ++i ) { @@ -556,7 +556,7 @@ namespace { outIndices.reserve(::reserveCount); - thread_local uf::stl::stack stack; + static thread_local uf::stl::stack stack; //stack.clear(); // there is no stack.clear(), and the stack should already be cleared by the end of this function stack.push(0); @@ -603,7 +603,7 @@ namespace { if ( bvh.nodes.empty() ) return; outIndices.reserve(::reserveCount); - thread_local uf::stl::stack stack; + static thread_local uf::stl::stack stack; //stack.clear(); // there is no stack.clear(), and the stack should already be cleared by the end of this function stack.push(0); diff --git a/engine/src/utils/math/physics/impl.cpp b/engine/src/utils/math/physics/impl.cpp index 0e8e1982..ec8ccb77 100644 --- a/engine/src/utils/math/physics/impl.cpp +++ b/engine/src/utils/math/physics/impl.cpp @@ -3,7 +3,6 @@ #include #include #include -#include namespace { bool warmupSolver = true; // cache manifold data to warm up the solver @@ -33,7 +32,6 @@ namespace { float baumgarteCorrectionPercent = 0.4f; float baumgarteCorrectionSlop = 0.01f; - std::mutex cacheMutex; uf::stl::unordered_map manifoldsCache; uint32_t manifoldCacheLifetime = 6; // to-do: find a good value for this @@ -167,6 +165,8 @@ void uf::physics::impl::step( pod::World& world, float dt ) { uf::stl::vector islands; ::buildIslands( pairs, bodies, islands ); + if ( ::warmupSolver ) ::prepareManifoldCache( ::manifoldsCache, islands, bodies ); + // iterate islands #pragma omp parallel for schedule(dynamic) for ( auto& island : islands ) { @@ -190,7 +190,10 @@ void uf::physics::impl::step( pod::World& world, float dt ) { for ( auto& c : manifold.points ) c.normal = ::orientNormalToAB( a, b, c.normal ); } // retrieve accumulated impulses - if ( ::warmupSolver ) ::retrieveContacts( manifold, ::manifoldsCache[::makePairKey( a, b )] ); + if ( ::warmupSolver ) { + auto it = ::manifoldsCache.find( ::makePairKey( a, b ) ); + if ( it != ::manifoldsCache.end() ) ::retrieveContacts( manifold, it->second ); + } // merge similar contacts from a mesh to ensure continuity if ( a.collider.type == pod::ShapeType::MESH || b.collider.type == pod::ShapeType::MESH ) { ::mergeContacts( manifold ); @@ -221,11 +224,12 @@ void uf::physics::impl::step( pod::World& world, float dt ) { ::solvePositions( manifolds, dt ); // cache manifold positions if ( ::warmupSolver ) { - std::lock_guard lock(::cacheMutex); - ::storeManifolds( manifolds, ::manifoldsCache ); + ::updateManifoldCache( manifolds, ::manifoldsCache ); } } + if ( ::warmupSolver ) ::pruneManifoldCache( ::manifoldsCache ); + for ( auto* b : bodies ) { if ( b->isStatic ) continue; ::snapVelocity( *b, dt ); @@ -294,18 +298,8 @@ void uf::physics::impl::updateInertia( pod::PhysicsBody& body ) { */ pod::Vector3f dims = (body.collider.aabb.max - body.collider.aabb.min); pod::Vector3f dimsSq = dims * dims; - - float massFactor = body.mass / 12.0f; - body.inertiaTensor = { - massFactor * (dimsSq.y + dimsSq.z), - massFactor * (dimsSq.x + dimsSq.z), - massFactor * (dimsSq.x + dimsSq.y), - }; - - body.inertiaTensor.x = std::max(body.inertiaTensor.x, EPS(1e-6f)); - body.inertiaTensor.y = std::max(body.inertiaTensor.y, EPS(1e-6f)); - body.inertiaTensor.z = std::max(body.inertiaTensor.z, EPS(1e-6f)); - + body.inertiaTensor = pod::Vector3f{ dimsSq.y + dimsSq.z, dimsSq.x + dimsSq.z, dimsSq.x + dimsSq.y } * (body.mass / 12.0f); + body.inertiaTensor = uf::vector::max( body.inertiaTensor, { EPS(1.0e-6f), EPS(1.0e-6f), EPS(1.0e-6f) } ); body.inverseInertiaTensor = 1.0f / body.inertiaTensor; } break; case pod::ShapeType::SPHERE: { @@ -341,37 +335,42 @@ void uf::physics::impl::updateInertia( pod::PhysicsBody& body ) { totalVolume += extents.x * extents.y * extents.z; } - // accumulate inertia - for ( size_t i = 0; i < bvh.nodes.size(); ++i ) { - if ( bvh.nodes[i].getCount() == 0 ) continue; - const auto& box = bvh.bounds[i]; + if ( totalVolume < EPS(1.0e-6) ) { + body.inertiaTensor = { FLT_MAX, FLT_MAX, FLT_MAX }; + body.inverseInertiaTensor = { 0.0f, 0.0f, 0.0f }; + } else { + // accumulate inertia + for ( size_t i = 0; i < bvh.nodes.size(); ++i ) { + if ( bvh.nodes[i].getCount() == 0 ) continue; + const auto& box = bvh.bounds[i]; - auto extents = box.max - box.min; - float mass = body.mass * extents.x * extents.y * extents.z / totalVolume; + auto extents = box.max - box.min; + float mass = body.mass * extents.x * extents.y * extents.z / totalVolume; - // inertia tensor of a box about its center - float x2 = extents.x * extents.x; - float y2 = extents.y * extents.y; - float z2 = extents.z * extents.z; + // inertia tensor of a box about its center + float x2 = extents.x * extents.x; + float y2 = extents.y * extents.y; + float z2 = extents.z * extents.z; - pod::Matrix3f Ibox; - Ibox(0,0) = (1.0f/12.0f) * mass * (y2 + z2); - Ibox(1,1) = (1.0f/12.0f) * mass * (x2 + z2); - Ibox(2,2) = (1.0f/12.0f) * mass * (x2 + y2); + pod::Matrix3f Ibox; + Ibox(0,0) = (1.0f/12.0f) * mass * (y2 + z2); + Ibox(1,1) = (1.0f/12.0f) * mass * (x2 + z2); + Ibox(2,2) = (1.0f/12.0f) * mass * (x2 + y2); - // parallel axis theorem - pod::Vector3f center = (box.min + box.max) * 0.5f; - pod::Vector3f d = center; // relative to mesh COM (assume COM at origin for now) - float dist2 = uf::vector::magnitude( d ); + // parallel axis theorem + pod::Vector3f center = (box.min + box.max) * 0.5f; + pod::Vector3f d = center; // relative to mesh COM (assume COM at origin for now) + float dist2 = uf::vector::magnitude( d ); - pod::Matrix3f pat = uf::matrix::identityi() * (mass * dist2); - pat -= uf::matrix::outerProduct(d, d) * mass; + pod::Matrix3f pat = uf::matrix::identityi() * (mass * dist2); + pat -= uf::matrix::outerProduct(d, d) * mass; - inertia += Ibox + pat; + inertia += Ibox + pat; + } + + body.inertiaTensor = { inertia(0,0), inertia(1,1), inertia(2,2) }; + body.inverseInertiaTensor = 1.0f / body.inertiaTensor; } - - body.inertiaTensor = { inertia(0,0), inertia(1,1), inertia(2,2) }; - body.inverseInertiaTensor = 1.0f / body.inertiaTensor; } break; // to-do: add others default: { @@ -560,7 +559,7 @@ pod::RayQuery uf::physics::impl::rayCast( const pod::Ray& ray, const pod::World& auto& staticBvh = world.staticBvh; auto& bodies = world.bodies; - thread_local uf::stl::vector candidates; + static thread_local uf::stl::vector candidates; candidates.clear(); ::queryBVH( dynamicBvh, ray, candidates ); if ( ::useSplitBvhs ) ::queryBVH( staticBvh, ray, candidates ); diff --git a/engine/src/utils/math/physics/integration.inl b/engine/src/utils/math/physics/integration.inl index c042a8b8..9c31f009 100644 --- a/engine/src/utils/math/physics/integration.inl +++ b/engine/src/utils/math/physics/integration.inl @@ -175,24 +175,37 @@ namespace { } } - void storeManifolds( uf::stl::vector& manifolds, uf::stl::unordered_map& manifoldsCache ){ - // update cache - for ( auto& manifold : manifolds ) { - manifoldsCache[::makePairKey( *manifold.a, *manifold.b )] = manifold; - } + void prepareManifoldCache( uf::stl::unordered_map& cache, const uf::stl::vector& islands, const uf::stl::vector& bodies ) { + for ( const auto& island : islands ) { + for ( const auto& pair : island.pairs ) { + auto& a = *bodies[pair.first]; + auto& b = *bodies[pair.second]; - // prune if too old / empty - for ( auto itCache = manifoldsCache.begin(); itCache != manifoldsCache.end(); ) { + cache[ ::makePairKey( a, b ) ]; + } + } + } + + void updateManifoldCache( const uf::stl::vector& manifolds, uf::stl::unordered_map& cache ) { + for ( const auto& m : manifolds ) { + auto it = cache.find( ::makePairKey( *m.a, *m.b ) ); + if ( it == cache.end() ) continue; // assert + it->second = m; + } + } + + void pruneManifoldCache( uf::stl::unordered_map& cache ) { + for ( auto itCache = cache.begin(); itCache != cache.end(); ) { auto& manifold = itCache->second; - // prune manifolds that are X frames old + // prune points that are too old for ( auto it = manifold.points.begin(); it != manifold.points.end(); ) { if ( it->lifetime > ::manifoldCacheLifetime ) it = manifold.points.erase(it); else ++it; } // empty manifold, kill it - if ( manifold.points.empty() ) itCache = manifoldsCache.erase(itCache); + if ( manifold.points.empty() ) itCache = cache.erase(itCache); else ++itCache; } } diff --git a/engine/src/utils/math/physics/mesh.inl b/engine/src/utils/math/physics/mesh.inl index ad01c02f..58b9ee00 100644 --- a/engine/src/utils/math/physics/mesh.inl +++ b/engine/src/utils/math/physics/mesh.inl @@ -24,7 +24,7 @@ namespace { // transform to local space for BVH query auto bounds = ::transformAabbToLocal( aabb.bounds, ::getTransform( mesh ) ); - thread_local uf::stl::vector candidates; + static thread_local uf::stl::vector candidates; candidates.clear(); ::queryBVH( bvh, bounds, candidates ); @@ -48,7 +48,7 @@ namespace { // transform to local space for BVH query auto bounds = ::transformAabbToLocal( sphere.bounds, ::getTransform( mesh ) ); - thread_local uf::stl::vector candidates; + static thread_local uf::stl::vector candidates; candidates.clear(); ::queryBVH( bvh, bounds, candidates ); @@ -74,7 +74,7 @@ namespace { // transform to local space for BVH query auto bounds = ::transformAabbToLocal( plane.bounds, ::getTransform( mesh ) ); - thread_local uf::stl::vector candidates; + static thread_local uf::stl::vector candidates; candidates.clear(); ::queryBVH( bvh, bounds, candidates ); @@ -99,7 +99,7 @@ namespace { // transform to local space for BVH query auto bounds = ::transformAabbToLocal( capsule.bounds, ::getTransform( mesh ) ); - thread_local uf::stl::vector candidates; + static thread_local uf::stl::vector candidates; candidates.clear(); ::queryBVH( bvh, bounds, candidates ); @@ -124,13 +124,13 @@ namespace { const auto& meshB = *b.collider.mesh.mesh; // compute overlaps between one BVH and another BVH - thread_local pod::BVH::pairs_t pairs; + static thread_local pod::BVH::pairs_t pairs; pairs.clear(); ::queryOverlaps( bvhA, bvhB, pairs ); bool hit = false; // do collision per triangle - for (auto [ idA, idB] : pairs ) { + for (auto [idA, idB] : pairs ) { auto tA = ::fetchTriangle( meshA, idA, a ); // transform triangles to world space auto tB = ::fetchTriangle( meshB, idB, b ); diff --git a/engine/src/utils/math/physics/solvers.inl b/engine/src/utils/math/physics/solvers.inl index bd588956..c8862a99 100644 --- a/engine/src/utils/math/physics/solvers.inl +++ b/engine/src/utils/math/physics/solvers.inl @@ -144,12 +144,12 @@ namespace { float penetrationBias = std::max(contact.penetration - ::baumgarteCorrectionSlop, 0.0f) * (::baumgarteCorrectionPercent / dt); penetrationBias = std::min(penetrationBias, 2.0f / dt); // clamp - float maxPenetrationRecovery = 2.0f; // Limit to 2 units per second + float maxPenetrationRecovery = 2.0f; // limit to 2 units per second if ( penetrationBias > maxPenetrationRecovery ) penetrationBias = maxPenetrationRecovery; float cDot = vRel + penetrationBias; - rhs[i] = (cDot < 0.0f) ? -cDot : 0.0f; // rHS is magnitude of correction needed + rhs[i] = (cDot < 0.0f) ? -cDot : 0.0f; // RHS is magnitude of correction needed lambda[i] = contact.accumulatedNormalImpulse; } diff --git a/engine/src/utils/math/physics/triangle.inl b/engine/src/utils/math/physics/triangle.inl index 9fddad07..e833e28c 100644 --- a/engine/src/utils/math/physics/triangle.inl +++ b/engine/src/utils/math/physics/triangle.inl @@ -303,6 +303,7 @@ namespace { float planeDist = uf::vector::dot(triNormal, v0); if ( uf::vector::dot(bestAxis, triNormal) < 0.0f ) bestAxis = -bestAxis; pod::Vector3f contact = boxCenter - bestAxis * (boxHalf.x * fabs(bestAxis.x) + boxHalf.y * fabs(bestAxis.y) + boxHalf.z * fabs(bestAxis.z)); + //pod::Vector3f contact = boxCenter - triNormal * planeDist; /*