Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
724 changes: 6 additions & 718 deletions .github/workflows/pr.yaml

Large diffs are not rendered by default.

55 changes: 45 additions & 10 deletions ci/test_cpp.sh
Original file line number Diff line number Diff line change
Expand Up @@ -52,23 +52,58 @@ nvidia-smi
# RAPIDS_DATASET_ROOT_DIR is used by test scripts
RAPIDS_DATASET_ROOT_DIR=${RAPIDS_TESTS_DIR}/dataset
export RAPIDS_DATASET_ROOT_DIR
./ci/get_test_data.sh --NEIGHBORS_ANN_VAMANA_TEST
# skipped for the reproducer: ./ci/get_test_data.sh --NEIGHBORS_ANN_VAMANA_TEST

EXITCODE=0
trap "EXITCODE=1" ERR
set +e

# Run Python build utilities tests (once, in the first shard)
if [[ "${SHARD}" == "1" ]]; then
rapids-logger "Run libcuvs Python build utilities tests"
pytest cpp/tests/python
fi

# Run libcuvs gtests from libcuvs-tests package
rapids-logger "Run libcuvs tests (shard ${SHARD} of ${NUM_SHARDS})"
# Flaky spectral clustering reproducer (DO NOT MERGE)
echo "Ignoring shard ${SHARD} of ${NUM_SHARDS}"
pushd "$CONDA_PREFIX"/bin/gtests/libcuvs
timeout -v --signal=SIGINT --kill-after=60s 100m ctest -j8 --output-on-failure -I "${SHARD},,${NUM_SHARDS}"
LOGDIR="${RAPIDS_TESTS_DIR}/spectral"
mkdir -p "${LOGDIR}"

rapids-logger "Full CLUSTER_TEST binary in fresh processes"
FULL_RUNS=${FULL_RUNS:-5}
full_fail=0
for i in $(seq 1 "${FULL_RUNS}"); do
if ! ./CLUSTER_TEST --gtest_filter='-SpectralClusteringDiag*' > "${LOGDIR}/full_${i}.log" 2>&1; then
full_fail=$((full_fail + 1))
grep -E "FAILED|Score|Failure|actual|Expected|Actual|eigensolver" "${LOGDIR}/full_${i}.log" || true
fi
done
echo "SUMMARY full CLUSTER_TEST: ${full_fail}/${FULL_RUNS} runs failed"

rapids-logger "Spectral clustering diagnostics"
SPECTRAL_DIAG_REPEATS=${SPECTRAL_DIAG_REPEATS:-300} ./CLUSTER_TEST --gtest_filter='SpectralClusteringDiag*' 2>&1 | tee "${LOGDIR}/diag.log" | grep -v "graph diff" || true

rapids-logger "Spectral clustering tests repeated in one process"
./CLUSTER_TEST --gtest_filter='SpectralClusteringTest*' --gtest_repeat=300 --gtest_brief=1 > "${LOGDIR}/repeat.log" 2>&1 || true
grep -E "FAILED|Score|actual" "${LOGDIR}/repeat.log" | sort | uniq -c | sort -rn | head -50 || true
echo "SUMMARY in-process repeat: $(grep -c '^\[ FAILED \] .*ms)$' "${LOGDIR}/repeat.log" || true) failed test instances (300 repeats x 22 tests)"

rapids-logger "SpectralClusteringTestF.Result/7 in fresh processes"
FRESH_RUNS=${FRESH_RUNS:-200}
fresh_fail=0
for i in $(seq 1 "${FRESH_RUNS}"); do
if ! ./CLUSTER_TEST --gtest_filter='SpectralClusteringTests/SpectralClusteringTestF.Result/7' --gtest_brief=1 > "${LOGDIR}/fresh_${i}.log" 2>&1; then
fresh_fail=$((fresh_fail + 1))
grep -E "Score|actual|eigensolver" "${LOGDIR}/fresh_${i}.log" || true
fi
done
echo "SUMMARY fresh-process Result/7: ${fresh_fail}/${FRESH_RUNS} runs failed"

rapids-logger "Distance and 1-NN tests"
other_fail=0
./DISTANCE_TEST --gtest_brief=1 || other_fail=$((other_fail + 1))
./NEIGHBORS_TEST --gtest_filter='NNTest*:Top1nn*' --gtest_brief=1 || other_fail=$((other_fail + 1))
echo "SUMMARY distance / 1-NN test binaries failed: ${other_fail}"
popd

if [[ ${full_fail} -gt 0 || ${fresh_fail} -gt 0 || ${other_fail} -gt 0 ]] || grep -q '^\[ FAILED \]' "${LOGDIR}/repeat.log"; then
EXITCODE=1
fi

rapids-logger "Test script exiting with value: $EXITCODE"
exit ${EXITCODE}
39 changes: 22 additions & 17 deletions cpp/src/distance/detail/distance_ops/l2_exp.cuh
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

Expand Down Expand Up @@ -28,6 +28,23 @@ __device__ constexpr AccT get_clamp_precision()
}
}

/**
* Whether an expanded L2 distance is round-off error of a self-distance and should be clamped to
* zero. Self-neighboring points have (aNorm == bNorm) == accVal, but the dot product (accVal) can
* have round-off errors, which will cause (aNorm == bNorm) ~ accVal instead.
*
* The round-off error scales with the norms, so the tolerance is relative to them for squared
* norms below one. Otherwise, distinct points with equal norms and a small magnitude (for example,
* a spectral embedding) would have their distances incorrectly clamped to zero.
*/
template <typename DataT, typename AccT>
DI bool is_self_distance_round_off(AccT dist, AccT aNorm, AccT bNorm)
{
const AccT norm_product = aNorm * bNorm;
const AccT norm_scale = norm_product < AccT(1) ? norm_product : AccT(1);
return (aNorm == bNorm) && (dist * dist < get_clamp_precision<DataT, AccT>() * norm_scale);
}

// Epilogue operator for CUTLASS based kernel
template <typename DataT, typename AccT>
struct l2_exp_cutlass_op {
Expand All @@ -38,13 +55,7 @@ struct l2_exp_cutlass_op {
inline __device__ AccT operator()(AccT aNorm, AccT bNorm, AccT accVal) const noexcept
{
AccT outVal = aNorm + bNorm - AccT(2.0) * accVal;

/**
* Self-neighboring points should have (aNorm == bNorm) == accVal and the dot product (accVal)
* can sometimes have round-off errors, which will cause (aNorm == bNorm) ~ accVal instead.
*/
outVal =
outVal * AccT(!((outVal * outVal < get_clamp_precision<DataT, AccT>()) * (aNorm == bNorm)));
outVal = outVal * AccT(!is_self_distance_round_off<DataT, AccT>(outVal, aNorm, bNorm));
return sqrt ? raft::sqrt(outVal * static_cast<AccT>(outVal > AccT(0))) : outVal;
}

Expand Down Expand Up @@ -112,15 +123,9 @@ struct l2_exp_distance_op {
for (int j = 0; j < Policy::AccColsPerTh; ++j) {
AccT accVal = acc[i][j];
AccT val = regxn[i] + regyn[j] - (AccT)2.0 * accVal;

/**
* Self-neighboring points should have (aNorm == bNorm) == accVal and the dot product
* (accVal) can sometimes have round-off errors, which will cause (aNorm == bNorm) ~ accVal
* instead.
*/
acc[i][j] = val * static_cast<AccT>((val > AccT(0))) *
static_cast<AccT>(
!((val * val < get_clamp_precision<DataT, AccT>()) * (regxn[i] == regyn[j])));
acc[i][j] =
val * static_cast<AccT>((val > AccT(0))) *
static_cast<AccT>(!is_self_distance_round_off<DataT, AccT>(val, regxn[i], regyn[j]));
}
}
if (sqrt) {
Expand Down
1 change: 1 addition & 0 deletions cpp/tests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -141,6 +141,7 @@ ConfigureTest(
cluster/linkage.cu
cluster/connect_knn.cu
cluster/spectral.cu
cluster/spectral_diag.cu
GPUS 1
PERCENT 100
)
Expand Down
Loading
Loading