-
Notifications
You must be signed in to change notification settings - Fork 2
/
hairTilesPass.perStrand.wgsl.ts
133 lines (111 loc) · 3.87 KB
/
hairTilesPass.perStrand.wgsl.ts
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
import { BUFFER_HAIR_DATA } from '../../scene/hair/hairDataBuffer.ts';
import { BUFFER_HAIR_POINTS_POSITIONS } from '../../scene/hair/hairPointsPositionsBuffer.ts';
import { BUFFER_HAIR_TANGENTS } from '../../scene/hair/hairTangentsBuffer.ts';
import { RenderUniformsBuffer } from '../renderUniformsBuffer.ts';
import * as SHADER_SNIPPETS from '../_shaderSnippets/shaderSnippets.wgls.ts';
import { SW_RASTERIZE_HAIR } from './shaderImpl/swRasterizeHair.wgsl.ts';
import { BUFFER_HAIR_TILES_RESULT } from './shared/hairTilesResultBuffer.ts';
import { BUFFER_HAIR_TILE_SEGMENTS } from './shared/hairTileSegmentsBuffer.ts';
import { SHADER_TILE_UTILS } from './shaderImpl/tileUtils.wgsl.ts';
import {
TILE_PASSES_BINDINGS,
TILE_PASSES_SHARED,
} from './shaderImpl/tilePassesShared.wgsl.ts';
/*
This is an alternative implementation. In practise it seems much slower/harder
to optimize than the per-segment. Leaving the code cause this fact seems weird.
TODO [IGNORE] remove this impl? It is actually quite slow
*/
export const SHADER_PARAMS = {
workgroupSizeX: 32,
bindings: TILE_PASSES_BINDINGS,
};
///////////////////////////
/// SHADER CODE
///////////////////////////
const c = SHADER_PARAMS;
const b = SHADER_PARAMS.bindings;
export const SHADER_CODE = () => /* wgsl */ `
${SHADER_SNIPPETS.GET_MVP_MAT}
${SHADER_SNIPPETS.GENERIC_UTILS}
${SW_RASTERIZE_HAIR}
${SHADER_TILE_UTILS}
${TILE_PASSES_SHARED}
${RenderUniformsBuffer.SHADER_SNIPPET(b.renderUniforms)}
${BUFFER_HAIR_DATA(b.hairData)}
${BUFFER_HAIR_POINTS_POSITIONS(b.hairPositions)}
${BUFFER_HAIR_TANGENTS(b.hairTangents)}
${BUFFER_HAIR_TILES_RESULT(b.tilesBuffer, 'read_write')}
${BUFFER_HAIR_TILE_SEGMENTS(b.tileSegmentsBuffer, 'read_write')}
@group(0) @binding(${b.depthTexture})
var _depthTexture: texture_depth_2d;
@compute
@workgroup_size(${c.workgroupSizeX}, 1, 1)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
) {
let viewportSize: vec2f = _uniforms.viewport.xy;
let viewportSizeU32: vec2u = vec2u(viewportSize);
let maxDrawnSegments: u32 = _uniforms.maxDrawnHairSegments;
let mvMatrix = _uniforms.modelViewMat;
let projMatrixInv = _uniforms.projMatrixInv;
let strandsCount: u32 = _hairData.strandsCount;
let pointsPerStrand: u32 = _hairData.pointsPerStrand;
let hairDepthBoundsVS = getHairDepthBoundsVS(mvMatrix);
let strandIdx = global_id.x;
if (strandIdx >= strandsCount) { return; }
// get rasterize data
let swHairRasterizeParams = SwHairRasterizeParams(
pointsPerStrand,
viewportSize,
_uniforms.fiberRadius
);
var v00: vec2f;
var v01: vec2f;
var v10: vec2f;
var v11: vec2f;
var depthsProj0: vec2f;
var depthsProj1: vec2f;
// strand root
swRasterizeHairPoint(
swHairRasterizeParams,
strandIdx, 0u,
&v00, &v01, &depthsProj0
);
for (var segmentIdx: u32 = 0u; segmentIdx < pointsPerStrand - 1; segmentIdx += 1u) {
// raster segment end
swRasterizeHairPoint(
swHairRasterizeParams,
strandIdx, segmentIdx + 1u,
&v10, &v11, &depthsProj1
);
let sw = SwRasterizedHair(
v00, v01, v10, v11,
vec4f(depthsProj0, depthsProj1)
);
// get segment bounds and convert to tiles
let bounds4f = getRasterizedHairBounds(sw, viewportSize);
let boundRectMin = bounds4f.xy;
let boundRectMax = bounds4f.zw;
let tileMinXY: vec2u = getHairTileXY_FromPx(vec2u(boundRectMin));
let tileMaxXY: vec2u = getHairTileXY_FromPx(vec2u(boundRectMax));
// for each affected tile
for (var tileY: u32 = tileMinXY.y; tileY <= tileMaxXY.y; tileY += 1u) {
for (var tileX: u32 = tileMinXY.x; tileX <= tileMaxXY.x; tileX += 1u) {
processTile(
sw,
viewportSizeU32,
projMatrixInv,
maxDrawnSegments,
hairDepthBoundsVS,
vec2u(tileX, tileY),
strandIdx, segmentIdx
);
}}
// move to next segment
v00 = v10;
v01 = v11;
depthsProj0 = depthsProj1;
}
}
`;