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)