1//
2// Copyright 2019 The ANGLE Project. 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
7#include "common.h"
8
9constant bool kSourceBufferAligned[[function_constant(0)]];
10constant bool kSourceIndexIsU8[[function_constant(1)]];
11constant bool kSourceIndexIsU16[[function_constant(2)]];
12constant bool kSourceIndexIsU32[[function_constant(3)]];
13constant bool kSourceBufferUnaligned = !kSourceBufferAligned;
14constant bool kUseSourceBufferU8     = kSourceIndexIsU8 || kSourceBufferUnaligned;
15constant bool kUseSourceBufferU16    = kSourceIndexIsU16 && kSourceBufferAligned;
16constant bool kUseSourceBufferU32    = kSourceIndexIsU32 && kSourceBufferAligned;
17
18struct IndexConversionParams
19{
20    uint32_t srcOffset;  // offset in bytes
21    uint32_t indexCount;
22};
23
24#define ANGLE_IDX_CONVERSION_GUARD(IDX, OPTS) ANGLE_KERNEL_GUARD(IDX, OPTS.indexCount)
25
26inline ushort getIndexAligned(constant ushort *inputAligned, uint offset, uint idx)
27{
28    return inputAligned[offset / 2 + idx];
29}
30inline uint getIndexAligned(constant uint *inputAligned, uint offset, uint idx)
31{
32    return inputAligned[offset / 4 + idx];
33}
34inline uchar getIndexAligned(constant uchar *input, uint offset, uint idx)
35{
36    return input[offset + idx];
37}
38inline ushort getIndexUnalignedU16(constant uchar *input, uint offset, uint idx)
39{
40    ushort inputLo = input[offset + 2 * idx];
41    ushort inputHi = input[offset + 2 * idx + 1];
42    // Little endian conversion:
43    return inputLo | (inputHi << 8);
44}
45inline uint getIndexUnalignedU32(constant uchar *input, uint offset, uint idx)
46{
47    uint input0 = input[offset + 4 * idx];
48    uint input1 = input[offset + 4 * idx + 1];
49    uint input2 = input[offset + 4 * idx + 2];
50    uint input3 = input[offset + 4 * idx + 3];
51    // Little endian conversion:
52    return input0 | (input1 << 8) | (input2 << 16) | (input3 << 24);
53}
54
55kernel void convertIndexU8ToU16(uint idx[[thread_position_in_grid]],
56                                constant IndexConversionParams &options[[buffer(0)]],
57                                constant uchar *input[[buffer(1)]],
58                                device ushort *output[[buffer(2)]])
59{
60    ANGLE_IDX_CONVERSION_GUARD(idx, options);
61    output[idx] = getIndexAligned(input, options.srcOffset, idx);
62}
63
64kernel void convertIndexU16(
65    uint idx[[thread_position_in_grid]],
66    constant IndexConversionParams &options[[buffer(0)]],
67    constant uchar *input[[ buffer(1), function_constant(kSourceBufferUnaligned) ]],
68    constant ushort *inputAligned[[ buffer(1), function_constant(kSourceBufferAligned) ]],
69    device ushort *output[[buffer(2)]])
70{
71    ANGLE_IDX_CONVERSION_GUARD(idx, options);
72
73    ushort value;
74    if (kSourceBufferAligned)
75    {
76        value = getIndexAligned(inputAligned, options.srcOffset, idx);
77    }
78    else
79    {
80        value = getIndexUnalignedU16(input, options.srcOffset, idx);
81    }
82    output[idx] = value;
83}
84
85kernel void convertIndexU32(
86    uint idx[[thread_position_in_grid]],
87    constant IndexConversionParams &options[[buffer(0)]],
88    constant uchar *input[[ buffer(1), function_constant(kSourceBufferUnaligned) ]],
89    constant uint *inputAligned[[ buffer(1), function_constant(kSourceBufferAligned) ]],
90    device uint *output[[buffer(2)]])
91{
92    ANGLE_IDX_CONVERSION_GUARD(idx, options);
93
94    uint value;
95    if (kSourceBufferAligned)
96    {
97        value = getIndexAligned(inputAligned, options.srcOffset, idx);
98    }
99    else
100    {
101        value = getIndexUnalignedU32(input, options.srcOffset, idx);
102    }
103    output[idx] = value;
104}
105
106struct TriFanArrayParams
107{
108    uint firstVertex;
109    uint vertexCountFrom3rd;  // vertex count excluding the 1st & 2nd vertices.
110};
111kernel void genTriFanIndicesFromArray(uint idx[[thread_position_in_grid]],
112                                      constant TriFanArrayParams &options[[buffer(0)]],
113                                      device uint *output[[buffer(2)]])
114{
115    ANGLE_KERNEL_GUARD(idx, options.vertexCountFrom3rd);
116
117    uint vertexIdx = options.firstVertex + 2 + idx;
118
119    output[3 * idx]     = options.firstVertex;
120    output[3 * idx + 1] = vertexIdx - 1;
121    output[3 * idx + 2] = vertexIdx;
122}
123
124inline uint getIndexU32(uint offset,
125                        uint idx,
126                        constant uchar *inputU8[[function_constant(kUseSourceBufferU8)]],
127                        constant ushort *inputU16[[function_constant(kUseSourceBufferU16)]],
128                        constant uint *inputU32[[function_constant(kUseSourceBufferU32)]])
129{
130    if (kUseSourceBufferU8)
131    {
132        if (kSourceIndexIsU16)
133        {
134            return getIndexUnalignedU16(inputU8, offset, idx);
135        }
136        else if (kSourceIndexIsU32)
137        {
138            return getIndexUnalignedU32(inputU8, offset, idx);
139        }
140        return getIndexAligned(inputU8, offset, idx);
141    }
142    else if (kUseSourceBufferU16)
143    {
144        return getIndexAligned(inputU16, offset, idx);
145    }
146    else if (kUseSourceBufferU32)
147    {
148        return getIndexAligned(inputU32, offset, idx);
149    }
150    return 0;
151}
152
153// Generate triangle fan indices from an indices buffer. indexCount options indicates number
154// of indices starting from the 3rd.
155kernel void genTriFanIndicesFromElements(
156    uint idx[[thread_position_in_grid]],
157    constant IndexConversionParams &options[[buffer(0)]],
158    constant uchar *inputU8[[ buffer(1), function_constant(kUseSourceBufferU8) ]],
159    constant ushort *inputU16[[ buffer(1), function_constant(kUseSourceBufferU16) ]],
160    constant uint *inputU32[[ buffer(1), function_constant(kUseSourceBufferU32) ]],
161    device uint *output[[buffer(2)]])
162{
163    ANGLE_IDX_CONVERSION_GUARD(idx, options);
164
165    uint elemIdx = 2 + idx;
166
167    output[3 * idx]     = getIndexU32(options.srcOffset, 0, inputU8, inputU16, inputU32);
168    output[3 * idx + 1] = getIndexU32(options.srcOffset, elemIdx - 1, inputU8, inputU16, inputU32);
169    output[3 * idx + 2] = getIndexU32(options.srcOffset, elemIdx, inputU8, inputU16, inputU32);
170}