return 1.0 - (D - B) / t;
}
Here:
- empirically 4 tiles of 2048 are enough
- min over measurements is less noisy than mean
- We'll run the first few tiles with calibration
- measure bsl, because it always computes `f`, so `T = t_bsl - B`
- std::max here protects against divide-by-zero and against `t < 0` when `T` is very cheap
### pilot v1
When `d_max < 0` bsl is always faster:
void pilot_v1(float* dst, const size_t n, auto f, auto mask) {
const auto w = vld1q_u32(weights.data());
const float d_max = bsl_calibrate<tile / 2>(dst, 2 * tile, f, mask);
dst += 2 * tile;
for (size_t i = 2 * tile; i < n; i += tile) {
if (d_max < 0)
bsl<false>(dst, tile, f, mask);
else
detour<false>(dst, tile, w, f, mask);
dst += tile;
}
}
| thd | tiled detour sqrt | pilot v1 sqrt | BSL sqrt | tiled detour sin35 | pilot v1 sin35 | BSL sin35 |
|----:|------------------:|--------------:|---------:|-------------------:|---------------:|----------:|
| 0 | 38.37\* | 32.19 | 32.24 | 38.26\* | 38.16 | 9.75 |
| 0.3 | 16.89 | 32.18 | 32.24\* | 12.72\* | 12.71 | 9.75 |
| 0.6 | 14.54 | 32.18\* | 32.18\* | 9.25 | 9.36 | 9.75\* |
| 1 | 12.29 | 32.21 | 32.25\* | 6.87 | 6.86 | 9.75\* |
The algorithm got sqrt right.
But for sin35 at high thd, detour is selected, and we lose 30%:
v1 switches to bsl only when it's faster at every density.
### pilot v2
It's expensive to calculate the density of the whole tile, so we'll use the first 256
(It reads 6% of the tile, which is noise on pow, but noticeable on sqrt)
For a more or less uniform distribution this is enough:
size_t density(float* dst, const size_t n, auto mask) {
std::array<uint32x4_t, 4> acc;
acc.fill(vdupq_n_u32(0));
for (size_t i = 0; i < n; i += 16) {
for (size_t j = 0; j < 4; ++j)
acc[j] = vsubq_u32(acc[j], mask(vld1q_f32(dst + i + 4 * j)));
}
return vaddvq_u32(vaddq_u32(vaddq_u32(acc[0], acc[1]), vaddq_u32(acc[2], acc[3])));
}
Trick here: true lane of bitmask = 0xFFFFFFFF = -1, subtracting the lane actually adds.
Skip wins when the predictor rarely mispredicts, i.e. an 80% chance that all 4 registers are empty.
The density is approximately `0.014` (`(1 - x)^16 = 0.8`)
void pilot_v2(float* dst, const size_t n, auto f, auto mask) {
const auto w = vld1q_u32(weights.data());
const float d_max = bsl_calibrate<tile / 2>(dst, 2 * tile, f, mask);
constexpr size_t probe = 256;
constexpr size_t xlo = 0.014 * probe;
const long long hi = d_max * probe;
dst += 2 * tile;
for (size_t i = 2 * tile; i < n; i += tile) {
const long long cnt = density(dst, probe, mask);
if (cnt > hi) {
if (cnt < xlo)
bsl<true>(dst, tile, f, mask);
else
bsl<false>(dst, tile, f, mask);
} else if (cnt < xlo)
detour<true>(dst, tile, w, f, mask);
else
detour<false>(dst, tile, w, f, mask);
dst += tile;
}
}
`hi` and `xlo` here are `d_max` and `0.014` cutoffs, but multiplied by the probe length (256).
| thd | 0 | 0.1 | 0.2 | 0.3 | 0.4 | 0.5 | 0.6 | 0.7 | 0.8 | 0.9 | 1 |
|:---------------|--------:|--------:|--------:|--------:|--------:|--------:|-------:|-------:|-------:|-------:|-------:|
| pilot_v1 sin35 | 38.25 | 16.72\* | 14.48\* | 12.77\* | 11.41\* | 10.33\* | 9.40 | 8.62 | 7.93 | 7.36 | 6.91 |
| pilot_v2 sin35 | 77.29\* | 16.55 | 14.34 | 12.65 | 11.33 | 10.14 | 9.71 | 9.73 | 9.74 | 9.73 | 9.73 |
| BSL sin35 | 9.78 | 9.77 | 9.77 | 9.76 | 9.76 | 9.78 | 9.78\* | 9.77\* | 9.78\* | 9.77\* | 9.78\* |
|
Post #25680
4