@@ -5,7 +5,6 @@ using namespace metal;
5
5
6
6
#include " include/vger.h"
7
7
#include " sdf.h"
8
- #include " commands.h"
9
8
#include " paint.h"
10
9
11
10
#define SQRT_2 1.414213562373095
@@ -157,211 +156,3 @@ fragment float4 vger_fragment(VertexOut in [[ stage_in ]],
157
156
return mix (float4 (color.rgb ,0.0 ), color, 1.0 -smoothstep (-fw/2 ,fw/2 ,d) );
158
157
159
158
}
160
-
161
- kernel void vger_tile_clear (device uint *tileLengths [[buffer(0 )]],
162
- uint2 gid [[thread_position_in_grid]]) {
163
-
164
- tileLengths[gid.y * MAX_TILES_WIDTH + gid.x ] = 0 ;
165
- }
166
-
167
- fragment float4 vger_tile_fragment (VertexOut in [[ stage_in ]],
168
- const device vgerPrim* prims,
169
- const device float2* cvs,
170
- const device vgerPaint* paints,
171
- device Tile *tiles,
172
- device uint *tileLengths [[ /* raster_order_group(0)*/ ]]) {
173
-
174
- device auto & prim = prims[in.primIndex ];
175
- uint x = (uint ) in.position .x ;
176
- uint y = MAX_TILES_WIDTH - (uint ) in.position .y - 1 ;
177
- uint tileIx = y * MAX_TILES_WIDTH + x;
178
-
179
- // Always accessing tileLengths seems to work around the compiler bug.
180
- uint length = tileLengths[tileIx];
181
-
182
- device Tile& tile = tiles[tileIx];
183
-
184
- switch (prim.type ) {
185
- case vgerRect:
186
- tile.append (vgerCmdRect{vgerOpRect, prim.cvs [0 ], prim.cvs [1 ], prim.radius }, length);
187
- break ;
188
- case vgerPathFill:
189
- for (int i=0 ; i<prim.count ; i++) {
190
- int j = prim.start + 3 *i;
191
- auto a = cvs[j];
192
- auto b = cvs[j+1 ];
193
- auto c = cvs[j+2 ];
194
-
195
- auto m = in.t .y - TILE_SIZE_PIXELS/2 ;
196
- if (a.y < m and b.y < m and c.y < m) {
197
- continue ;
198
- }
199
- m = in.t .y + TILE_SIZE_PIXELS/2 ;
200
- if (a.y > m and b.y > m and c.y > m) {
201
- continue ;
202
- }
203
-
204
- m = in.t .x - TILE_SIZE_PIXELS/2 ;
205
- if (a.x < m and b.x < m and c.x < m) {
206
- continue ;
207
- }
208
-
209
- tile.append (vgerCmdBezFillIndirect{vgerOpBezIndirect,j}, length);
210
- }
211
- break ;
212
- case vgerSegment:
213
- tile.append (vgerCmdSegment{vgerOpSegment, prim.cvs [0 ], prim.cvs [1 ], prim.width }, length);
214
- break ;
215
- case vgerCircle:
216
- tile.append (vgerCmdCircle{vgerOpCircle, prim.cvs [0 ], prim.radius }, length);
217
- break ;
218
- case vgerBezier:
219
- tile.append (vgerCmdBezStroke{vgerOpBezStroke, prim.cvs [0 ], prim.cvs [1 ], prim.cvs [2 ], prim.width }, length);
220
- default :
221
- break ;
222
- }
223
-
224
- tile.append (vgerCmdSolid{vgerOpSolid,
225
- pack_float_to_srgb_unorm4x8 (paints[prim.paint ].innerColor )},
226
- length);
227
-
228
- tileLengths[tileIx] = length;
229
-
230
- // This is just for debugging so we can see what was rendered
231
- // in the coarse rasterization.
232
- return float4 (in.position .x /32 , 0 , (MAX_TILES_WIDTH-in.position .y )/32 , 1 );
233
-
234
- }
235
-
236
- kernel void vger_tile_render (texture2d<half, access::write> outTexture [[texture(0 )]],
237
- const device Tile *tiles [[buffer(0 )]],
238
- device uint *tileLengths [[buffer(1 )]],
239
- device float2 *cvs,
240
- uint2 gid [[thread_position_in_grid]],
241
- uint2 tgid [[threadgroup_position_in_grid]]) {
242
-
243
- uint tileIx = tgid.y * MAX_TILES_WIDTH + tgid.x ;
244
- const device char *src = tiles[tileIx].commands ;
245
- const device char *end = src + tileLengths[tileIx];
246
- uint x = gid.x ;
247
- uint y = gid.y ;
248
- float2 xy = float2 (x, y);
249
-
250
- if (x >= outTexture.get_width () || y >= outTexture.get_height ()) {
251
- return ;
252
- }
253
-
254
- // Show overflows.
255
- if (tileLengths[tileIx] > 4000 ) {
256
- outTexture.write (half4 (1.0 , 0.0 , 1.0 , 1.0 ), uint2{gid.x , outTexture.get_height () - gid.y - 1 });
257
- return ;
258
- }
259
-
260
- half3 rgb = half3 (0.0 );
261
- float d = 1e9 ;
262
-
263
- while (src < end) {
264
- vgerOp op = *(device vgerOp*) src;
265
-
266
- if (op == vgerOpEnd) {
267
- break ;
268
- }
269
-
270
- switch (op) {
271
- case vgerOpSegment: {
272
- vgerCmdSegment cmd = *(device vgerCmdSegment*) src;
273
-
274
- d = sdSegment2 (xy, cmd.a , cmd.b , cmd.width );
275
-
276
- src += sizeof (vgerCmdSegment);
277
- break ;
278
- }
279
-
280
- case vgerOpRect: {
281
- vgerCmdRect cmd = *(device vgerCmdRect*) src;
282
- auto center = .5 *(cmd.a + cmd.b );
283
- auto size = cmd.b - cmd.a ;
284
- d = sdBox (xy - center, .5 *size, cmd.radius );
285
-
286
- src += sizeof (vgerCmdRect);
287
- break ;
288
- }
289
-
290
- case vgerOpLine: {
291
- vgerCmdLineFill cmd = *(device vgerCmdLineFill*) src;
292
-
293
- if (lineTest (xy, cmd.a , cmd.b )) {
294
- d = -d;
295
- }
296
-
297
- src += sizeof (vgerCmdLineFill);
298
- break ;
299
- }
300
-
301
- case vgerOpBez: {
302
- vgerCmdBezFill cmd = *(device vgerCmdBezFill*) src;
303
- d = copysign (min (abs (d), sdBezierApprox2 (xy, cmd.a , cmd.b , cmd.c )), d);
304
-
305
- if (lineTest (xy, cmd.a , cmd.c )) {
306
- d = -d;
307
- }
308
-
309
- if (bezierTest (xy, cmd.a , cmd.b , cmd.c )) {
310
- d = -d;
311
- }
312
-
313
- src += sizeof (vgerCmdBezFill);
314
- break ;
315
- }
316
-
317
- case vgerOpBezIndirect: {
318
- vgerCmdBezFillIndirect cmd = *(device vgerCmdBezFillIndirect*) src;
319
- auto a = cvs[cmd.index ];
320
- auto b = cvs[cmd.index +1 ];
321
- auto c = cvs[cmd.index +2 ];
322
- d = copysign (min (abs (d), sdBezierApprox2 (xy, a, b, c)), d);
323
-
324
- if (lineTest (xy, a, c)) {
325
- d = -d;
326
- }
327
-
328
- if (bezierTest (xy, a, b, c)) {
329
- d = -d;
330
- }
331
-
332
- src += sizeof (vgerCmdBezFillIndirect);
333
- break ;
334
- }
335
-
336
- case vgerOpFillTile: {
337
- vgerCmdSolid cmd = *(device vgerCmdSolid*) src;
338
- half4 c = unpack_unorm4x8_srgb_to_half (cmd.color );
339
- rgb = c.rgb ;
340
- d = 1e9 ;
341
- src += sizeof (vgerCmdSolid);
342
- break ;
343
- }
344
-
345
- case vgerOpSolid: {
346
- vgerCmdSolid cmd = *(device vgerCmdSolid*) src;
347
- half4 c = unpack_unorm4x8_srgb_to_half (cmd.color );
348
- rgb = mix (rgb, c.rgb , 1.0 -smoothstep (-.5 ,.5 ,d) );
349
- d = 1e9 ;
350
- src += sizeof (vgerCmdSolid);
351
- break ;
352
- }
353
-
354
- default :
355
- outTexture.write (half4 (1.0 , 0.0 , 1.0 , 1.0 ), gid);
356
- return ;
357
-
358
- }
359
- }
360
-
361
- // Linear to sRGB conversion. Note that if we had writable sRGB textures
362
- // we could let this be done in the write call.
363
- rgb = select (1.055 * pow (rgb, 1 /2.4 ) - 0.055 , 12.92 * rgb, rgb < 0.0031308 );
364
- half4 rgba = half4 (rgb, 1.0 );
365
- outTexture.write (rgba, uint2{gid.x , outTexture.get_height () - gid.y - 1 });
366
-
367
- }
0 commit comments