but there's one big problem we've overlooked: we're assuming the data is uniform.
So it's easy to build a test where v3 will fail:

for (size_t i = 0; i < n; i++)
dst[i] = i % 4096 >= 256;

In this case, v3 always prefers BSL, even though detour wins on 3840 elements of the tile.

### pilot v3.5

According to the first table, in the worst case detour is under 3x slower (it happens on sqrt thd = 1), but on pow,
thd = 0, detour is 37x faster.
So a wrong BSL costs much more than a wrong detour.

For bsl, we'll play it safe by running it in tile/8 blocks and checking the density. If it drops well below `d_max`,
we'll switch to detour.
It costs 1 instruction per register. `acc = vsubq_u32(acc, m)` works because the mask is 0/-1. And unlike the probe, the
density here is exact.

size_t bsl_verified(float* dst, const size_t n, float d_max, auto f, auto mask) {
for (size_t i0 = 0; i0 < 8; ++i0) {
std::array<uint32x4_t, 4> acc;
acc.fill(vdupq_n_u32(0));

for (size_t i = 0; i < n / 8; i += 16) {
std::array<float32x4_t, 4> v;
for (size_t j = 0; j < 4; ++j) v[j] = vld1q_f32(dst + 4 * j);
std::array<uint32x4_t, 4> m;
for (size_t j = 0; j < 4; ++j) m[j] = mask(v[j]);
for (size_t j = 0; j < 4; ++j) acc[j] = vsubq_u32(acc[j], m[j]);
for (size_t j = 0; j < 4; ++j)
vst1q_f32(dst + 4 * j, vbslq_f32(m[j], f(v[j]), v[j]));

dst += 16;
}
const size_t cur = vaddvq_u32(vaddq_u32(vaddq_u32(acc[0], acc[1]), vaddq_u32(acc[2], acc[3])));
if (static_cast<double>(cur) / (n / 8) < 0.75 * d_max) return (i0 + 1) * n / 8;
}
return n;
}

BSL bails out when the density is less than `0.75 * d_max`. I have no math behind the 0.75, it just won on average.

And pilot v3.5 will use bsl_verified if the density > hi:

void pilot_v3_5(float* dst, const size_t n, auto f, auto mask) {
// ... same as v3

for (size_t i = 2 * tile; i < n; i += tile) {
const long long cnt = density(dst, probe, mask);
if (cnt > hi) {
if (hi < 0) {
if (cnt < xlo)
bsl<true>(dst, tile, f, mask);
else
bsl<false>(dst, tile, f, mask);
} else {
auto done = bsl_verified(dst, tile, d_max, f, mask);
if (done < tile)
detour<false>(dst + done, tile - done, w, f, mask);
}
}
// ... same as v3
}
}

But here too it's easy to build a countertest:

for (size_t i = 0; i < n; ++i)
dst[i] = i % 4096 >= 390;

The first block will pass the check (> 75% zeros in it), but the second won't. BSL runs on it for nothing, and for pow
that's expensive. v3.5's problem: it has no memory.
After a miss the algo keeps trusting the first 256 and misses on every tile.

### pilot v4

v4 will fix this: if bsl bails out, we stop trusting the probe for the next 16 tiles, and instead we take the density of
the previous tile:

void pilot_v4(float* dst, const size_t n, auto f, auto mask) {
// ... same as v3

size_t distrust = 0;
size_t prev = 0;
for (size_t i = 2 * tile; i < n; i += tile) {
if (distrust) --distrust;

const long long cnt = distrust ? prev : density(dst, probe, mask);
if (cnt > hi) {
if (hi < 0) {
if (cnt < xlo) {
bsl<true>(dst, tile, f, mask);
} else {
bsl<false>(dst, tile, f, mask);
}
} else {
if (distrust == 0) {
auto done = bsl_verified(dst, tile, d_max, f, mask);
if (done < tile) {
distrust = 16;
const size_t rem =