/* Copyright (c) Microsoft Corporation Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #include "geometry.h" #include "tessellator.h" #define LIBAGX_TESS_MIN_ISOLINE_DENSITY_TESSELLATION_FACTOR 1.0f #define LIBAGX_TESS_MAX_ISOLINE_DENSITY_TESSELLATION_FACTOR 64.0f typedef unsigned int FXP; // fixed point number enum { U = 0, // points on a tri patch V = 1, }; enum { Ueq0 = 0, // edges on a tri patch Veq0 = 1, Weq0 = 2, }; enum { Ueq1 = 2, // edges on a quad patch: Ueq0, Veq0, Ueq1, Veq1 Veq1 = 3, }; #define QUAD_AXES 2 #define QUAD_EDGES 4 #define TRI_EDGES 3 // The interior can just use a simpler stitch. typedef enum DIAGONALS { DIAGONALS_INSIDE_TO_OUTSIDE, DIAGONALS_INSIDE_TO_OUTSIDE_EXCEPT_MIDDLE, DIAGONALS_MIRRORED } DIAGONALS; typedef struct TESS_FACTOR_CONTEXT { FXP fxpInvNumSegmentsOnFloorTessFactor; FXP fxpInvNumSegmentsOnCeilTessFactor; FXP fxpHalfTessFactorFraction; int numHalfTessFactorPoints; int splitPointOnFloorHalfTessFactor; } TESS_FACTOR_CONTEXT; struct INDEX_PATCH_CONTEXT { int insidePointIndexDeltaToRealValue; int insidePointIndexBadValue; int insidePointIndexReplacementValue; int outsidePointIndexPatchBase; int outsidePointIndexDeltaToRealValue; int outsidePointIndexBadValue; int outsidePointIndexReplacementValue; }; struct INDEX_PATCH_CONTEXT2 { int baseIndexToInvert; int indexInversionEndPoint; int cornerCaseBadValue; int cornerCaseReplacementValue; }; struct CHWTessellator { enum libagx_tess_mode mode; uint index_bias; // array where we will store u/v's for the points we generate global struct libagx_tess_point *Point; // array where we will store index topology global void *Index; // A second index patch we have to do handles the leftover strip of quads in // the middle of an odd quad patch after finishing all the concentric rings. // This also handles the leftover strip of points in the middle of an even // quad patch, when stitching the row of triangles up the left side (V major // quad) or bottom (U major quad) of the inner ring bool bUsingPatchedIndices; bool bUsingPatchedIndices2; struct INDEX_PATCH_CONTEXT IndexPatchCtx; struct INDEX_PATCH_CONTEXT2 IndexPatchCtx2; }; #define FXP_INTEGER_BITS 15 #define FXP_FRACTION_BITS 16 #define FXP_FRACTION_MASK 0x0000ffff #define FXP_INTEGER_MASK 0x7fff0000 #define FXP_ONE (1 << FXP_FRACTION_BITS) #define FXP_ONE_THIRD 0x00005555 #define FXP_TWO_THIRDS 0x0000aaaa #define FXP_ONE_HALF 0x00008000 static global float * tess_factors(constant struct libagx_tess_args *p, uint patch) { return p->tcs_buffer + (patch * p->tcs_stride_el); } static inline uint libagx_heap_alloc(global struct agx_geometry_state *heap, uint size_B) { // TODO: drop align to 4 I think return atomic_fetch_add((volatile atomic_uint *)(&heap->heap_bottom), align(size_B, 8)); } /* * Generate an indexed draw for a patch with the computed number of indices. * This allocates heap memory for the index buffer, returning the allocated * memory. */ static global void * libagx_draw(constant struct libagx_tess_args *p, enum libagx_tess_mode mode, bool lines, uint patch, uint count) { if (mode == LIBAGX_TESS_MODE_COUNT) { p->counts[patch] = count; } if (mode == LIBAGX_TESS_MODE_WITH_COUNTS) { /* The index buffer is already allocated, get a pointer inside it. * p->counts has had an inclusive prefix sum hence the subtraction. */ uint offset_el = p->counts[sub_sat(patch, 1u)]; if (patch == 0) offset_el = 0; return &p->index_buffer[offset_el]; } return NULL; } static void libagx_draw_points(private struct CHWTessellator *ctx, constant struct libagx_tess_args *p, uint patch, uint count) { /* For points mode with a single draw, we need to generate a trivial index * buffer to stuff in the patch ID in the right place. */ global uint32_t *indices = libagx_draw(p, ctx->mode, false, patch, count); if (ctx->mode == LIBAGX_TESS_MODE_COUNT) return; for (int i = 0; i < count; ++i) { indices[i] = ctx->index_bias + i; } } static void libagx_draw_empty(constant struct libagx_tess_args *p, enum libagx_tess_mode mode, uint patch) { if (mode == LIBAGX_TESS_MODE_COUNT) { p->counts[patch] = 0; } } /* * Allocate heap memory for domain points for a patch. The allocation * is recorded in the coord_allocs[] array, which is in elements. */ static global struct libagx_tess_point * libagx_heap_alloc_points(constant struct libagx_tess_args *p, uint patch, uint count) { /* If we're recording statistics, increment now. The statistic is for * tessellation evaluation shader invocations, which is equal to the number * of domain points generated. */ if (p->statistic) { atomic_fetch_add((volatile atomic_uint *)(p->statistic), count); } uint32_t elsize_B = sizeof(struct libagx_tess_point); uint32_t alloc_B = libagx_heap_alloc(p->heap, elsize_B * count); uint32_t alloc_el = alloc_B / elsize_B; p->coord_allocs[patch] = alloc_el; return (global struct libagx_tess_point *)(((uintptr_t)p->heap->heap) + alloc_B); } // Microsoft D3D11 Fixed Function Tessellator Reference - May 7, 2012 // amar.patel@microsoft.com #define LIBAGX_TESS_MIN_ODD_TESSELLATION_FACTOR 1 #define LIBAGX_TESS_MAX_ODD_TESSELLATION_FACTOR 63 #define LIBAGX_TESS_MIN_EVEN_TESSELLATION_FACTOR 2 #define LIBAGX_TESS_MAX_EVEN_TESSELLATION_FACTOR 64 // 2^(-16), min positive fixed point fraction #define EPSILON 0.0000152587890625f #define MIN_ODD_TESSFACTOR_PLUS_HALF_EPSILON \ (LIBAGX_TESS_MIN_ODD_TESSELLATION_FACTOR + EPSILON / 2) static float clamp_factor(float factor, enum libagx_tess_partitioning partitioning, float maxf) { float lower = (partitioning == LIBAGX_TESS_PARTITIONING_FRACTIONAL_EVEN) ? LIBAGX_TESS_MIN_EVEN_TESSELLATION_FACTOR : LIBAGX_TESS_MIN_ODD_TESSELLATION_FACTOR; float upper = (partitioning == LIBAGX_TESS_PARTITIONING_FRACTIONAL_ODD) ? LIBAGX_TESS_MAX_ODD_TESSELLATION_FACTOR : LIBAGX_TESS_MAX_EVEN_TESSELLATION_FACTOR; // If any TessFactor will end up > 1 after floatToFixed conversion later, // then force the inside TessFactors to be > 1 so there is a picture frame. if (partitioning == LIBAGX_TESS_PARTITIONING_FRACTIONAL_ODD && maxf > MIN_ODD_TESSFACTOR_PLUS_HALF_EPSILON) { lower = LIBAGX_TESS_MIN_ODD_TESSELLATION_FACTOR + EPSILON; } factor = clamp(factor, lower, upper); if (partitioning == LIBAGX_TESS_PARTITIONING_INTEGER) { factor = ceil(factor); } return factor; } static FXP floatToFixed(const float input) { return mad(input, FXP_ONE, 0.5f); } static bool isOdd(const float input) { return ((int)input) & 1; } static FXP fxpCeil(const FXP input) { if (input & FXP_FRACTION_MASK) { return (input & FXP_INTEGER_MASK) + FXP_ONE; } return input; } static FXP fxpFloor(const FXP input) { return (input & FXP_INTEGER_MASK); } static int PatchIndexValue(private struct CHWTessellator *ctx, int index) { if (ctx->bUsingPatchedIndices) { // assumed remapped outide indices are > remapped inside vertices if (index >= ctx->IndexPatchCtx.outsidePointIndexPatchBase) { if (index == ctx->IndexPatchCtx.outsidePointIndexBadValue) return ctx->IndexPatchCtx.outsidePointIndexReplacementValue; else return index + ctx->IndexPatchCtx.outsidePointIndexDeltaToRealValue; } else { if (index == ctx->IndexPatchCtx.insidePointIndexBadValue) return ctx->IndexPatchCtx.insidePointIndexReplacementValue; else return index + ctx->IndexPatchCtx.insidePointIndexDeltaToRealValue; } } else if (ctx->bUsingPatchedIndices2) { if (index == ctx->IndexPatchCtx2.cornerCaseBadValue) { return ctx->IndexPatchCtx2.cornerCaseReplacementValue; } else if (index >= ctx->IndexPatchCtx2.baseIndexToInvert) { return ctx->IndexPatchCtx2.indexInversionEndPoint - index; } } return index; } static void DefinePoint(global struct libagx_tess_point *out, FXP fxpU, FXP fxpV) { out->u = fxpU; out->v = fxpV; } static void DefineIndex(private struct CHWTessellator *ctx, int index, int indexStorageOffset) { global uint32_t *indices = (global uint32_t *)ctx->Index; indices[indexStorageOffset] = ctx->index_bias + PatchIndexValue(ctx, index); } static void DefineTriangle(private struct CHWTessellator *ctx, int index0, int index1, int index2, int indexStorageBaseOffset) { index0 = PatchIndexValue(ctx, index0); index1 = PatchIndexValue(ctx, index1); index2 = PatchIndexValue(ctx, index2); vstore3(ctx->index_bias + (uint3)(index0, index1, index2), 0, (global uint *)ctx->Index + indexStorageBaseOffset); } static uint32_t RemoveMSB(uint32_t val) { uint32_t bit = val ? (1 << (31 - clz(val))) : 0; return val & ~bit; } static int NumPointsForTessFactor(bool odd, FXP fxpTessFactor) { // Add epsilon for rounding and add 1 for odd FXP f = fxpTessFactor + (odd ? (FXP_ONE + 1) : 1); int r = fxpCeil(f / 2) >> (FXP_FRACTION_BITS - 1); return odd ? r : r + 1; } static void ComputeTessFactorCtx(bool odd, FXP fxpTessFactor, private TESS_FACTOR_CONTEXT *TessFactorCtx) { // fxpHalfTessFactor == 1/2 if TessFactor is 1, // but we're pretending we are even. FXP fxpHalfTessFactor = (fxpTessFactor + 1 /*round*/) / 2; if (odd || (fxpHalfTessFactor == FXP_ONE_HALF)) { fxpHalfTessFactor += FXP_ONE_HALF; } FXP fxpFloorHalfTessFactor = fxpFloor(fxpHalfTessFactor); FXP fxpCeilHalfTessFactor = fxpCeil(fxpHalfTessFactor); TessFactorCtx->fxpHalfTessFactorFraction = fxpHalfTessFactor - fxpFloorHalfTessFactor; TessFactorCtx->numHalfTessFactorPoints = (fxpCeilHalfTessFactor >> FXP_FRACTION_BITS); // for EVEN, we don't include the point always // fixed at the midpoint of the TessFactor if (fxpCeilHalfTessFactor == fxpFloorHalfTessFactor) { TessFactorCtx->splitPointOnFloorHalfTessFactor = /*pick value to cause this to be ignored*/ TessFactorCtx->numHalfTessFactorPoints + 1; } else if (odd) { if (fxpFloorHalfTessFactor == FXP_ONE) { TessFactorCtx->splitPointOnFloorHalfTessFactor = 0; } else { TessFactorCtx->splitPointOnFloorHalfTessFactor = (RemoveMSB((fxpFloorHalfTessFactor >> FXP_FRACTION_BITS) - 1) << 1) + 1; } } else { TessFactorCtx->splitPointOnFloorHalfTessFactor = (RemoveMSB(fxpFloorHalfTessFactor >> FXP_FRACTION_BITS) << 1) + 1; } int numFloorSegments = (fxpFloorHalfTessFactor * 2) >> FXP_FRACTION_BITS; int numCeilSegments = (fxpCeilHalfTessFactor * 2) >> FXP_FRACTION_BITS; if (odd) { numFloorSegments -= 1; numCeilSegments -= 1; } TessFactorCtx->fxpInvNumSegmentsOnFloorTessFactor = floatToFixed(1.0f / (float)numFloorSegments); TessFactorCtx->fxpInvNumSegmentsOnCeilTessFactor = floatToFixed(1.0f / (float)numCeilSegments); } static FXP PlacePointIn1D(private const TESS_FACTOR_CONTEXT *TessFactorCtx, bool odd, int point) { bool bFlip = point >= TessFactorCtx->numHalfTessFactorPoints; if (bFlip) { point = (TessFactorCtx->numHalfTessFactorPoints << 1) - point - odd; } // special casing middle since 16 bit fixed math below can't reproduce 0.5 exactly if (point == TessFactorCtx->numHalfTessFactorPoints) return FXP_ONE_HALF; unsigned int indexOnCeilHalfTessFactor = point; unsigned int indexOnFloorHalfTessFactor = indexOnCeilHalfTessFactor; if (point > TessFactorCtx->splitPointOnFloorHalfTessFactor) { indexOnFloorHalfTessFactor -= 1; } // For the fixed point multiplies below, we know the results are <= 16 bits // because the locations on the halfTessFactor are <= half the number of // segments for the total TessFactor. So a number divided by a number that // is at least twice as big will give a result no bigger than 0.5 (which in // fixed point is 16 bits in our case) FXP fxpLocationOnFloorHalfTessFactor = indexOnFloorHalfTessFactor * TessFactorCtx->fxpInvNumSegmentsOnFloorTessFactor; FXP fxpLocationOnCeilHalfTessFactor = indexOnCeilHalfTessFactor * TessFactorCtx->fxpInvNumSegmentsOnCeilTessFactor; // Since we know the numbers calculated above are <= fixed point 0.5, and the // equation below is just lerping between two values <= fixed point 0.5 // (0x00008000), then we know that the final result before shifting by 16 bits // is no larger than 0x80000000. Once we shift that down by 16, we get the // result of lerping 2 numbers <= 0.5, which is obviously at most 0.5 // (0x00008000) FXP fxpLocation = fxpLocationOnFloorHalfTessFactor * (FXP_ONE - TessFactorCtx->fxpHalfTessFactorFraction) + fxpLocationOnCeilHalfTessFactor * (TessFactorCtx->fxpHalfTessFactorFraction); fxpLocation = (fxpLocation + FXP_ONE_HALF /*round*/) >> FXP_FRACTION_BITS; // get back to n.16 if (bFlip) { fxpLocation = FXP_ONE - fxpLocation; } return fxpLocation; } static void StitchRegular(private struct CHWTessellator *ctx, bool bTrapezoid, DIAGONALS diagonals, int baseIndexOffset, int numInsideEdgePoints, int insideEdgePointBaseOffset, int outsideEdgePointBaseOffset) { int insidePoint = insideEdgePointBaseOffset; int outsidePoint = outsideEdgePointBaseOffset; if (bTrapezoid) { DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; outsidePoint++; } int p; switch (diagonals) { case DIAGONALS_INSIDE_TO_OUTSIDE: // Diagonals pointing from inside edge forward towards outside edge for (p = 0; p < numInsideEdgePoints - 1; p++) { DefineTriangle(ctx, insidePoint, outsidePoint, outsidePoint + 1, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, insidePoint, outsidePoint + 1, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; } break; case DIAGONALS_INSIDE_TO_OUTSIDE_EXCEPT_MIDDLE: // Assumes ODD tessellation // Diagonals pointing from outside edge forward towards inside edge // First half for (p = 0; p < numInsideEdgePoints / 2 - 1; p++) { DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, insidePoint, outsidePoint + 1, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; } // Middle DefineTriangle(ctx, outsidePoint, insidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; p += 2; // Second half for (; p < numInsideEdgePoints; p++) { DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, insidePoint, outsidePoint + 1, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; } break; case DIAGONALS_MIRRORED: // First half, diagonals pointing from outside of outside edge to inside of // inside edge for (p = 0; p < numInsideEdgePoints / 2; p++) { DefineTriangle(ctx, outsidePoint, insidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; } // Second half, diagonals pointing from inside of inside edge to outside of // outside edge for (; p < numInsideEdgePoints - 1; p++) { DefineTriangle(ctx, insidePoint, outsidePoint, outsidePoint + 1, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, insidePoint, outsidePoint + 1, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; } break; } if (bTrapezoid) { DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; } } // loop_start and loop_end give optimal loop bounds for // the stitching algorithm further below, for any given halfTssFactor. There // is probably a better way to encode this... // // Return the FIRST entry in finalPointPositionTable awhich is less than // halfTessFactor, except entry 0 and 1 which are set up to skip the loop. static int loop_start(int N) { if (N < 2) return 1; else if (N == 2) return 17; else if (N < 5) return 9; else if (N < 9) return 5; else if (N < 17) return 3; else return 2; } // Return the LAST entry in finalPointPositionTable[] which is less than // halfTessFactor, except entry 0 and 1 which are set up to skip the loop. static int loop_end(int N) { if (N < 2) return 0; else if (N < 4) return 17; else if (N < 8) return 25; else if (N < 16) return 29; else if (N < 32) return 31; else return 32; } // Tables to assist in the stitching of 2 rows of points having arbitrary // TessFactors. The stitching order is governed by Ruler Function vertex // split ordering (see external documentation). // // The contents of the finalPointPositionTable are where vertex i [0..33] // ends up on the half-edge at the max tessellation amount given // ruler-function split order. Recall the other half of an edge is mirrored, // so we only need to deal with one half. This table is used to decide when // to advance a point on the interior or exterior. It supports odd TessFactor // up to 65 and even TessFactor up to 64. /* TODO: Is this actually faster than a LUT? */ static uint32_t finalPointPositionTable(uint32_t x) { if (x == 0) return 0; if (x == 1) return 0x20; uint32_t shift; if ((x & 1) == 0) { shift = 1; } else if ((x & 3) == 3) { shift = 2; } else if ((x & 7) == 5) { shift = 3; } else if (x != 17) { shift = 4; } else { shift = 5; } // SWAR vectorized right-shift of (0x20, x) // We're calculating `min(0xf, 0x20 >> shift) + (x >> shift)`. uint32_t items_to_shift = x | (0x20 << 16); uint32_t shifted = items_to_shift >> shift; uint32_t bias = min(0xfu, shifted >> 16); return bias + (shifted & 0xffff); } static void StitchTransition(private struct CHWTessellator *ctx, int baseIndexOffset, int insideEdgePointBaseOffset, int insideNumHalfTessFactorPoints, bool insideEdgeTessFactorOdd, int outsideEdgePointBaseOffset, int outsideNumHalfTessFactorPoints, bool outsideTessFactorOdd) { if (insideEdgeTessFactorOdd) { insideNumHalfTessFactorPoints -= 1; } if (outsideTessFactorOdd) { outsideNumHalfTessFactorPoints -= 1; } // Walk first half int outsidePoint = outsideEdgePointBaseOffset; int insidePoint = insideEdgePointBaseOffset; // iStart,iEnd are a small optimization so the loop below doesn't have to go // from 0 up to 31 int iStart = min(loop_start(insideNumHalfTessFactorPoints), loop_start(outsideNumHalfTessFactorPoints)); int iEnd = loop_end( max(insideNumHalfTessFactorPoints, outsideNumHalfTessFactorPoints)); // since we don't start the loop at 0 below, we need a special case. if (0 < outsideNumHalfTessFactorPoints) { // Advance outside DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; outsidePoint++; } for (int i = iStart; i <= iEnd; i++) { int bound = finalPointPositionTable(i); if (bound < insideNumHalfTessFactorPoints) { // Advance inside DefineTriangle(ctx, insidePoint, outsidePoint, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; } if (bound < outsideNumHalfTessFactorPoints) { // Advance outside DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; outsidePoint++; } } if ((insideEdgeTessFactorOdd != outsideTessFactorOdd) || insideEdgeTessFactorOdd) { if (insideEdgeTessFactorOdd == outsideTessFactorOdd) { // Quad in the middle DefineTriangle(ctx, insidePoint, outsidePoint, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; DefineTriangle(ctx, insidePoint + 1, outsidePoint, outsidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; outsidePoint++; } else if (!insideEdgeTessFactorOdd) { // Triangle pointing inside DefineTriangle(ctx, insidePoint, outsidePoint, outsidePoint + 1, baseIndexOffset); baseIndexOffset += 3; outsidePoint++; } else { // Triangle pointing outside DefineTriangle(ctx, insidePoint, outsidePoint, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; } } // Walk second half. for (int i = iEnd; i >= iStart; i--) { int bound = finalPointPositionTable(i); if (bound < outsideNumHalfTessFactorPoints) { // Advance outside DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; outsidePoint++; } if (bound < insideNumHalfTessFactorPoints) { // Advance inside DefineTriangle(ctx, insidePoint, outsidePoint, insidePoint + 1, baseIndexOffset); baseIndexOffset += 3; insidePoint++; } } // Below case is not needed if we didn't optimize loop above and made it run // from 31 down to 0. if (0 < outsideNumHalfTessFactorPoints) { DefineTriangle(ctx, outsidePoint, outsidePoint + 1, insidePoint, baseIndexOffset); baseIndexOffset += 3; outsidePoint++; } } KERNEL(64) libagx_tess_isoline(constant struct libagx_tess_args *p, enum libagx_tess_mode mode__2) { enum libagx_tess_mode mode = mode__2; uint patch = get_global_id(0); enum libagx_tess_partitioning partitioning = p->partitioning; bool lineDensityOdd; bool lineDetailOdd; TESS_FACTOR_CONTEXT lineDensityTessFactorCtx; TESS_FACTOR_CONTEXT lineDetailTessFactorCtx; global float *factors = tess_factors(p, patch); float TessFactor_V_LineDensity = factors[0]; float TessFactor_U_LineDetail = factors[1]; // Is the patch culled? NaN will pass. if (!(TessFactor_V_LineDensity > 0) || !(TessFactor_U_LineDetail > 0)) { libagx_draw_empty(p, mode, patch); return; } // Clamp edge TessFactors TessFactor_V_LineDensity = clamp(TessFactor_V_LineDensity, LIBAGX_TESS_MIN_ISOLINE_DENSITY_TESSELLATION_FACTOR, LIBAGX_TESS_MAX_ISOLINE_DENSITY_TESSELLATION_FACTOR); TessFactor_U_LineDetail = clamp_factor(TessFactor_U_LineDetail, partitioning, 0); // Process tessFactors if (partitioning == LIBAGX_TESS_PARTITIONING_INTEGER) { lineDetailOdd = isOdd(TessFactor_U_LineDetail); } else { lineDetailOdd = (partitioning == LIBAGX_TESS_PARTITIONING_FRACTIONAL_ODD); } FXP fxpTessFactor_U_LineDetail = floatToFixed(TessFactor_U_LineDetail); ComputeTessFactorCtx(lineDetailOdd, fxpTessFactor_U_LineDetail, &lineDetailTessFactorCtx); int numPointsPerLine = NumPointsForTessFactor(lineDetailOdd, fxpTessFactor_U_LineDetail); TessFactor_V_LineDensity = ceil(TessFactor_V_LineDensity); lineDensityOdd = isOdd(TessFactor_V_LineDensity); FXP fxpTessFactor_V_LineDensity = floatToFixed(TessFactor_V_LineDensity); ComputeTessFactorCtx(lineDensityOdd, fxpTessFactor_V_LineDensity, &lineDensityTessFactorCtx); // don't draw last line at V == 1. int numLines = NumPointsForTessFactor(lineDensityOdd, fxpTessFactor_V_LineDensity) - 1; /* Points */ uint num_points = numPointsPerLine * numLines; if (mode != LIBAGX_TESS_MODE_COUNT) { global struct libagx_tess_point *points = libagx_heap_alloc_points(p, patch, num_points); for (int line = 0, pointOffset = 0; line < numLines; line++) { FXP fxpV = PlacePointIn1D(&lineDensityTessFactorCtx, lineDensityOdd, line); for (int point = 0; point < numPointsPerLine; point++) { FXP fxpU = PlacePointIn1D(&lineDetailTessFactorCtx, lineDetailOdd, point); DefinePoint(&points[pointOffset++], fxpU, fxpV); } } } struct CHWTessellator ctx = { .mode = mode, .index_bias = patch * LIBAGX_TES_PATCH_ID_STRIDE, }; /* Connectivity */ if (!p->points_mode) { uint num_indices = numLines * (numPointsPerLine - 1) * 2; ctx.Index = libagx_draw(p, mode, true, patch, num_indices); if (mode == LIBAGX_TESS_MODE_COUNT) return; for (int line = 0, pointOffset = 0, indexOffset = 0; line < numLines; line++) { pointOffset++; for (int point = 1; point < numPointsPerLine; point++) { DefineIndex(&ctx, pointOffset - 1, indexOffset++); DefineIndex(&ctx, pointOffset, indexOffset++); pointOffset++; } } } else { libagx_draw_points(&ctx, p, patch, num_points); } } KERNEL(64) libagx_tess_tri(constant struct libagx_tess_args *p, enum libagx_tess_mode mode__2) { enum libagx_tess_mode mode = mode__2; uint patch = get_global_id(0); enum libagx_tess_partitioning partitioning = p->partitioning; global float *factors = tess_factors(p, patch); float tessFactor_Ueq0 = factors[0]; float tessFactor_Veq0 = factors[1]; float tessFactor_Weq0 = factors[2]; float insideTessFactor_f = factors[4]; struct CHWTessellator ctx = { .mode = mode, .index_bias = patch * LIBAGX_TES_PATCH_ID_STRIDE, }; // Is the patch culled? NaN will pass. if (!(tessFactor_Ueq0 > 0) || !(tessFactor_Veq0 > 0) || !(tessFactor_Weq0 > 0)) { libagx_draw_empty(p, mode, patch); return; } FXP outsideTessFactor[TRI_EDGES]; FXP insideTessFactor; bool outsideTessFactorOdd[TRI_EDGES]; bool insideTessFactorOdd; TESS_FACTOR_CONTEXT outsideTessFactorCtx[TRI_EDGES]; TESS_FACTOR_CONTEXT insideTessFactorCtx; // Stuff below is just specific to the traversal order // this code happens to use to generate points/lines int numPointsForOutsideEdge[TRI_EDGES]; int numPointsForInsideTessFactor; int insideEdgePointBaseOffset; // Clamp TessFactors tessFactor_Ueq0 = clamp_factor(tessFactor_Ueq0, partitioning, 0); tessFactor_Veq0 = clamp_factor(tessFactor_Veq0, partitioning, 0); tessFactor_Weq0 = clamp_factor(tessFactor_Weq0, partitioning, 0); float maxf = max(max(tessFactor_Ueq0, tessFactor_Veq0), tessFactor_Weq0); insideTessFactor_f = clamp_factor(insideTessFactor_f, partitioning, maxf); // Note the above clamps map NaN to the lower bound // Process tessFactors float outsideTessFactor_f[TRI_EDGES] = {tessFactor_Ueq0, tessFactor_Veq0, tessFactor_Weq0}; if (partitioning == LIBAGX_TESS_PARTITIONING_INTEGER) { for (int edge = 0; edge < TRI_EDGES; edge++) { outsideTessFactorOdd[edge] = isOdd(outsideTessFactor_f[edge]); } insideTessFactorOdd = isOdd(insideTessFactor_f) && (1.0f != insideTessFactor_f); } else { bool odd = (partitioning == LIBAGX_TESS_PARTITIONING_FRACTIONAL_ODD); for (int edge = 0; edge < TRI_EDGES; edge++) { outsideTessFactorOdd[edge] = odd; } insideTessFactorOdd = odd; } // Save fixed point TessFactors for (int edge = 0; edge < TRI_EDGES; edge++) { outsideTessFactor[edge] = floatToFixed(outsideTessFactor_f[edge]); } insideTessFactor = floatToFixed(insideTessFactor_f); if (partitioning != LIBAGX_TESS_PARTITIONING_FRACTIONAL_EVEN) { // Special case if all TessFactors are 1 if ((FXP_ONE == insideTessFactor) && (FXP_ONE == outsideTessFactor[Ueq0]) && (FXP_ONE == outsideTessFactor[Veq0]) && (FXP_ONE == outsideTessFactor[Weq0])) { /* Just do minimum tess factor */ if (mode == LIBAGX_TESS_MODE_COUNT) { p->counts[patch] = 3; return; } global struct libagx_tess_point *points = libagx_heap_alloc_points(p, patch, 3); DefinePoint(&points[0], 0, FXP_ONE); // V=1 (beginning of Ueq0 edge VW) DefinePoint(&points[1], 0, 0); // W=1 (beginning of Veq0 edge WU) DefinePoint(&points[2], FXP_ONE, 0); // U=1 (beginning of Weq0 edge UV) if (!p->points_mode) { ctx.Index = libagx_draw(p, mode, false, patch, 3); DefineTriangle(&ctx, 0, 1, 2, /*indexStorageBaseOffset*/ 0); } else { libagx_draw_points(&ctx, p, patch, 3); } return; } } // Compute per-TessFactor metadata for (int edge = 0; edge < TRI_EDGES; edge++) { ComputeTessFactorCtx(outsideTessFactorOdd[edge], outsideTessFactor[edge], &outsideTessFactorCtx[edge]); } ComputeTessFactorCtx(insideTessFactorOdd, insideTessFactor, &insideTessFactorCtx); // Compute some initial data. int NumPoints = 0; // outside edge offsets and storage for (int edge = 0; edge < TRI_EDGES; edge++) { numPointsForOutsideEdge[edge] = NumPointsForTessFactor( outsideTessFactorOdd[edge], outsideTessFactor[edge]); NumPoints += numPointsForOutsideEdge[edge]; } NumPoints -= 3; // inside edge offsets numPointsForInsideTessFactor = NumPointsForTessFactor(insideTessFactorOdd, insideTessFactor); { int pointCountMin = insideTessFactorOdd ? 4 : 3; // max() allows degenerate transition regions when inside TessFactor == 1 numPointsForInsideTessFactor = max(pointCountMin, numPointsForInsideTessFactor); } insideEdgePointBaseOffset = NumPoints; // inside storage, including interior edges above { int interiorRings = (numPointsForInsideTessFactor >> 1) - 1; int even = insideTessFactorOdd ? 0 : 1; NumPoints += TRI_EDGES * (interiorRings * (interiorRings + even)) + even; } /* GENERATE POINTS */ if (mode != LIBAGX_TESS_MODE_COUNT) { ctx.Point = libagx_heap_alloc_points(p, patch, NumPoints); // Generate exterior ring edge points, clockwise starting from point V // (VW, the U==0 edge) int pointOffset = 0; for (int edge = 0; edge < TRI_EDGES; edge++) { int odd = edge & 0x1; int endPoint = numPointsForOutsideEdge[edge] - 1; // don't include end, since next edge starts with it. for (int p = 0; p < endPoint; p++, pointOffset++) { // whether to reverse point order given we are defining V or U (W // implicit): edge0, VW, has V decreasing, so reverse 1D points // below edge1, WU, has U increasing, so don't reverse 1D points // below edge2, UV, has U decreasing, so reverse 1D points below int q = odd ? p : endPoint - p; FXP fxpParam = PlacePointIn1D(&outsideTessFactorCtx[edge], outsideTessFactorOdd[edge], q); DefinePoint(&ctx.Point[pointOffset], (edge == 0) ? 0 : fxpParam, (edge == 0) ? fxpParam : (edge == 2) ? FXP_ONE - fxpParam : 0); } } // Generate interior ring points, clockwise spiralling in int numRings = (numPointsForInsideTessFactor >> 1); for (int ring = 1; ring < numRings; ring++) { int startPoint = ring; int endPoint = numPointsForInsideTessFactor - 1 - startPoint; int perpendicularAxisPoint = startPoint; FXP fxpPerpParam = PlacePointIn1D( &insideTessFactorCtx, insideTessFactorOdd, perpendicularAxisPoint); // Map location to the right size in // barycentric space. We know this fixed // point math won't over/underflow fxpPerpParam *= FXP_TWO_THIRDS; fxpPerpParam = (fxpPerpParam + FXP_ONE_HALF /*round*/) >> FXP_FRACTION_BITS; // get back to n.16 for (int edge = 0; edge < TRI_EDGES; edge++) { int odd = edge & 0x1; // don't include end: next edge starts with it. for (int p = startPoint; p < endPoint; p++, pointOffset++) { // whether to reverse point given we are defining V or U (W // implicit): edge0, VW, has V decreasing, so reverse 1D points // below edge1, WU, has U increasing, so don't reverse 1D points // below edge2, UV, has U decreasing, so reverse 1D points below int q = odd ? p : endPoint - (p - startPoint); FXP fxpParam = PlacePointIn1D(&insideTessFactorCtx, insideTessFactorOdd, q); // edge0 VW, has perpendicular parameter U constant // edge1 WU, has perpendicular parameter V constant // edge2 UV, has perpendicular parameter W constant // reciprocal is the rate of change of edge-parallel parameters // as they are pushed into the triangle const unsigned int deriv = 2; // we know this fixed point math won't over/underflow FXP tmp = fxpParam - (fxpPerpParam + 1 /*round*/) / deriv; DefinePoint(&ctx.Point[pointOffset], edge > 0 ? tmp : fxpPerpParam, edge == 0 ? tmp : edge == 1 ? fxpPerpParam : FXP_ONE - tmp - fxpPerpParam); } } } if (!insideTessFactorOdd) { // Last point is the point at the center. DefinePoint(&ctx.Point[pointOffset], FXP_ONE_THIRD, FXP_ONE_THIRD); } } if (p->points_mode) { libagx_draw_points(&ctx, p, patch, NumPoints); return; } { // Generate primitives for all the concentric rings, one side at a time // for each ring +1 is so even tess includes the center point, which we // want to now int numRings = ((numPointsForInsideTessFactor + 1) >> 1); int NumIndices = 0; { int OuterPoints = numPointsForOutsideEdge[0] + numPointsForOutsideEdge[1] + numPointsForOutsideEdge[2]; int numRings18 = numRings * 18; NumIndices = ((numRings18 - 27) * numPointsForInsideTessFactor) + (3 * OuterPoints) - (numRings18 * (numRings - 1)) + (insideTessFactorOdd ? 3 : 0); } // Generate the draw and allocate the index buffer now that we know the size ctx.Index = libagx_draw(p, mode, false, patch, NumIndices); if (mode == LIBAGX_TESS_MODE_COUNT) return; int insideOffset = insideEdgePointBaseOffset; int outsideEdgePointBaseOffset = 0; NumIndices = 0; for (int ring = 1; ring < numRings; ring++) { int numPointsForInsideEdge = numPointsForInsideTessFactor - 2 * ring; int edge0InsidePointBaseOffset = insideOffset; int edge0OutsidePointBaseOffset = outsideEdgePointBaseOffset; for (int edge = 0; edge < TRI_EDGES; edge++) { int outsidePoints = ring == 1 ? numPointsForOutsideEdge[edge] : (numPointsForInsideEdge + 2); int numTriangles = numPointsForInsideEdge + outsidePoints - 2; int insideBaseOffset; int outsideBaseOffset; if (edge == 2) { ctx.IndexPatchCtx.insidePointIndexDeltaToRealValue = insideOffset; ctx.IndexPatchCtx.insidePointIndexBadValue = numPointsForInsideEdge - 1; ctx.IndexPatchCtx.insidePointIndexReplacementValue = edge0InsidePointBaseOffset; ctx.IndexPatchCtx.outsidePointIndexPatchBase = ctx.IndexPatchCtx.insidePointIndexBadValue + 1; // past inside patched index range ctx.IndexPatchCtx.outsidePointIndexDeltaToRealValue = outsideEdgePointBaseOffset - ctx.IndexPatchCtx.outsidePointIndexPatchBase; ctx.IndexPatchCtx.outsidePointIndexBadValue = ctx.IndexPatchCtx.outsidePointIndexPatchBase + outsidePoints - 1; ctx.IndexPatchCtx.outsidePointIndexReplacementValue = edge0OutsidePointBaseOffset; ctx.bUsingPatchedIndices = true; insideBaseOffset = 0; outsideBaseOffset = ctx.IndexPatchCtx.outsidePointIndexPatchBase; } else { insideBaseOffset = insideOffset; outsideBaseOffset = outsideEdgePointBaseOffset; } if (ring == 1) { StitchTransition( &ctx, /*baseIndexOffset: */ NumIndices, insideBaseOffset, insideTessFactorCtx.numHalfTessFactorPoints, insideTessFactorOdd, outsideBaseOffset, outsideTessFactorCtx[edge].numHalfTessFactorPoints, outsideTessFactorOdd[edge]); } else { StitchRegular(&ctx, /*bTrapezoid*/ true, DIAGONALS_MIRRORED, /*baseIndexOffset: */ NumIndices, numPointsForInsideEdge, insideBaseOffset, outsideBaseOffset); } if (2 == edge) { ctx.bUsingPatchedIndices = false; } NumIndices += numTriangles * 3; outsideEdgePointBaseOffset += outsidePoints - 1; insideOffset += numPointsForInsideEdge - 1; } } if (insideTessFactorOdd) { // Triangulate center (a single triangle) DefineTriangle(&ctx, outsideEdgePointBaseOffset, outsideEdgePointBaseOffset + 1, outsideEdgePointBaseOffset + 2, NumIndices); NumIndices += 3; } } } KERNEL(64) libagx_tess_quad(constant struct libagx_tess_args *p, enum libagx_tess_mode mode__2) { enum libagx_tess_mode mode = mode__2; uint patch = get_global_id(0); enum libagx_tess_partitioning partitioning = p->partitioning; global float *factors = tess_factors(p, patch); float tessFactor_Ueq0 = factors[0]; float tessFactor_Veq0 = factors[1]; float tessFactor_Ueq1 = factors[2]; float tessFactor_Veq1 = factors[3]; float insideTessFactor_U = factors[4]; float insideTessFactor_V = factors[5]; struct CHWTessellator ctx = { .mode = mode, .index_bias = patch * LIBAGX_TES_PATCH_ID_STRIDE, }; // Is the patch culled? if (!(tessFactor_Ueq0 > 0) || // NaN will pass !(tessFactor_Veq0 > 0) || !(tessFactor_Ueq1 > 0) || !(tessFactor_Veq1 > 0)) { libagx_draw_empty(p, mode, patch); return; } FXP outsideTessFactor[QUAD_EDGES]; FXP insideTessFactor[QUAD_AXES]; bool outsideTessFactorOdd[QUAD_EDGES]; bool insideTessFactorOdd[QUAD_AXES]; TESS_FACTOR_CONTEXT outsideTessFactorCtx[QUAD_EDGES]; TESS_FACTOR_CONTEXT insideTessFactorCtx[QUAD_AXES]; // Stuff below is just specific to the traversal order // this code happens to use to generate points/lines int numPointsForOutsideEdge[QUAD_EDGES]; int numPointsForInsideTessFactor[QUAD_AXES]; int insideEdgePointBaseOffset; // Clamp edge TessFactors tessFactor_Ueq0 = clamp_factor(tessFactor_Ueq0, partitioning, 0); tessFactor_Veq0 = clamp_factor(tessFactor_Veq0, partitioning, 0); tessFactor_Ueq1 = clamp_factor(tessFactor_Ueq1, partitioning, 0); tessFactor_Veq1 = clamp_factor(tessFactor_Veq1, partitioning, 0); float maxf = max(max(max(tessFactor_Ueq0, tessFactor_Veq0), max(tessFactor_Ueq1, tessFactor_Veq1)), max(insideTessFactor_U, insideTessFactor_V)); insideTessFactor_U = clamp_factor(insideTessFactor_U, partitioning, maxf); insideTessFactor_V = clamp_factor(insideTessFactor_V, partitioning, maxf); // Note the above clamps map NaN to lowerBound // Process tessFactors float outsideTessFactor_f[QUAD_EDGES] = {tessFactor_Ueq0, tessFactor_Veq0, tessFactor_Ueq1, tessFactor_Veq1}; float insideTessFactor_f[QUAD_AXES] = {insideTessFactor_U, insideTessFactor_V}; if (partitioning == LIBAGX_TESS_PARTITIONING_INTEGER) { for (int edge = 0; edge < QUAD_EDGES; edge++) { outsideTessFactorOdd[edge] = isOdd(outsideTessFactor_f[edge]); } for (int axis = 0; axis < QUAD_AXES; axis++) { insideTessFactorOdd[axis] = isOdd(insideTessFactor_f[axis]) && (1.0f != insideTessFactor_f[axis]); } } else { bool odd = (partitioning == LIBAGX_TESS_PARTITIONING_FRACTIONAL_ODD); for (int edge = 0; edge < QUAD_EDGES; edge++) { outsideTessFactorOdd[edge] = odd; } insideTessFactorOdd[U] = insideTessFactorOdd[V] = odd; } // Save fixed point TessFactors for (int edge = 0; edge < QUAD_EDGES; edge++) { outsideTessFactor[edge] = floatToFixed(outsideTessFactor_f[edge]); } for (int axis = 0; axis < QUAD_AXES; axis++) { insideTessFactor[axis] = floatToFixed(insideTessFactor_f[axis]); } if (partitioning != LIBAGX_TESS_PARTITIONING_FRACTIONAL_EVEN) { // Special case if all TessFactors are 1 if ((FXP_ONE == insideTessFactor[U]) && (FXP_ONE == insideTessFactor[V]) && (FXP_ONE == outsideTessFactor[Ueq0]) && (FXP_ONE == outsideTessFactor[Veq0]) && (FXP_ONE == outsideTessFactor[Ueq1]) && (FXP_ONE == outsideTessFactor[Veq1])) { /* Just do minimum tess factor */ if (!p->points_mode) { ctx.Index = libagx_draw(p, mode, false, patch, 6); if (mode == LIBAGX_TESS_MODE_COUNT) return; DefineTriangle(&ctx, 0, 1, 3, /*indexStorageOffset*/ 0); DefineTriangle(&ctx, 1, 2, 3, /*indexStorageOffset*/ 3); } else { libagx_draw_points(&ctx, p, patch, 4); if (mode == LIBAGX_TESS_MODE_COUNT) return; } global struct libagx_tess_point *points = libagx_heap_alloc_points(p, patch, 4); DefinePoint(&points[0], 0, 0); DefinePoint(&points[1], FXP_ONE, 0); DefinePoint(&points[2], FXP_ONE, FXP_ONE); DefinePoint(&points[3], 0, FXP_ONE); return; } } // Compute TessFactor-specific metadata for (int edge = 0; edge < QUAD_EDGES; edge++) { ComputeTessFactorCtx(outsideTessFactorOdd[edge], outsideTessFactor[edge], &outsideTessFactorCtx[edge]); } for (int axis = 0; axis < QUAD_AXES; axis++) { ComputeTessFactorCtx(insideTessFactorOdd[axis], insideTessFactor[axis], &insideTessFactorCtx[axis]); } int NumPoints = 0; // outside edge offsets and storage for (int edge = 0; edge < QUAD_EDGES; edge++) { numPointsForOutsideEdge[edge] = NumPointsForTessFactor( outsideTessFactorOdd[edge], outsideTessFactor[edge]); NumPoints += numPointsForOutsideEdge[edge]; } NumPoints -= 4; // inside edge offsets for (int axis = 0; axis < QUAD_AXES; axis++) { numPointsForInsideTessFactor[axis] = NumPointsForTessFactor( insideTessFactorOdd[axis], insideTessFactor[axis]); int pointCountMin = insideTessFactorOdd[axis] ? 4 : 3; // max() allows degenerate transition regions when inside TessFactor == 1 numPointsForInsideTessFactor[axis] = max(pointCountMin, numPointsForInsideTessFactor[axis]); } insideEdgePointBaseOffset = NumPoints; // inside storage, including interior edges above int numInteriorPoints = (numPointsForInsideTessFactor[U] - 2) * (numPointsForInsideTessFactor[V] - 2); NumPoints += numInteriorPoints; if (mode != LIBAGX_TESS_MODE_COUNT) { ctx.Point = libagx_heap_alloc_points(p, patch, NumPoints); // Generate exterior ring edge points, clockwise from top-left int pointOffset = 0; for (int edge = 0; edge < QUAD_EDGES; edge++) { int odd = edge & 0x1; // don't include end, since next edge starts with it. int endPoint = numPointsForOutsideEdge[edge] - 1; for (int p = 0; p < endPoint; p++, pointOffset++) { int q = ((edge == 1) || (edge == 2)) ? p : endPoint - p; // reverse order FXP fxpParam = PlacePointIn1D(&outsideTessFactorCtx[edge], outsideTessFactorOdd[edge], q); FXP u = odd ? fxpParam : ((edge == 2) ? FXP_ONE : 0); FXP v = odd ? ((edge == 3) ? FXP_ONE : 0) : fxpParam; DefinePoint(&ctx.Point[pointOffset], u, v); } } // Generate interior ring points, clockwise from (U==0,V==1) (bottom-left) // spiralling toward center int minNumPointsForTessFactor = min(numPointsForInsideTessFactor[U], numPointsForInsideTessFactor[V]); // note for even tess we aren't counting center point here. int numRings = (minNumPointsForTessFactor >> 1); for (int ring = 1; ring < numRings; ring++) { int startPoint = ring; int endPoint[QUAD_AXES] = { numPointsForInsideTessFactor[U] - 1 - startPoint, numPointsForInsideTessFactor[V] - 1 - startPoint, }; for (int edge = 0; edge < QUAD_EDGES; edge++) { int odd[QUAD_AXES] = {edge & 0x1, ((edge + 1) & 0x1)}; int perpendicularAxisPoint = (edge < 2) ? startPoint : endPoint[odd[0]]; FXP fxpPerpParam = PlacePointIn1D(&insideTessFactorCtx[odd[0]], insideTessFactorOdd[odd[0]], perpendicularAxisPoint); for (int p = startPoint; p < endPoint[odd[1]]; p++, pointOffset++) // don't include end: next edge starts with // it. { bool odd_ = odd[1]; int q = ((edge == 1) || (edge == 2)) ? p : endPoint[odd_] - (p - startPoint); FXP fxpParam = PlacePointIn1D(&insideTessFactorCtx[odd_], insideTessFactorOdd[odd_], q); DefinePoint(&ctx.Point[pointOffset], odd_ ? fxpPerpParam : fxpParam, odd_ ? fxpParam : fxpPerpParam); } } } // For even tessellation, the inner "ring" is degenerate - a row of points if ((numPointsForInsideTessFactor[U] > numPointsForInsideTessFactor[V]) && !insideTessFactorOdd[V]) { int startPoint = numRings; int endPoint = numPointsForInsideTessFactor[U] - 1 - startPoint; for (int p = startPoint; p <= endPoint; p++, pointOffset++) { FXP fxpParam = PlacePointIn1D(&insideTessFactorCtx[U], insideTessFactorOdd[U], p); DefinePoint(&ctx.Point[pointOffset], fxpParam, FXP_ONE_HALF); } } else if ((numPointsForInsideTessFactor[V] >= numPointsForInsideTessFactor[U]) && !insideTessFactorOdd[U]) { int startPoint = numRings; int endPoint = numPointsForInsideTessFactor[V] - 1 - startPoint; for (int p = endPoint; p >= startPoint; p--, pointOffset++) { FXP fxpParam = PlacePointIn1D(&insideTessFactorCtx[V], insideTessFactorOdd[V], p); DefinePoint(&ctx.Point[pointOffset], FXP_ONE_HALF, fxpParam); } } } if (p->points_mode) { libagx_draw_points(&ctx, p, patch, NumPoints); return; } /* CONNECTIVITY */ { // Generate primitives for all the concentric rings, one side at a time // for each ring. +1 is so even tess includes the center point int numPointRowsToCenter[QUAD_AXES] = { (numPointsForInsideTessFactor[U] + 1) >> 1, (numPointsForInsideTessFactor[V] + 1) >> 1, }; int numRings = min(numPointRowsToCenter[U], numPointRowsToCenter[V]); /* Calculate # of indices so we can allocate */ { /* Handle main case */ int OuterPoints = numPointsForOutsideEdge[0] + numPointsForOutsideEdge[1] + numPointsForOutsideEdge[2] + numPointsForOutsideEdge[3]; int InnerPoints = numPointsForInsideTessFactor[U] + numPointsForInsideTessFactor[V]; int NumIndices = (OuterPoints * 3) + (12 * numRings * InnerPoints) - (InnerPoints * 18) - (24 * numRings * (numRings - 1)); /* Determine major/minor axes */ bool U_major = (numPointsForInsideTessFactor[U] > numPointsForInsideTessFactor[V]); unsigned M = U_major ? U : V; unsigned m = U_major ? V : U; /* Handle degenerate ring */ if (insideTessFactorOdd[m]) { NumIndices += 12 * ((numPointsForInsideTessFactor[M] >> 1) - (numPointsForInsideTessFactor[m] >> 1)); NumIndices += (insideTessFactorOdd[M] ? 6 : 12); } // Generate the draw and allocate the index buffer with the size ctx.Index = libagx_draw(p, mode, false, patch, NumIndices); } if (mode == LIBAGX_TESS_MODE_COUNT) return; int degeneratePointRing[QUAD_AXES] = { // Even partitioning causes degenerate row of points, // which results in exceptions to the point ordering conventions // when travelling around the rings counterclockwise. !insideTessFactorOdd[V] ? numPointRowsToCenter[V] - 1 : -1, !insideTessFactorOdd[U] ? numPointRowsToCenter[U] - 1 : -1, }; int numPointsForOutsideEdge_[QUAD_EDGES] = { numPointsForOutsideEdge[Ueq0], numPointsForOutsideEdge[Veq0], numPointsForOutsideEdge[Ueq1], numPointsForOutsideEdge[Veq1], }; int insideEdgePointBaseOffset_ = insideEdgePointBaseOffset; int outsideEdgePointBaseOffset = 0; int NumIndices = 0; for (int ring = 1; ring < numRings; ring++) { int numPointsForInsideEdge[QUAD_AXES] = { numPointsForInsideTessFactor[U] - 2 * ring, numPointsForInsideTessFactor[V] - 2 * ring}; int edge0InsidePointBaseOffset = insideEdgePointBaseOffset_; int edge0OutsidePointBaseOffset = outsideEdgePointBaseOffset; for (int edge = 0; edge < QUAD_EDGES; edge++) { int odd = (edge + 1) & 0x1; int numTriangles = numPointsForInsideEdge[odd] + numPointsForOutsideEdge_[edge] - 2; int insideBaseOffset; int outsideBaseOffset; // We need to patch the indexing so Stitch() can think it sees 2 // sequentially increasing rows of points, even though we have // wrapped around to the end of the inner and outer ring's points, // so the last point is really the first point for the ring. We make // it so that when Stitch() calls AddIndex(), that function will do // any necessary index adjustment. if (edge == 3) { if (ring == degeneratePointRing[odd]) { ctx.IndexPatchCtx2.baseIndexToInvert = insideEdgePointBaseOffset_ + 1; ctx.IndexPatchCtx2.cornerCaseBadValue = outsideEdgePointBaseOffset + numPointsForOutsideEdge_[edge] - 1; ctx.IndexPatchCtx2.cornerCaseReplacementValue = edge0OutsidePointBaseOffset; ctx.IndexPatchCtx2.indexInversionEndPoint = (ctx.IndexPatchCtx2.baseIndexToInvert << 1) - 1; insideBaseOffset = ctx.IndexPatchCtx2.baseIndexToInvert; outsideBaseOffset = outsideEdgePointBaseOffset; ctx.bUsingPatchedIndices2 = true; } else { ctx.IndexPatchCtx.insidePointIndexDeltaToRealValue = insideEdgePointBaseOffset_; ctx.IndexPatchCtx.insidePointIndexBadValue = numPointsForInsideEdge[odd] - 1; ctx.IndexPatchCtx.insidePointIndexReplacementValue = edge0InsidePointBaseOffset; ctx.IndexPatchCtx.outsidePointIndexPatchBase = ctx.IndexPatchCtx.insidePointIndexBadValue + 1; // past inside patched index range ctx.IndexPatchCtx.outsidePointIndexDeltaToRealValue = outsideEdgePointBaseOffset - ctx.IndexPatchCtx.outsidePointIndexPatchBase; ctx.IndexPatchCtx.outsidePointIndexBadValue = ctx.IndexPatchCtx.outsidePointIndexPatchBase + numPointsForOutsideEdge_[edge] - 1; ctx.IndexPatchCtx.outsidePointIndexReplacementValue = edge0OutsidePointBaseOffset; insideBaseOffset = 0; outsideBaseOffset = ctx.IndexPatchCtx.outsidePointIndexPatchBase; ctx.bUsingPatchedIndices = true; } } else if ((edge == 2) && (ring == degeneratePointRing[odd])) { ctx.IndexPatchCtx2.baseIndexToInvert = insideEdgePointBaseOffset_; ctx.IndexPatchCtx2.cornerCaseBadValue = -1; // unused ctx.IndexPatchCtx2.cornerCaseReplacementValue = -1; // unused ctx.IndexPatchCtx2.indexInversionEndPoint = ctx.IndexPatchCtx2.baseIndexToInvert << 1; insideBaseOffset = ctx.IndexPatchCtx2.baseIndexToInvert; outsideBaseOffset = outsideEdgePointBaseOffset; ctx.bUsingPatchedIndices2 = true; } else { insideBaseOffset = insideEdgePointBaseOffset_; outsideBaseOffset = outsideEdgePointBaseOffset; } if (ring == 1) { StitchTransition( &ctx, /*baseIndexOffset: */ NumIndices, insideBaseOffset, insideTessFactorCtx[odd].numHalfTessFactorPoints, insideTessFactorOdd[odd], outsideBaseOffset, outsideTessFactorCtx[edge].numHalfTessFactorPoints, outsideTessFactorOdd[edge]); } else { StitchRegular(&ctx, /*bTrapezoid*/ true, DIAGONALS_MIRRORED, /*baseIndexOffset: */ NumIndices, numPointsForInsideEdge[odd], insideBaseOffset, outsideBaseOffset); } ctx.bUsingPatchedIndices = false; ctx.bUsingPatchedIndices2 = false; NumIndices += numTriangles * 3; outsideEdgePointBaseOffset += numPointsForOutsideEdge_[edge] - 1; if ((edge == 2) && (ring == degeneratePointRing[odd])) { insideEdgePointBaseOffset_ -= numPointsForInsideEdge[odd] - 1; } else { insideEdgePointBaseOffset_ += numPointsForInsideEdge[odd] - 1; } numPointsForOutsideEdge_[edge] = numPointsForInsideEdge[odd]; } } // Triangulate center - a row of quads if odd // This triangulation may be producing diagonals that are asymmetric about // the center of the patch in this region. if ((numPointsForInsideTessFactor[U] > numPointsForInsideTessFactor[V]) && insideTessFactorOdd[V]) { ctx.bUsingPatchedIndices2 = true; int stripNumQuads = (((numPointsForInsideTessFactor[U] >> 1) - (numPointsForInsideTessFactor[V] >> 1)) << 1) + (insideTessFactorOdd[U] ? 1 : 2); ctx.IndexPatchCtx2.baseIndexToInvert = outsideEdgePointBaseOffset + stripNumQuads + 2; ctx.IndexPatchCtx2.cornerCaseBadValue = ctx.IndexPatchCtx2.baseIndexToInvert; ctx.IndexPatchCtx2.cornerCaseReplacementValue = outsideEdgePointBaseOffset; ctx.IndexPatchCtx2.indexInversionEndPoint = ctx.IndexPatchCtx2.baseIndexToInvert + ctx.IndexPatchCtx2.baseIndexToInvert + stripNumQuads; StitchRegular( &ctx, /*bTrapezoid*/ false, DIAGONALS_INSIDE_TO_OUTSIDE, /*baseIndexOffset: */ NumIndices, /*numInsideEdgePoints:*/ stripNumQuads + 1, /*insideEdgePointBaseOffset*/ ctx.IndexPatchCtx2.baseIndexToInvert, outsideEdgePointBaseOffset + 1); ctx.bUsingPatchedIndices2 = false; NumIndices += stripNumQuads * 6; } else if ((numPointsForInsideTessFactor[V] >= numPointsForInsideTessFactor[U]) && insideTessFactorOdd[U]) { ctx.bUsingPatchedIndices2 = true; int stripNumQuads = (((numPointsForInsideTessFactor[V] >> 1) - (numPointsForInsideTessFactor[U] >> 1)) << 1) + (insideTessFactorOdd[V] ? 1 : 2); ctx.IndexPatchCtx2.baseIndexToInvert = outsideEdgePointBaseOffset + stripNumQuads + 1; ctx.IndexPatchCtx2.cornerCaseBadValue = -1; // unused ctx.IndexPatchCtx2.indexInversionEndPoint = ctx.IndexPatchCtx2.baseIndexToInvert + ctx.IndexPatchCtx2.baseIndexToInvert + stripNumQuads; DIAGONALS diag = insideTessFactorOdd[V] ? DIAGONALS_INSIDE_TO_OUTSIDE_EXCEPT_MIDDLE : DIAGONALS_INSIDE_TO_OUTSIDE; StitchRegular( &ctx, /*bTrapezoid*/ false, diag, /*baseIndexOffset: */ NumIndices, /*numInsideEdgePoints:*/ stripNumQuads + 1, /*insideEdgePointBaseOffset*/ ctx.IndexPatchCtx2.baseIndexToInvert, outsideEdgePointBaseOffset); ctx.bUsingPatchedIndices2 = false; NumIndices += stripNumQuads * 6; } } }