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}