• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1//
2// Copyright 2021 The ANGLE Project Authors. All rights reserved.
3// Use of this source code is governed by a BSD-style license that can be
4// found in the LICENSE file.
5//
6// rewrite_indices.metal:
7//    Contains utility methods for rewriting indices for provoking vertex usecases.
8//
9
10#include "common.h"
11#include "rewrite_indices_shared.h"
12using namespace metal;
13
14constant uint fixIndexBufferKey [[ function_constant(2000) ]];
15constant bool indexBufferIsUint16 = (((fixIndexBufferKey >> MtlFixIndexBufferKeyInShift) & MtlFixIndexBufferKeyTypeMask) == MtlFixIndexBufferKeyUint16);
16constant bool indexBufferIsUint32 = (((fixIndexBufferKey >> MtlFixIndexBufferKeyInShift) & MtlFixIndexBufferKeyTypeMask) == MtlFixIndexBufferKeyUint32);
17constant bool outIndexBufferIsUint16 = (((fixIndexBufferKey >> MtlFixIndexBufferKeyOutShift) & MtlFixIndexBufferKeyTypeMask) == MtlFixIndexBufferKeyUint16);
18constant bool outIndexBufferIsUint32 = (((fixIndexBufferKey >> MtlFixIndexBufferKeyOutShift) & MtlFixIndexBufferKeyTypeMask) == MtlFixIndexBufferKeyUint32);
19constant bool doPrimRestart = (fixIndexBufferKey & MtlFixIndexBufferKeyPrimRestart);
20constant uint fixIndexBufferMode = (fixIndexBufferKey >> MtlFixIndexBufferKeyModeShift) & MtlFixIndexBufferKeyModeMask;
21
22
23static inline uint readIdx(
24                           const device ushort *indexBufferUint16,
25                           const device uint   *indexBufferUint32,
26                           const uint restartIndex,
27                           const uint indexCount,
28                           uint idx,
29                           thread bool &foundRestart,
30                           thread uint &indexThatRestartedFirst
31                           )
32{
33    uint inIndex = idx;
34    if(inIndex < indexCount)
35    {
36        if(indexBufferIsUint16)
37        {
38            inIndex = indexBufferUint16[inIndex];
39        }
40        else if(indexBufferIsUint32)
41        {
42            inIndex = indexBufferUint32[inIndex];
43        }
44    }
45    else
46    {
47        foundRestart = true;
48        indexThatRestartedFirst = idx;
49    }
50    if(doPrimRestart && !foundRestart && inIndex == restartIndex)
51    {
52        foundRestart = true;
53        indexThatRestartedFirst = idx;
54    }
55    return inIndex;
56}
57
58static inline void outputPrimitive(
59                                   const device ushort *indexBufferUint16,
60                                   const device uint   *indexBufferUint32,
61                                   device ushort *outIndexBufferUint16,
62                                   device uint   *outIndexBufferUint32,
63                                   const uint restartIndex,
64                                   const uint indexCount,
65                                   thread uint &baseIndex,
66                                   uint onIndex,
67                                   thread uint &onOutIndex
68                                   )
69{
70    if(baseIndex > onIndex) return; // skipped indices while processing
71    bool foundRestart = false;
72    uint indexThatRestartedFirst = 0;
73#define READ_IDX(_idx) readIdx(indexBufferUint16, indexBufferUint32, restartIndex, indexCount, _idx, foundRestart, indexThatRestartedFirst)
74#define WRITE_IDX(_idx, _val) \
75({ \
76    if(outIndexBufferIsUint16) \
77    { \
78        outIndexBufferUint16[(_idx)] = _val; \
79    } \
80    if(outIndexBufferIsUint32) \
81    { \
82        outIndexBufferUint32[(_idx)] = _val; \
83    } \
84    _idx++; \
85})
86    switch(fixIndexBufferMode)
87    {
88        case MtlFixIndexBufferKeyPoints:
89        {
90            auto tmpIndex = READ_IDX(onIndex);
91            if(foundRestart)
92            {
93                baseIndex = indexThatRestartedFirst + 1;
94                return;
95            }
96
97            WRITE_IDX(onOutIndex, tmpIndex);
98        }
99        break;
100        case MtlFixIndexBufferKeyLines:
101        {
102            auto tmpIndex0 = READ_IDX(onIndex + 0);
103            auto tmpIndex1 = READ_IDX(onIndex + 1);
104            if(foundRestart)
105            {
106                baseIndex = indexThatRestartedFirst + 1;
107                return;
108            }
109            if((onIndex - baseIndex) & 1) return; // skip this index...
110
111            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
112            {
113                WRITE_IDX(onOutIndex, tmpIndex1);
114                WRITE_IDX(onOutIndex, tmpIndex0);
115            }
116            else
117            {
118                WRITE_IDX(onOutIndex, tmpIndex0);
119                WRITE_IDX(onOutIndex, tmpIndex1);
120            }
121        }
122        break;
123        case MtlFixIndexBufferKeyLineStrip:
124        {
125            auto tmpIndex0 = READ_IDX(onIndex + 0);
126            auto tmpIndex1 = READ_IDX(onIndex + 1);
127            if(foundRestart)
128            {
129                baseIndex = indexThatRestartedFirst + 1;
130                return;
131            }
132
133            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
134            {
135                WRITE_IDX(onOutIndex, tmpIndex1);
136                WRITE_IDX(onOutIndex, tmpIndex0);
137            }
138            else
139            {
140                WRITE_IDX(onOutIndex, tmpIndex0);
141                WRITE_IDX(onOutIndex, tmpIndex1);
142            }
143        }
144        break;
145        case MtlFixIndexBufferKeyTriangles:
146        {
147            auto tmpIndex0 = READ_IDX(onIndex + 0);
148            auto tmpIndex1 = READ_IDX(onIndex + 1);
149            auto tmpIndex2 = READ_IDX(onIndex + 2);
150            if(foundRestart)
151            {
152                baseIndex = indexThatRestartedFirst + 1;
153                return;
154            }
155            if(((onIndex - baseIndex) % 3) != 0) return; // skip this index...
156
157            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
158            {
159                WRITE_IDX(onOutIndex, tmpIndex2);
160                WRITE_IDX(onOutIndex, tmpIndex0);
161                WRITE_IDX(onOutIndex, tmpIndex1);
162            }
163            else
164            {
165                WRITE_IDX(onOutIndex, tmpIndex0);
166                WRITE_IDX(onOutIndex, tmpIndex1);
167                WRITE_IDX(onOutIndex, tmpIndex2);
168            }
169        }
170        break;
171        case MtlFixIndexBufferKeyTriangleStrip:
172        {
173            uint isOdd = ((onIndex - baseIndex) & 1); // fixes winding (but not provoking...)
174            auto tmpIndex0 = READ_IDX(onIndex + 0 + isOdd);
175            auto tmpIndex1 = READ_IDX(onIndex + 1 - isOdd);
176            auto tmpIndex2 = READ_IDX(onIndex + 2);
177            if(foundRestart)
178            {
179                baseIndex = indexThatRestartedFirst + 1;
180                return;
181            }
182
183            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
184            {
185                WRITE_IDX(onOutIndex, tmpIndex2); // 2 is always the provoking vertex .: do not need to do anything special with isOdd
186                WRITE_IDX(onOutIndex, tmpIndex0);
187                WRITE_IDX(onOutIndex, tmpIndex1);
188            }
189            else
190            {
191                // NOTE: this case is trivially supported in Metal
192                if(isOdd)
193                {
194                    WRITE_IDX(onOutIndex, tmpIndex1); // in the case of odd this is REALLY (onIndex + 0) // provoking vertex
195                    WRITE_IDX(onOutIndex, tmpIndex2);
196                    WRITE_IDX(onOutIndex, tmpIndex0);
197                }
198                else
199                {
200                    WRITE_IDX(onOutIndex, tmpIndex0); // in the case of even this is (onIndex + 0) // provoking vertex
201                    WRITE_IDX(onOutIndex, tmpIndex1);
202                    WRITE_IDX(onOutIndex, tmpIndex2);
203                }
204            }
205            // assert never worse that worst-case expansion
206            assert(onOutIndex <= (onIndex + 1) * 3);
207            assert(onOutIndex <= (indexCount - 2) * 3);
208        }
209        break;
210
211    }
212#undef READ_IDX
213#undef WRITE_IDX
214}
215
216kernel void fixIndexBuffer(
217                           const device ushort *indexBufferUint16 [[ buffer(0), function_constant(indexBufferIsUint16) ]],
218                           const device uint   *indexBufferUint32 [[ buffer(0), function_constant(indexBufferIsUint32) ]],
219                           device ushort *outIndexBufferUint16 [[ buffer(1), function_constant(outIndexBufferIsUint16) ]],
220                           device uint   *outIndexBufferUint32 [[ buffer(1), function_constant(outIndexBufferIsUint32) ]],
221                           constant uint &indexCount [[ buffer(2) ]],
222                           constant uint &primCount [[ buffer(3) ]],
223                           uint prim [[thread_position_in_grid]])
224{
225    constexpr uint restartIndex = 0xFFFFFFFF; // unused
226    uint baseIndex = 0;
227    uint onIndex = onIndex;
228    uint onOutIndex = onOutIndex;
229    if(prim < primCount)
230    {
231        switch(fixIndexBufferMode)
232        {
233            case MtlFixIndexBufferKeyPoints:
234                onIndex = prim;
235                onOutIndex = prim;
236                break;
237            case MtlFixIndexBufferKeyLines:
238                onIndex = prim * 2;
239                onOutIndex = prim * 2;
240                break;
241            case MtlFixIndexBufferKeyLineStrip:
242                onIndex = prim;
243                onOutIndex = prim * 2;
244                break;
245            case MtlFixIndexBufferKeyTriangles:
246                onIndex = prim * 3;
247                onOutIndex = prim * 3;
248                break;
249            case MtlFixIndexBufferKeyTriangleStrip:
250                onIndex = prim;
251                onOutIndex = prim * 3;
252                break;
253        }
254        outputPrimitive(indexBufferUint16, indexBufferUint32, outIndexBufferUint16, outIndexBufferUint32, restartIndex, indexCount, baseIndex, onIndex, onOutIndex);
255    }
256}
257
258
259
260static inline void generatePrimitive(
261                                   device ushort *outIndexBufferUint16,
262                                   device uint   *outIndexBufferUint32,
263                                   const uint firstVertex,
264                                   const uint indexCount,
265                                   thread uint &baseIndex,
266                                   uint onIndex,
267                                   uint primCount,
268                                   thread uint &onOutIndex
269                                   )
270{
271    if(baseIndex > onIndex) return; // skipped indices while processing
272#define WRITE_IDX(_idx, _val) \
273({ \
274    if(outIndexBufferIsUint16) \
275    { \
276        outIndexBufferUint16[(_idx)] = _val + firstVertex; \
277    } \
278    if(outIndexBufferIsUint32) \
279    { \
280        outIndexBufferUint32[(_idx)] = _val + firstVertex; \
281    } \
282    _idx++; \
283})
284    switch(fixIndexBufferMode)
285    {
286        case MtlFixIndexBufferKeyPoints:
287        {
288            WRITE_IDX(onOutIndex, onIndex);
289        }
290        break;
291        case MtlFixIndexBufferKeyLines:
292        {
293            auto tmpIndex0 = onIndex + 0;
294            auto tmpIndex1 = onIndex + 1;
295            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
296            {
297                WRITE_IDX(onOutIndex, tmpIndex1);
298                WRITE_IDX(onOutIndex, tmpIndex0);
299            }
300            else
301            {
302                WRITE_IDX(onOutIndex, tmpIndex0);
303                WRITE_IDX(onOutIndex, tmpIndex1);
304            }
305        }
306        break;
307        case MtlFixIndexBufferKeyLineLoop:
308        {
309            auto tmpIndex0 = onIndex + 0;
310            auto tmpIndex1 = (onIndex + 1) % primCount;
311            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
312            {
313                WRITE_IDX(onOutIndex, tmpIndex1);
314                WRITE_IDX(onOutIndex, tmpIndex0);
315            }
316            else
317            {
318                WRITE_IDX(onOutIndex, tmpIndex0);
319                WRITE_IDX(onOutIndex, tmpIndex1);
320            }
321        }
322        break;
323        case MtlFixIndexBufferKeyLineStrip:
324        {
325            auto tmpIndex0 = onIndex + 0;
326            auto tmpIndex1 = onIndex + 1;
327            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
328            {
329                WRITE_IDX(onOutIndex, tmpIndex1);
330                WRITE_IDX(onOutIndex, tmpIndex0);
331            }
332            else
333            {
334                WRITE_IDX(onOutIndex, tmpIndex0);
335                WRITE_IDX(onOutIndex, tmpIndex1);
336            }
337        }
338        break;
339        case MtlFixIndexBufferKeyTriangles:
340        {
341            auto tmpIndex0 = onIndex + 0;
342            auto tmpIndex1 = onIndex + 1;
343            auto tmpIndex2 = onIndex + 2;
344            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
345            {
346                WRITE_IDX(onOutIndex, tmpIndex2);
347                WRITE_IDX(onOutIndex, tmpIndex0);
348                WRITE_IDX(onOutIndex, tmpIndex1);
349            }
350            else
351            {
352                WRITE_IDX(onOutIndex, tmpIndex0);
353                WRITE_IDX(onOutIndex, tmpIndex1);
354                WRITE_IDX(onOutIndex, tmpIndex2);
355            }
356        }
357        break;
358        case MtlFixIndexBufferKeyTriangleStrip:
359        {
360            uint isOdd = ((onIndex - baseIndex) & 1); // fixes winding. provoking fixed later.
361            auto tmpIndex0 = onIndex + 0 + isOdd;
362            auto tmpIndex1 = onIndex + 1 - isOdd;
363            auto tmpIndex2 = onIndex + 2;
364            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
365            {
366                WRITE_IDX(onOutIndex, tmpIndex2); // 2 is always the provoking vertex .: do not need to do anything special with isOdd
367                WRITE_IDX(onOutIndex, tmpIndex0);
368                WRITE_IDX(onOutIndex, tmpIndex1);
369            }
370            else
371            {
372                if(isOdd)
373                {
374                    WRITE_IDX(onOutIndex, tmpIndex1); // in the case of odd this is REALLY (onIndex + 0) // provoking vertex
375                    WRITE_IDX(onOutIndex, tmpIndex2);
376                    WRITE_IDX(onOutIndex, tmpIndex0);
377                }
378                else
379                {
380                    WRITE_IDX(onOutIndex, tmpIndex0); // in the case of even this is (onIndex + 0) // provoking vertex
381                    WRITE_IDX(onOutIndex, tmpIndex1);
382                    WRITE_IDX(onOutIndex, tmpIndex2);
383                }
384            }
385            // assert never worse that worst-case expansion
386            assert(onOutIndex <= (onIndex + 1) * 3);
387            assert(onOutIndex <= (indexCount - 2) * 3);
388            break;
389        }
390        case MtlFixIndexBufferKeyTriangleFan:
391        {
392            auto tmpIndex0 = 0;
393            auto tmpIndex1 = onIndex + 1;
394            auto tmpIndex2 = onIndex + 2;
395            // Provoking Vertex for triangle fans does not use the pivot index for flat shading data.
396            if(fixIndexBufferKey & MtlFixIndexBufferKeyProvokingVertexLast)
397            {
398                WRITE_IDX(onOutIndex, tmpIndex2);
399                WRITE_IDX(onOutIndex, tmpIndex0);
400                WRITE_IDX(onOutIndex, tmpIndex1);
401            }
402            else
403            {
404                WRITE_IDX(onOutIndex, tmpIndex1);
405                WRITE_IDX(onOutIndex, tmpIndex2);
406                WRITE_IDX(onOutIndex, tmpIndex0);
407            }
408        }
409        break;
410
411    }
412#undef WRITE_IDX
413}
414
415
416
417kernel void genIndexBuffer(
418                           device ushort *outIndexBufferUint16 [[ buffer(1), function_constant(outIndexBufferIsUint16) ]],
419                           device uint   *outIndexBufferUint32 [[ buffer(1), function_constant(outIndexBufferIsUint32) ]],
420                           constant uint &indexCount [[ buffer(2) ]],
421                           constant uint &primCount [[ buffer(3) ]],
422                           constant uint &firstVertex [[ buffer(4) ]],
423                           uint prim [[thread_position_in_grid]])
424{
425    uint baseIndex = 0;
426    uint onIndex = onIndex;
427    uint onOutIndex = onOutIndex;
428    if(prim < primCount)
429    {
430        switch(fixIndexBufferMode)
431        {
432            case MtlFixIndexBufferKeyPoints:
433                onIndex = prim;
434                onOutIndex = prim;
435                break;
436            case MtlFixIndexBufferKeyLines:
437                onIndex = prim * 2;
438                onOutIndex = prim * 2;
439                break;
440            case MtlFixIndexBufferKeyLineStrip:
441                onIndex = prim;
442                onOutIndex = prim * 2;
443                break;
444            case MtlFixIndexBufferKeyLineLoop:
445                onIndex = prim;
446                onOutIndex = prim * 2;
447                break;
448            case MtlFixIndexBufferKeyTriangles:
449                onIndex = prim * 3;
450                onOutIndex = prim * 3;
451                break;
452            case MtlFixIndexBufferKeyTriangleStrip:
453                onIndex = prim;
454                onOutIndex = prim * 3;
455                break;
456            case MtlFixIndexBufferKeyTriangleFan:
457                onIndex = prim;
458                onOutIndex = prim * 3;
459                break;
460        }
461        generatePrimitive(outIndexBufferUint16, outIndexBufferUint32, firstVertex, indexCount, baseIndex, onIndex, primCount, onOutIndex);
462    }
463}