IVF-Flat sample filter is not applied in all required places

[achirkin]

It seems like the filtering in ivf-flat is not entirely correct, it's only applied to one if case out of the few.
Perhaps, it makes sense to shift the condition further below to the place where the found value is added to the queue.

raft/cpp/include/raft/neighbors/detail/ivf_flat_interleaved_scan-inl.cuh

Lines 740 to 768 in 1034a41

	const bool valid = vec_id < list_length;
	// Process first shm_assisted_dim dimensions (always using shared memory)
	if (valid && sample_filter(queries_offset + blockIdx.y, probe_id, vec_id)) {
	loadAndComputeDist<kUnroll, decltype(compute_dist), Veclen, T, AccT> lc(dist,
	compute_dist);
	for (int pos = 0; pos < shm_assisted_dim;
	pos += WarpSize, data += kIndexGroupSize * WarpSize) {
	lc.runLoadShmemCompute(data, query_shared, lane_id, pos);
	}
	}
	if (dim > query_smem_elems) {
	// The default path - using shfl ops - for dimensions beyond query_smem_elems
	loadAndComputeDist<kUnroll, decltype(compute_dist), Veclen, T, AccT> lc(dist,
	compute_dist);
	for (int pos = shm_assisted_dim; pos < full_warps_along_dim; pos += WarpSize) {
	lc.runLoadShflAndCompute(data, query, pos, lane_id);
	}
	lc.runLoadShflAndComputeRemainder(data, query, lane_id, dim, full_warps_along_dim);
	} else {
	// when shm_assisted_dim == full_warps_along_dim < dim
	if (valid) {
	loadAndComputeDist<1, decltype(compute_dist), Veclen, T, AccT> lc(dist, compute_dist);
	for (int pos = full_warps_along_dim; pos < dim;
	pos += Veclen, data += kIndexGroupSize * Veclen) {
	lc.runLoadShmemCompute(data, query_shared, lane_id, pos);
	}
	}

[lowener]

The call to sample_filter is not right.
It should be sample_filter(queries_offset + blockIdx.y, list_id, vec_id).

[lowener]

True, the filtering was still adding neighbors that didn't pass the test to the queue. I am fixing this in PR #1831