From e4dab6b9c38425cd03c47e265c9bc0211d1b11b8 Mon Sep 17 00:00:00 2001 From: Rye Date: Mon, 26 Oct 2026 01:20:14 -0400 Subject: [PATCH] A sculpt's pinch test trusts its samples' codes and writes kept triangles by the unit A sculpt side finds the triangles two of whose corners are one point from each point's code, unit_cube_codes' 24 bits. It confirmed every flagged triangle against its positions, wrote each row with a flag a quad at a time, and tested every flag word of a row again for each band. It writes the same indices, word for word, for less: - A sculpt built from its map has the map's samples for points, each axis one of 256 values, and unit_cube_codes gives such a point its bytes as its code (now checked for every byte, mirrored or not): two samples have one code exactly where they are one point. Only a placeholder's points, the sphere's and the empty one's, are confirmed against positions now. - grid_coincident_triangles counts the flagged triangles and marks the rows that have any, which it had the figures for, gathering each word of a row's flags in registers rather than or-ing into memory. - The codes and flags take one block, not three: the room a doubled flat inner side keeps for its lefts, which a sculpt's side never uses, where they fit (LOD 0 and 1), and one scratch block past that. createSide's frame is as it was. - kept_bands writes a band's rows between those with a flag as units in a run, found from the marked rows. kept_columns writes each run of columns that keep the same triangles of a grid of up to seven rows, a sculpt's LOD 0, as one unit. The per-quad tests take unsigned bit arithmetic. - unit_bands, unit_columns, grid_bands and grid_indices for KeepAll and Kept, which every prim side runs, compile to the same instructions as on develop. Not done: a band of seven quads for unpinched sculpts. It would win back their 2% (THE BAND-OF-7 THRASH), but it changes their index order. Finding pinches once per sculpt from the map buys nothing over once per side, a sculpt having one, and would have to repeat the sampler's /fp:fast arithmetic to know which pixel each vertex read. The new test lists each sculpt side's indices from a model of the side order and the exact position test and compares them word for word: 4,560 sides of every stitching and flag at every LOD, from maps whose samples meet along rows, columns and folds, at single points, in no pattern or not at all, the placeholders, and parameters with a hole or a cut. It passes on develop's llvolume.cpp and on this one. It fails when placeholders skip the confirmation, when the column runs are taken as one, and when the marked rows are ignored; the existing test, which compares triangle sets, catches only the last. A probe over 24,552 volumes and 49,192 faces (16 map sizes; few-colour, stepped, pinched, folded, pinwheel, one-colour and noise maps; eight parameter sets; placeholders; unique volumes) found every face's indices identical to develop's. llvolume_bench gains a pinched plane's rows. Cycles of a whole sculpt(): rdtsc medians, each run pinned to one core with no build running, the mean of four code layouts by three rounds; "no test" is develop with the pinch test switched off, "of test" the share of its cost this removes: develop no test this change of test sphere sculpt, LOD 0 4786 4202 4556 -4.8% 39% sphere sculpt, LOD 1 8586 7664 8370 -2.5% 23% sphere sculpt, LOD 2 22577 20873 21960 -2.7% 36% sphere sculpt, LOD 3 69354 64470 66665 -3.9% 55% torus sculpt, LOD 0 4444 4208 4362 -1.8% 35% torus sculpt, LOD 1 7537 7077 7378 -2.1% 35% torus sculpt, LOD 2 21266 20034 20830 -2.1% 35% torus sculpt, LOD 3 67488 63864 65831 -2.5% 46% plane sculpt, LOD 0 4376 4139 4325 -1.2% 22% plane sculpt, LOD 1 7478 7021 7362 -1.6% 25% plane sculpt, LOD 2 21100 20035 20701 -1.9% 37% plane sculpt, LOD 3 67307 63743 65476 -2.7% 51% cylinder sculpt, LOD 0 4400 4155 4306 -2.1% 38% cylinder sculpt, LOD 1 7487 7044 7358 -1.7% 29% cylinder sculpt, LOD 2 21171 19966 20723 -2.1% 37% cylinder sculpt, LOD 3 67527 63811 65576 -2.9% 53% pinched plane sculpt, LOD 0 4994 4334 4761 -4.7% 35% pinched plane sculpt, LOD 1 8705 7423 8115 -6.8% 46% pinched plane sculpt, LOD 2 25750 22817 24171 -6.1% 54% pinched plane sculpt, LOD 3 90941 81255 84686 -6.9% 65% Whole builds move 2-3% with code placement even so: a variant of this change that differed only in KeepAll's band loop measured 33-95% "of test". Timed in place, the pinch path alone (codes, flags, emission) takes 24-36% less on the sphere sculpt at every LOD. Prims are within 2% of develop either way, as the no-test build, which leaves prims as they are, is too. llvolume_bench itself, pinned, eight rounds of each binary, reads -4.8% to +6.4% for sculpts and -2.2% to +5.2% for prims; develop relinked at another layout reads -13.8% to +7.5% for prims against itself. Co-Authored-By: Claude Opus 5.5 --- indra/llmath/alsimdkernels.cpp | 4 +- indra/llmath/alsimdkernels.h | 8 +- indra/llmath/alsimdkernels.inl | 95 ++++--- indra/llmath/llvolume.cpp | 326 +++++++++++++++++++--- indra/llmath/tests/alsimdkernels_test.cpp | 66 ++++- indra/llmath/tests/llvolume_bench.cpp | 32 +++ indra/llmath/tests/llvolume_test.cpp | 196 +++++++++++++ 7 files changed, 625 insertions(+), 102 deletions(-) diff --git a/indra/llmath/alsimdkernels.cpp b/indra/llmath/alsimdkernels.cpp index 50f9445124..b714c4f97b 100644 --- a/indra/llmath/alsimdkernels.cpp +++ b/indra/llmath/alsimdkernels.cpp @@ -143,9 +143,9 @@ void unit_cube_codes(const LLVector4a* src, size_t n, U32* dst) kernels::unit_cube_codes_impl(src, n, dst); } -size_t grid_coincident_triangles(const U32* codes, size_t columns, size_t rows, U64* upper, U64* lower, size_t words) +size_t grid_coincident_triangles(const U32* codes, size_t columns, size_t rows, U64* upper, U64* lower, size_t words, U64* flagged_rows) { - return kernels::grid_coincident_triangles_impl(codes, columns, rows, upper, lower, words); + return kernels::grid_coincident_triangles_impl(codes, columns, rows, upper, lower, words, flagged_rows); } } // namespace alsimd diff --git a/indra/llmath/alsimdkernels.h b/indra/llmath/alsimdkernels.h index 3164fb441b..eb79d0126d 100644 --- a/indra/llmath/alsimdkernels.h +++ b/indra/llmath/alsimdkernels.h @@ -128,9 +128,11 @@ void unit_cube_codes(const LLVector4a* src, size_t n, U32* dst); // s + columns * t, that have two corners of one code. The quad at column s // and row t is bottom left, top right, top left, flagged in bit s % 64 of // upper[t * words + s / 64], and bottom left, bottom right, top right, in -// lower; words is at least (columns + 62) / 64. codes is read up to eight -// past its columns * rows. Returns how many rows of quads have a flag. -size_t grid_coincident_triangles(const U32* codes, size_t columns, size_t rows, U64* upper, U64* lower, size_t words); +// lower; words is at least (columns + 62) / 64. Bit t % 64 of +// flagged_rows[t / 64] is set where row t of quads has a flag, and the +// rest of its (rows + 62) / 64 words are cleared. codes is read up to eight +// past its columns * rows. Returns how many triangles are flagged. +size_t grid_coincident_triangles(const U32* codes, size_t columns, size_t rows, U64* upper, U64* lower, size_t words, U64* flagged_rows); // Three little-endian 16-bit values per vertex, six bytes apart, to // {x, y, z, 0} * scale + bias; the fourth lane of scale is ignored, so w diff --git a/indra/llmath/alsimdkernels.inl b/indra/llmath/alsimdkernels.inl index 50f71ad356..511eeac452 100644 --- a/indra/llmath/alsimdkernels.inl +++ b/indra/llmath/alsimdkernels.inl @@ -28,6 +28,8 @@ #include "alsimdkernels.h" +#include +#include #include #include #include @@ -643,13 +645,18 @@ void unit_cube_codes_impl(const LLVector4a* src, size_t n, U32* dst) } } -// A row of quads at a time, four quads a step, eight under AVX2. A step -// reads its bottom left, bottom right, top left and top right corners' -// codes four or eight at once, which runs past the row's last quad: the -// flags past it are cleared. +// A row of quads at a time and a word of its flags at a time, gathered in +// registers, four quads a step, eight under AVX2. A step reads its bottom +// left, bottom right, top left and top right corners' codes four or eight +// at once, which runs past the row's last quad: the flags past it are +// cleared. template -size_t grid_coincident_triangles_impl(const U32* codes, size_t columns, size_t rows, U64* upper, U64* lower, size_t words) +size_t grid_coincident_triangles_impl(const U32* codes, size_t columns, size_t rows, U64* upper, U64* lower, size_t words, U64* flagged_rows) { + for (size_t w = 0; w < (rows + 62) / 64; ++w) + { + flagged_rows[w] = 0; + } if (columns < 2 || rows < 2) { return 0; @@ -662,51 +669,53 @@ size_t grid_coincident_triangles_impl(const U32* codes, size_t columns, size_t r { const int32_t* bottom = reinterpret_cast(codes + columns * t); const int32_t* top = bottom + columns; - U64* up = upper + words * t; - U64* lo = lower + words * t; + U64 any = 0; for (size_t w = 0; w < words; ++w) { - up[w] = 0; - lo[w] = 0; - } - size_t s = 0; + U64 up = 0; + U64 lo = 0; + const size_t end = std::min(quads, (w + 1) * 64); + size_t s = w * 64; #if AL_SIMD_AVX2 - if constexpr (W::lanes >= 2) - { - for (; s < quads; s += 8) + if constexpr (W::lanes >= 2) { - const __m256i bl = _mm256_loadu_si256(reinterpret_cast(bottom + s)); - const __m256i br = _mm256_loadu_si256(reinterpret_cast(bottom + s + 1)); - const __m256i tl = _mm256_loadu_si256(reinterpret_cast(top + s)); - const __m256i tr = _mm256_loadu_si256(reinterpret_cast(top + s + 1)); - const __m256i diagonal = _mm256_cmpeq_epi32(bl, tr); - const __m256i u = _mm256_or_si256(_mm256_or_si256(diagonal, _mm256_cmpeq_epi32(tl, tr)), _mm256_cmpeq_epi32(bl, tl)); - const __m256i l = _mm256_or_si256(_mm256_or_si256(diagonal, _mm256_cmpeq_epi32(bl, br)), _mm256_cmpeq_epi32(br, tr)); - up[s / 64] |= U64(_mm256_movemask_ps(_mm256_castsi256_ps(u))) << (s % 64); - lo[s / 64] |= U64(_mm256_movemask_ps(_mm256_castsi256_ps(l))) << (s % 64); + for (; s < end; s += 8) + { + const __m256i bl = _mm256_loadu_si256(reinterpret_cast(bottom + s)); + const __m256i br = _mm256_loadu_si256(reinterpret_cast(bottom + s + 1)); + const __m256i tl = _mm256_loadu_si256(reinterpret_cast(top + s)); + const __m256i tr = _mm256_loadu_si256(reinterpret_cast(top + s + 1)); + const __m256i diagonal = _mm256_cmpeq_epi32(bl, tr); + const __m256i u = _mm256_or_si256(_mm256_or_si256(diagonal, _mm256_cmpeq_epi32(tl, tr)), _mm256_cmpeq_epi32(bl, tl)); + const __m256i l = _mm256_or_si256(_mm256_or_si256(diagonal, _mm256_cmpeq_epi32(bl, br)), _mm256_cmpeq_epi32(br, tr)); + up |= U64(_mm256_movemask_ps(_mm256_castsi256_ps(u))) << (s % 64); + lo |= U64(_mm256_movemask_ps(_mm256_castsi256_ps(l))) << (s % 64); + } } - } #endif - for (; s < quads; s += 4) - { - const i32x4 bl = alsimd::loadu_i32(bottom + s); - const i32x4 br = alsimd::loadu_i32(bottom + s + 1); - const i32x4 tl = alsimd::loadu_i32(top + s); - const i32x4 tr = alsimd::loadu_i32(top + s + 1); - const mask4 diagonal = alsimd::cmpeq_i32(bl, tr); - const mask4 u = alsimd::or_(alsimd::or_(diagonal, alsimd::cmpeq_i32(tl, tr)), alsimd::cmpeq_i32(bl, tl)); - const mask4 l = alsimd::or_(alsimd::or_(diagonal, alsimd::cmpeq_i32(bl, br)), alsimd::cmpeq_i32(br, tr)); - up[s / 64] |= U64(alsimd::bits(u)) << (s % 64); - lo[s / 64] |= U64(alsimd::bits(l)) << (s % 64); - } - up[last_word] &= last_bits; - lo[last_word] &= last_bits; - U64 any = 0; - for (size_t w = 0; w <= last_word; ++w) - { - any |= up[w] | lo[w]; + for (; s < end; s += 4) + { + const i32x4 bl = alsimd::loadu_i32(bottom + s); + const i32x4 br = alsimd::loadu_i32(bottom + s + 1); + const i32x4 tl = alsimd::loadu_i32(top + s); + const i32x4 tr = alsimd::loadu_i32(top + s + 1); + const mask4 diagonal = alsimd::cmpeq_i32(bl, tr); + const mask4 u = alsimd::or_(alsimd::or_(diagonal, alsimd::cmpeq_i32(tl, tr)), alsimd::cmpeq_i32(bl, tl)); + const mask4 l = alsimd::or_(alsimd::or_(diagonal, alsimd::cmpeq_i32(bl, br)), alsimd::cmpeq_i32(br, tr)); + up |= U64(alsimd::bits(u)) << (s % 64); + lo |= U64(alsimd::bits(l)) << (s % 64); + } + if (w == last_word) + { + up &= last_bits; + lo &= last_bits; + } + upper[words * t + w] = up; + lower[words * t + w] = lo; + any |= up | lo; + flagged += std::popcount(up) + std::popcount(lo); } - flagged += any != 0; + flagged_rows[t / 64] |= U64(any != 0) << (t % 64); } return flagged; } diff --git a/indra/llmath/llvolume.cpp b/indra/llmath/llvolume.cpp index d90eb37dca..fd38969364 100644 --- a/indra/llmath/llvolume.cpp +++ b/indra/llmath/llvolume.cpp @@ -6774,11 +6774,13 @@ namespace // as having two corners in one place. whole(t) is whether every // triangle of row t is kept. WHOLE_ROWS is whether rows are whole often // enough to write a band's row at once: a sphere's collapsed poles - // leave a triangle out of every row. + // leave a triangle out of every row. COLUMN_UNITS is whether a grid of + // a few rows writes a column at once. struct KeepAll { static constexpr bool ALL = true; static constexpr bool WHOLE_ROWS = true; + static constexpr bool COLUMN_UNITS = true; bool upper(S32, S32) const { return true; } bool lower(S32, S32) const { return true; } @@ -6789,6 +6791,7 @@ namespace { static constexpr bool ALL = false; static constexpr bool WHOLE_ROWS = false; + static constexpr bool COLUMN_UNITS = false; bool first_row; bool last_row; @@ -6802,27 +6805,101 @@ namespace bool whole(S32 t) const { return !first_column && !last_column && !(first_row && t == 0) && !(last_row && t == last_quad_row); } }; + // Bit i of a set held in 64-bit words. + LL_FORCE_INLINE bool bit_set(const U64* set, U32 i) + { + return (set[i / 64] >> (i % 64)) & 1; + } + + // A keep that names the rows that are not whole, rather than saying + // whole(t) of each: partial_from(t, rows) is the first from row t up, + // or rows, and partial_to(t) the last from row t down, or -1. column(s, + // rows) is what it keeps of column s, bit 2t for the first triangle of + // its quad at row t and 2t + 1 for the second, and column_change(s, end, + // rows) the first column from s on that keeps other triangles than the + // one before it, or end. struct Unpinched { static constexpr bool ALL = false; static constexpr bool WHOLE_ROWS = true; + static constexpr bool COLUMN_UNITS = true; - // Bit s % 64 of word t * words + s / 64 for the quad at column s - // and row t. + // Bit s of the words from t * words on for the quad at column s and + // row t, and bit t of flagged_rows for a row with any of them set. const U64* upper_pinched; const U64* lower_pinched; - S32 words; + const U64* flagged_rows; + U32 words; - bool upper(S32 s, S32 t) const { return !((upper_pinched[t * words + s / 64] >> (s % 64)) & 1); } - bool lower(S32 s, S32 t) const { return !((lower_pinched[t * words + s / 64] >> (s % 64)) & 1); } - bool whole(S32 t) const + bool upper(S32 s, S32 t) const { return !bit_set(upper_pinched + (U32)t * words, (U32)s); } + bool lower(S32 s, S32 t) const { return !bit_set(lower_pinched + (U32)t * words, (U32)s); } + + S32 partial_from(S32 t, S32 rows) const { - U64 any = 0; - for (S32 w = 0; w < words; ++w) + U32 w = (U32)t / 64; + U64 bits = flagged_rows[w] & (~U64(0) << ((U32)t % 64)); + while (bits == 0) + { + if (++w * 64 >= (U32)rows) + { + return rows; + } + bits = flagged_rows[w]; + } + return (S32)(w * 64 + std::countr_zero(bits)); + } + + S32 partial_to(S32 t) const + { + S32 w = t / 64; + U64 bits = flagged_rows[w] & (~U64(0) >> (63 - t % 64)); + while (bits == 0) + { + if (--w < 0) + { + return -1; + } + bits = flagged_rows[w]; + } + return w * 64 + 63 - std::countl_zero(bits); + } + + U32 column(S32 s, S32 rows) const + { + const U32 w = (U32)s / 64; + const U32 b = (U32)s % 64; + U32 kept = 0; + for (S32 t = 0; t < rows; ++t) { - any |= upper_pinched[t * words + w] | lower_pinched[t * words + w]; + kept |= (U32)(~(upper_pinched[t * words + w] >> b) & 1) << (2 * t); + kept |= (U32)(~(lower_pinched[t * words + w] >> b) & 1) << (2 * t + 1); } - return any == 0; + return kept; + } + + // s is at least 1. + S32 column_change(S32 s, S32 end, S32 rows) const + { + for (U32 w = (U32)s / 64; w * 64 < (U32)end; ++w) + { + U64 changes = 0; + for (S32 t = 0; t < rows; ++t) + { + const U64* up = upper_pinched + t * words; + const U64* lo = lower_pinched + t * words; + changes |= up[w] ^ ((up[w] << 1) | (w > 0 ? up[w - 1] >> 63 : 0)); + changes |= lo[w] ^ ((lo[w] << 1) | (w > 0 ? lo[w - 1] >> 63 : 0)); + } + if (w == (U32)s / 64) + { + changes &= ~U64(0) << ((U32)s % 64); + } + if (changes != 0) + { + return llmin((S32)(w * 64 + std::countr_zero(changes)), end); + } + } + return end; } }; @@ -6884,6 +6961,32 @@ namespace std::fill(offsets + length, offsets + vectors * 8, (U16)0); } + // The same quads but for the triangles kept leaves out: bit 2q + // keeps quad q's first triangle, and 2q + 1 its second. A unit + // that keeps none has no length. + QuadUnit(S32 quad_count, S32 step, S32 num_s, U32 kept) + : quads(quad_count), length(0) + { + for (S32 q = 0; q < quads; ++q) + { + const S32 bottom = q * step; + const S32 top = bottom + num_s; + if (kept & (1u << (2 * q))) + { + offsets[length++] = (U16)bottom; //bottom left + offsets[length++] = (U16)(top + 1); //top right + offsets[length++] = (U16)top; //top left + } + if (kept & (2u << (2 * q))) + { + offsets[length++] = (U16)bottom; //bottom left + offsets[length++] = (U16)(bottom + 1); //bottom right + offsets[length++] = (U16)(top + 1); //top right + } + } + vectors = (length + 7) / 8; + std::fill(offsets + length, offsets + vectors * 8, (U16)0); + } }; // A unit while it is written, its offsets held as values: a store of an @@ -6960,6 +7063,30 @@ namespace return out; } + // The first count columns of rows quads, each a unit of what keep keeps + // of it: one unit for each run of columns that keep the same + // triangles. Nothing is written past limit. Returns the end of what it + // wrote. + template + VOLUME_NOINLINE U16* kept_columns(U16* out, const U16* limit, S32 count, S32 num_s, S32 rows, const KEEP& keep) + { + for (S32 q = 0; q < count;) + { + const QuadUnit column(rows, num_s, num_s, keep.column(q, rows)); + const UnitWriter writer(column); + const S32 end = keep.column_change(q + 1, count, rows); + if (column.length == 0) + { + q = end; + } + for (; q < end; ++q) + { + out = writer.write(out, q, limit); + } + } + return out; + } + // The quads from grid column begin on in bands, the first wide of them // a quad wider than narrow, of rows quads up: each band row by row, // every other band down, a unit a band's row where keep keeps the whole @@ -7006,6 +7133,59 @@ namespace return out; } + // The same bands for a keep that names the rows it leaves triangles out + // of: the whole rows between them a unit each, in a run. + template + VOLUME_NOINLINE U16* kept_bands(U16* out, const U16* limit, S32 begin, S32 bands, S32 narrow, S32 wide, S32 num_s, S32 rows, const KEEP& keep) + { + const QuadUnit narrow_row(narrow, 1, num_s); + const QuadUnit wide_row(wide > 0 ? narrow + 1 : narrow, 1, num_s); + S32 first = begin; + for (S32 band = 0; band < bands; ++band) + { + const S32 width = band < wide ? narrow + 1 : narrow; + const UnitWriter writer(band < wide ? wide_row : narrow_row); + const auto partial_row = [&](U16* at, S32 t) + { + for (S32 s = first; s < first + width; ++s) + { + at = quad_indices(at, s, t, num_s, keep); + } + return at; + }; + if (band & 1) + { + for (S32 t = rows - 1; t >= 0;) + { + for (const S32 partial = keep.partial_to(t); t > partial; --t) + { + out = writer.write(out, first + num_s * t, limit); + } + if (t >= 0) + { + out = partial_row(out, t--); + } + } + } + else + { + for (S32 t = 0; t < rows;) + { + for (const S32 partial = keep.partial_from(t, rows); t < partial; ++t) + { + out = writer.write(out, first + num_s * t, limit); + } + if (t < rows) + { + out = partial_row(out, t++); + } + } + } + first += width; + } + return out; + } + // A grid of up to COLUMN_ROWS rows, a column of quads at a time: the // count columns listed, or the first count. Nothing is written past // limit. Returns the end of what it wrote. @@ -7016,6 +7196,13 @@ namespace { return unit_columns(out, limit, lefts, count, num_s, rows); } + if constexpr (!KEEP::ALL && KEEP::COLUMN_UNITS && !LISTED) + { + if (rows >= UNIT_QUADS && count >= UNIT_USES) + { + return kept_columns(out, limit, count, num_s, rows, keep); + } + } if (rows == 1) { for (S32 q = 0; q < count; ++q) @@ -7049,7 +7236,14 @@ namespace { if (narrow >= UNIT_QUADS) { - return unit_bands(out, limit, begin, bands, narrow, wide, num_s, rows, keep); + if constexpr (KEEP::ALL) + { + return unit_bands(out, limit, begin, bands, narrow, wide, num_s, rows, keep); + } + else + { + return kept_bands(out, limit, begin, bands, narrow, wide, num_s, rows, keep); + } } } if (bands == 1) @@ -7153,6 +7347,74 @@ namespace return pinched; } + // Sets bit t of flagged, and clears the rest of its words, for each of + // rows rows of words words whose upper or lower flags have any set. + void flag_rows(const U64* upper, const U64* lower, S32 words, S32 rows, U64* flagged) + { + std::fill(flagged, flagged + (rows + 63) / 64, U64(0)); + for (S32 t = 0; t < rows; ++t) + { + U64 any = 0; + for (S32 w = 0; w < words; ++w) + { + any |= upper[t * words + w] | lower[t * words + w]; + } + flagged[t / 64] |= U64(any != 0) << (t % 64); + } + } + + // Whether a sculpt's points are its map's samples, each axis one of the + // 256 values a byte takes it to, mirrored or not. unit_cube_codes gives + // such points a code each, so two are one point exactly where their + // codes are one code. A shared sculpt's faces are built as + // LLVolume::sculpt ends, from the map or from a placeholder, whose + // points are not samples. + bool points_are_samples(const LLVolume& volume) + { + return volume.getSculptPixels() > 0 && !volume.isSculptSpherePlaceholder(); + } + + // A sculpt side's triangles, from its num_s by num_t grid of points, in + // the order above, but for those two of whose corners are one point: + // they have no area. The points' codes find them, and where the points + // are not the map's samples the points confirm them. A side keeps at + // least one triangle, of the triangles it has. Nothing is written past + // limit. The work space is room where it fits, as it does a small + // grid's, and the scratch arena where it does not. Returns the end of + // what it wrote. + LL_FORCE_INLINE U16* sculpt_indices(U16* out, const U16* limit, const LLVolume& volume, const LLVector4a* pos, S32 num_s, S32 num_t, S32 triangles, + std::byte* room, size_t room_bytes) + { + const S32 quads = num_s - 1; + const S32 rows = num_t - 1; + const S32 num_vertices = num_s * num_t; + const S32 words = (quads + 63) / 64; + const size_t flag_words = (size_t)words * rows; + // The codes, and the eight past them the kernel reads, in whole + // words, then the flags and the rows that have any. + const size_t code_words = ((size_t)num_vertices + 9) / 2; + const size_t block_words = code_words + 2 * flag_words + (rows + 63) / 64; + ALScratchArena::Scope scratch; + U64* block = block_words * sizeof(U64) <= room_bytes ? reinterpret_cast(room) : scratch.allocate(block_words); + U32* codes = reinterpret_cast(block); + U64* upper = block + code_words; + U64* lower = upper + flag_words; + U64* flagged_rows = lower + flag_words; + alsimd::unit_cube_codes(pos, num_vertices, codes); + std::fill(codes + num_vertices, codes + num_vertices + 8, 0u); + S32 pinched = (S32)alsimd::grid_coincident_triangles(codes, num_s, num_t, upper, lower, words, flagged_rows); + if (pinched > 0 && !points_are_samples(volume)) + { + pinched = confirm_pinches(pos, num_s, num_t, upper, lower, words); + flag_rows(upper, lower, words, rows, flagged_rows); + } + if (pinched > 0 && pinched < triangles) + { + return grid_indices(out, limit, num_s, num_t, nullptr, quads, Unpinched{ upper, lower, flagged_rows, (U32)words }); + } + return grid_indices(out, limit, num_s, num_t, nullptr, quads, KeepAll()); + } + // The path station a cap closes: the last for a top cap, the first for a // bottom one. S32 cap_station(const LLVolume* volume, bool top) @@ -8118,10 +8380,12 @@ bool LLVolumeFace::createSide(LLVolume* volume, bool partial_build) Column stack_columns[STACK_COLUMNS]; std::vector heap_columns; Column* columns = stack_columns; - // A flat inner side's quads with area, by their left columns. - S32 stack_lefts[STACK_COLUMNS]; + // A flat inner side's quads with area, by their left columns. A + // sculpt's side, which copies no corners, takes the room for its pinch + // test instead. + alignas(16) std::byte stack_room[STACK_COLUMNS * sizeof(S32)]; std::vector heap_lefts; - S32* lefts = stack_lefts; + S32* lefts = reinterpret_cast(stack_room); if (mNumS > STACK_COLUMNS) { heap_columns.resize(mNumS); @@ -8310,33 +8574,11 @@ bool LLVolumeFace::createSide(LLVolume* volume, bool partial_build) } else if (sculpt_stitching != LL_SCULPT_TYPE_NONE && !volume->isUnique() && !doubled && quads > 0 && rows > 0) { - // A sculpt's triangles two of whose corners are one point have - // no area, and are left out: the points' codes find them and the - // points confirm them. A flexible sculpt's points move after its - // full build, and it keeps them. A face keeps at least one - // triangle. The flat inner side of a sculpt whose parameters - // carry a square or triangle hole has its corners doubled, and - // keeps all of its quads. - ALScratchArena::Scope scratch; - U32* codes = scratch.allocate(num_vertices + 8); - alsimd::unit_cube_codes(pos, num_vertices, codes); - std::fill(codes + num_vertices, codes + num_vertices + 8, 0u); - const S32 words = (quads + 63) / 64; - U64* upper = scratch.allocate((size_t)words * rows); - U64* lower = scratch.allocate((size_t)words * rows); - S32 pinched = 0; - if (alsimd::grid_coincident_triangles(codes, mNumS, mNumT, upper, lower, words) > 0) - { - pinched = confirm_pinches(pos, mNumS, mNumT, upper, lower, words); - } - if (pinched > 0 && pinched < triangles) - { - end_index = grid_indices(mIndices, limit, mNumS, mNumT, lefts, quads, Unpinched{ upper, lower, words }); - } - else - { - end_index = grid_indices(mIndices, limit, mNumS, mNumT, lefts, quads, KeepAll()); - } + // A flexible sculpt's points move after its full build, and it + // keeps every triangle. The flat inner side of a sculpt whose + // parameters carry a square or triangle hole has its corners + // doubled, and keeps all of its quads. + end_index = sculpt_indices(mIndices, limit, *volume, pos, mNumS, mNumT, triangles, stack_room, sizeof(stack_room)); } else { diff --git a/indra/llmath/tests/alsimdkernels_test.cpp b/indra/llmath/tests/alsimdkernels_test.cpp index 10a879d6c6..d4d1bf9013 100644 --- a/indra/llmath/tests/alsimdkernels_test.cpp +++ b/indra/llmath/tests/alsimdkernels_test.cpp @@ -397,6 +397,38 @@ namespace std::uniform_int_distribution byte(0, 255); std::uniform_real_distribution near_cube(-0.7f, 0.7f); const U32 UNWRITTEN = 0xDEADBEEFu; + // Every byte on each axis, x mirrored and not: a code each, so two + // such points have one code only where they are one point. + { + std::vector points(512); + std::vector expected(512); + for (int v = 0; v < 256; ++v) + { + for (int mirrored = 0; mirrored < 2; ++mirrored) + { + const int r = v, g = 255 - v, b = (v * 7) & 255; + LLVector4a p((F32)r, (F32)g, (F32)b); + p.mul(1.f / 255.f); + p.sub(LLVector4a(0.5f, 0.5f, 0.5f)); + if (mirrored) + { + p.mul(LLVector4a(-1.f, 1.f, 1.f, 1.f)); + } + points[2 * v + mirrored] = p; + expected[2 * v + mirrored] = (U32)(mirrored ? 255 - r : r) | ((U32)g << 8) | ((U32)b << 16); + } + } + for_each_width([&](const char* width) + { + std::vector codes(points.size()); + unit_cube_codes_impl(points.data(), points.size(), codes.data()); + for (size_t i = 0; i < points.size(); ++i) + { + CAPTURE(width, i); + REQUIRE(codes[i] == expected[i]); + } + }); + } for (size_t n : COUNTS) { std::vector points(n + 1); @@ -450,19 +482,20 @@ namespace // A grid's triangles with two corners of one code, against every // pair of every quad's corners: codes from a few values so that many // corners match, rows of quads across one, two and three words, the - // flags past the row's last quad and in the words past it clear, and - // the rows counted. The codes past the grid are what the kernel may - // read and must not flag from. + // flags past the row's last quad and in the words past it clear, the + // triangles counted and the rows with any marked. The codes past the + // grid are what the kernel may read and must not flag from. TEST_CASE("alsimdkernels: grid_coincident_triangles flags every triangle with two corners of one code and no other", "[alsimdkernels]") { std::mt19937 rng(15); const size_t columns_tried[] = {2, 3, 4, 5, 8, 9, 16, 17, 33, 64, 65, 66, 70, 129, 130}; - const size_t rows_tried[] = {2, 3, 8}; + const size_t rows_tried[] = {2, 3, 8, 66, 130}; for (size_t columns : columns_tried) { for (size_t rows : rows_tried) { - for (int few : {3, 50}) + // Sparse flags, so that some rows have none, and dense ones. + for (int few : {3, 50, 5000}) { std::uniform_int_distribution value(0, few); std::vector codes(columns * rows + 8); @@ -472,11 +505,11 @@ namespace } const size_t quads = columns - 1; const size_t words = (columns + 62) / 64 + 1; - std::vector want_upper(words * rows, 0), want_lower(words * rows, 0); - size_t want_rows = 0; + const size_t row_words = (rows + 62) / 64; + std::vector want_upper(words * rows, 0), want_lower(words * rows, 0), want_rows(row_words, 0); + size_t want_flags = 0; for (size_t t = 0; t + 1 < rows; ++t) { - bool any = false; for (size_t s = 0; s < quads; ++s) { const U32 bl = codes[s + columns * t], br = codes[s + 1 + columns * t]; @@ -485,16 +518,25 @@ namespace const bool lower = bl == br || br == tr || bl == tr; want_upper[words * t + s / 64] |= U64(upper) << (s % 64); want_lower[words * t + s / 64] |= U64(lower) << (s % 64); - any = any || upper || lower; + want_rows[t / 64] |= U64(upper || lower) << (t % 64); + want_flags += (upper ? 1 : 0) + (lower ? 1 : 0); } - want_rows += any ? 1 : 0; } for_each_width([&](const char* width) { std::vector upper(words * rows, ~U64(0)), lower(words * rows, ~U64(0)); - const size_t got_rows = grid_coincident_triangles_impl(codes.data(), columns, rows, upper.data(), lower.data(), words); + // A word past the rows' is left as it is. + std::vector flagged_rows(row_words + 1, ~U64(0)); + const size_t got_flags = + grid_coincident_triangles_impl(codes.data(), columns, rows, upper.data(), lower.data(), words, flagged_rows.data()); CAPTURE(width, columns, rows, few); - CHECK(got_rows == want_rows); + CHECK(got_flags == want_flags); + for (size_t w = 0; w < row_words; ++w) + { + CAPTURE(w); + CHECK(flagged_rows[w] == want_rows[w]); + } + CHECK(flagged_rows[row_words] == ~U64(0)); for (size_t t = 0; t + 1 < rows; ++t) { for (size_t w = 0; w < words; ++w) diff --git a/indra/llmath/tests/llvolume_bench.cpp b/indra/llmath/tests/llvolume_bench.cpp index ee7e0436c6..f5f53c6009 100644 --- a/indra/llmath/tests/llvolume_bench.cpp +++ b/indra/llmath/tests/llvolume_bench.cpp @@ -162,6 +162,27 @@ namespace return map; } + // A 64 by 64 plane whose middle rows are one colour: the mesh rows + // sampled from them are one place, a thread, and their triangles have + // two corners at one point. + SculptMap pinchedSculptMap() + { + SculptMap map; + map.rgb.resize((size_t)map.width * map.height * 3); + for (U16 y = 0; y < map.height; ++y) + { + for (U16 x = 0; x < map.width; ++x) + { + const U16 row = (y >= 16 && y <= 48) ? 16 : y; + U8* out = &map.rgb[((size_t)y * map.width + x) * 3]; + out[0] = toByte((F32)x / (map.width - 1) - 0.5f); + out[1] = toByte((F32)row / (map.height - 1) - 0.5f); + out[2] = toByte(0.f); + } + } + return map; + } + LLPointer sculptVolume(U8 stitching, S32 lod) { LLVolumeParams p; @@ -640,6 +661,17 @@ namespace }; } } + const SculptMap pinched = pinchedSculptMap(); + for (S32 lod = 0; lod < LOD_COUNT; ++lod) + { + LLPointer volume = sculptVolume(LL_SCULPT_TYPE_PLANE, lod); + volume->sculpt(pinched.width, pinched.height, 3, pinched.rgb.data(), 0, false); + BENCHMARK(fmt::format("pinched plane sculpt, LOD {}, 64 by 64 map, {}", lod, counts(*volume))) + { + volume->sculpt(pinched.width, pinched.height, 3, pinched.rgb.data(), 0, false); + return volume->getNumVolumeFaces(); + }; + } } // createTangents sums meshoptimizer's default weighting a vertex at a diff --git a/indra/llmath/tests/llvolume_test.cpp b/indra/llmath/tests/llvolume_test.cpp index 385840607f..0b81b79957 100644 --- a/indra/llmath/tests/llvolume_test.cpp +++ b/indra/llmath/tests/llvolume_test.cpp @@ -3100,6 +3100,202 @@ namespace return sculptMap(32, 32, [](U16, U16) { return LLVector3(0.1f, 0.1f, 0.1f); }); } + // The quads of a num_s by num_t grid, as column and row, in the order a + // side lists them: a grid of up to seven rows a column at a time, every + // row of a column before the next; a taller one in bands, one of up to + // seven quads across or as many of at most six as near equal as the + // quads divide, the wider first, each band row by row and every other + // band from its last row back to its first. + std::vector> sideQuadOrder(S32 num_s, S32 num_t) + { + const S32 quads = num_s - 1; + const S32 rows = num_t - 1; + std::vector> order; + if (num_t <= 7) + { + for (S32 s = 0; s < quads; ++s) + { + for (S32 t = 0; t < rows; ++t) + { + order.push_back({ s, t }); + } + } + return order; + } + const S32 bands = quads <= 7 ? 1 : (quads + 5) / 6; + S32 first = 0; + for (S32 band = 0; band < bands; ++band) + { + const S32 width = quads / bands + (band < quads % bands ? 1 : 0); + for (S32 row = 0; row < rows; ++row) + { + const S32 t = (band & 1) ? rows - 1 - row : row; + for (S32 s = first; s < first + width; ++s) + { + order.push_back({ s, t }); + } + } + first += width; + } + return order; + } + + // A sculpt side's indices: its quads in the order above, each as bottom + // left, top right, top left and bottom left, bottom right, top right, + // and where leave_out is set, but for the triangles two of whose + // corners are one point, unless that would leave none. + std::vector sculptSideIndices(const LLVolumeFace& face, bool leave_out) + { + std::vector all, kept; + const LLVector4a* pos = face.mPositions; + for (const auto& [s, t] : sideQuadOrder(face.mNumS, face.mNumT)) + { + const U16 bl = (U16)gridIndex(face, s, t), br = (U16)(bl + 1); + const U16 tl = (U16)gridIndex(face, s, t + 1), tr = (U16)(tl + 1); + for (const std::array& tri : { std::array{ bl, tr, tl }, std::array{ bl, br, tr } }) + { + all.insert(all.end(), tri.begin(), tri.end()); + if (!samePoint(pos[tri[0]], pos[tri[1]]) && !samePoint(pos[tri[1]], pos[tri[2]]) && !samePoint(pos[tri[0]], pos[tri[2]])) + { + kept.insert(kept.end(), tri.begin(), tri.end()); + } + } + } + return leave_out && !kept.empty() ? kept : all; + } + + // A map of a few colours, scattered: samples meet everywhere, in no + // pattern. + SculptMap fewColoursMap(U32 colours) + { + return sculptMap(64, 64, [&](U16 x, U16 y) + { + U32 h = (U32(y) * 64u + x) * 0x9e3779b9u; + h ^= h >> 15; + const U32 c = h % colours; + return LLVector3((F32)c / colours - 0.5f, (F32)(c * 7 % colours) / colours - 0.5f, 0.25f); + }); + } + + // A plane drawn in steps, each value held for three columns and five + // rows: some samples meet their neighbours along a row, some down a + // column, depending on the detail. + SculptMap steppedMap() + { + return sculptMap(64, 64, [](U16 x, U16 y) + { + const F32 u = (F32)(x / 3 * 3) / 63, v = (F32)(y / 5 * 5) / 63; + return LLVector3(u - 0.5f, v - 0.5f, 0.1f * sinf(F_TWO_PI * u)); + }); + } + + // Every side of a sculpt lists its quads in the side order, a quad's two + // triangles in turn, and a shared volume's leave out the triangles two + // of whose corners are one point: the indices word for word, whatever + // the stitching, flags and detail, for maps whose samples meet along + // rows, columns and folds, at single points, in no pattern or not at + // all, for the placeholders, and for parameters with a hole or a cut. A + // flat inner side with its corners doubled lists its runs of quads + // apart, and keeps all of them; it is left out here. + TEST_CASE("llvolume: a sculpt lists its quads in order less the triangles with two corners at one point", "[llvolume]") + { + std::vector> maps; + const std::array sizes[] = { { 64, 64 }, { 128, 32 }, { 32, 128 }, { 64, 16 }, { 32, 32 }, { 8, 8 } }; + for (const auto& size : sizes) + { + maps.push_back({ fmt::format("own surface {} by {}", size[0], size[1]), {} }); + } + maps.push_back({ "pinched", pinchedMap() }); + maps.push_back({ "folded", foldedMap() }); + maps.push_back({ "pinwheel", pinwheelMap() }); + maps.push_back({ "stepped", steppedMap() }); + maps.push_back({ "three colours", fewColoursMap(3) }); + maps.push_back({ "seventeen colours", fewColoursMap(17) }); + maps.push_back({ "one colour", oneColourMap() }); + maps.push_back({ "no map, sphere placeholder", SculptMap() }); + maps.push_back({ "no map, empty placeholder", SculptMap() }); + + struct Shape + { + const char* name; + U8 profile; + F32 hollow; + F32 begin; + }; + const Shape shapes[] = { + { "", LL_PCODE_PROFILE_CIRCLE, 0.f, 0.f }, + { ", circle hole", LL_PCODE_PROFILE_CIRCLE | LL_PCODE_HOLE_CIRCLE, 0.5f, 0.f }, + { ", square hole", LL_PCODE_PROFILE_CIRCLE | LL_PCODE_HOLE_SQUARE, 0.4f, 0.f }, + { ", profile cut", LL_PCODE_PROFILE_CIRCLE, 0.f, 0.2f }, + }; + const U8 flags[] = { 0, LL_SCULPT_FLAG_MIRROR, LL_SCULPT_FLAG_INVERT, LL_SCULPT_FLAG_MIRROR | LL_SCULPT_FLAG_INVERT }; + + Report report; + S32 sides = 0; + S32 sides_left_out = 0; + for (U8 stitching : STITCHINGS) + { + for (size_t m = 0; m < maps.size(); ++m) + { + const SculptMap map = m < std::size(sizes) ? mapFor(stitching, sizes[m][0], sizes[m][1]) : maps[m].second; + const bool visible = maps[m].first != "no map, empty placeholder"; + for (const Shape& shape : shapes) + { + // A hole or a cut on the stitchings' own surfaces and the + // stepped map. + if (shape.name[0] && m >= std::size(sizes) && maps[m].first != "stepped") + { + continue; + } + for (U8 flag : flags) + { + for (S32 lod = 0; lod < LOD_COUNT; ++lod) + { + for (bool unique : { false, true }) + { + if (unique && flag != 0) + { + continue; + } + LLVolumeParams p; + p.setType(shape.profile, LL_PCODE_PATH_CIRCLE); + p.setRatio(1.f, 0.5f); + p.setHollow(shape.hollow); + p.setBeginAndEndS(shape.begin, 1.f); + p.setSculptID(LLUUID("c8b26d0b-6f5a-4f10-8d1e-5c3b3e4a9a01"), (U8)(stitching | flag)); + LLPointer volume = new LLVolume(p, LOD_DETAIL[lod], false, unique); + volume->sculpt(map.width, map.height, 3, map.rgb.empty() ? nullptr : map.rgb.data(), 0, visible); + for (S32 f = 0; f < volume->getNumVolumeFaces(); ++f) + { + const LLVolumeFace& face = volume->getVolumeFace(f); + if (!isGrid(face) || isDoubled(face)) + { + continue; + } + const std::vector want = sculptSideIndices(face, !unique); + const std::vector got(face.mIndices, face.mIndices + face.mNumIndices); + ++sides; + sides_left_out += want.size() < 6u * (face.mNumS - 1) * (face.mNumT - 1) ? 1 : 0; + if (got != want) + { + const size_t at = std::mismatch(got.begin(), got.end(), want.begin(), want.end()).first - got.begin(); + report.fail(fmt::format("{} sculpt, {}{}, flags {}, LOD {}{}, face {}: {} indices where {} are listed, first apart at {}", + stitchingName(stitching), maps[m].first, shape.name, flag, lod, unique ? ", unique" : "", f, + got.size(), want.size(), at)); + } + } + } + } + } + } + } + } + INFO(fmt::format("{} sides, {} with triangles left out", sides, sides_left_out)); + INFO(report.text()); + CHECK(sides_left_out > 0); + CHECK(report.failures == 0); + } + // The sphere a sculpt shows while its map is missing or has no area, // whatever its stitching, and the empty one it shows otherwise. void inspectPlaceholders(const Inspect& how, bool spheres_only, Report& report)