Merge branch 'branch-23.12' into imp-2312-output_aarch_yml

rapidsai · Oct 19, 2023 · d9eb308 · d9eb308
2 parents 71e6de3 + 747f6a6
commit d9eb308
Show file tree

Hide file tree

Showing 16 changed files with 115 additions and 108 deletions.
diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml
@@ -28,7 +28,7 @@ concurrency:
 jobs:
   cpp-build:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
@@ -37,7 +37,7 @@ jobs:
   python-build:
     needs: [cpp-build]
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
@@ -46,7 +46,7 @@ jobs:
   upload-conda:
     needs: [cpp-build, python-build]
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
@@ -57,7 +57,7 @@ jobs:
     if: github.ref_type == 'branch'
     needs: python-build
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       arch: "amd64"
       branch: ${{ inputs.branch }}
@@ -69,7 +69,7 @@ jobs:
       sha: ${{ inputs.sha }}
   wheel-build-pylibraft:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
@@ -79,7 +79,7 @@ jobs:
   wheel-publish-pylibraft:
     needs: wheel-build-pylibraft
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
@@ -89,7 +89,7 @@ jobs:
   wheel-build-raft-dask:
     needs: wheel-publish-pylibraft
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
@@ -99,7 +99,7 @@ jobs:
   wheel-publish-raft-dask:
     needs: wheel-build-raft-dask
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}

diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml
@@ -24,41 +24,41 @@ jobs:
       - wheel-tests-raft-dask
       - devcontainer
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
   checks:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       enable_check_generated_files: false
   conda-cpp-build:
     needs: checks
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
       node_type: cpu16
   conda-cpp-tests:
     needs: conda-cpp-build
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
   conda-python-build:
     needs: conda-cpp-build
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
   conda-python-tests:
     needs: conda-python-build
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
   docs-build:
     needs: conda-python-build
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
       node_type: "gpu-v100-latest-1"
@@ -68,34 +68,34 @@ jobs:
   wheel-build-pylibraft:
     needs: checks
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
       script: ci/build_wheel_pylibraft.sh
   wheel-tests-pylibraft:
     needs: wheel-build-pylibraft
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
       script: ci/test_wheel_pylibraft.sh
   wheel-build-raft-dask:
     needs: wheel-tests-pylibraft
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
       script: "ci/build_wheel_raft_dask.sh"
   wheel-tests-raft-dask:
     needs: wheel-build-raft-dask
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: pull-request
       script: ci/test_wheel_raft_dask.sh
   devcontainer:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_command: |
         sccache -z;

diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml
@@ -16,23 +16,23 @@ on:
 jobs:
   conda-cpp-tests:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: nightly
       branch: ${{ inputs.branch }}
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
   conda-python-tests:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: nightly
       branch: ${{ inputs.branch }}
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
   wheel-tests-pylibraft:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: nightly
       branch: ${{ inputs.branch }}
@@ -41,7 +41,7 @@ jobs:
       script: ci/test_wheel_pylibraft.sh
   wheel-tests-raft-dask:
     secrets: inherit
-    uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
+    uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       build_type: nightly
       branch: ${{ inputs.branch }}

diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh
@@ -82,7 +82,7 @@ for FILE in .github/workflows/*.yaml; do
 done
 
 for FILE in .github/workflows/*.yaml; do
-  sed_runner "/shared-action-workflows/ s/@.*/@branch-${NEXT_SHORT_TAG}/g" "${FILE}"
+  sed_runner "/shared-workflows/ s/@.*/@branch-${NEXT_SHORT_TAG}/g" "${FILE}"
 done
 sed_runner "s/RAPIDS_VERSION_NUMBER=\".*/RAPIDS_VERSION_NUMBER=\"${NEXT_SHORT_TAG}\"/g" ci/build_docs.sh
 

diff --git a/ci/test_cpp.sh b/ci/test_cpp.sh
@@ -36,6 +36,7 @@ trap "EXITCODE=1" ERR
 set +e
 
 # Run libraft gtests from libraft-tests package
+cd "$CONDA_PREFIX"/bin/gtests/libraft
 ctest -j8 --output-on-failure
 
 rapids-logger "Test script exiting with value: $EXITCODE"

diff --git a/cpp/include/raft/distance/detail/distance_ops/l2_exp.cuh b/cpp/include/raft/distance/detail/distance_ops/l2_exp.cuh
@@ -21,6 +21,22 @@
 
 namespace raft::distance::detail::ops {
 
+/**
+ * Reserve 1 digit of precision from each floating-point type
+ * for round-off error tolerance.
+ * @tparam DataT
+ */
+template <typename DataT>
+__device__ constexpr DataT get_clamp_precision()
+{
+  switch (sizeof(DataT)) {
+    case 2: return 1e-3;
+    case 4: return 1e-6;
+    case 8: return 1e-15;
+    default: return 0;
+  }
+}
+
 // Epilogue operator for CUTLASS based kernel
 template <typename DataT, typename AccT>
 struct l2_exp_cutlass_op {
@@ -31,11 +47,13 @@ struct l2_exp_cutlass_op {
   __device__ AccT operator()(DataT& aNorm, const DataT& bNorm, DataT& accVal) const noexcept
   {
     AccT outVal = aNorm + bNorm - DataT(2.0) * accVal;
-    // outVal could be negative due to numerical instability, especially when
-    // calculating self distance.
-    // clamp to 0 to avoid potential NaN in sqrt
-    outVal = outVal * (raft::abs(outVal) >= DataT(0.0001));
-    return sqrt ? raft::sqrt(outVal) : outVal;
+
+    /**
+     * 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 * !((outVal * outVal < get_clamp_precision<DataT>()) * (aNorm == bNorm));
+    return sqrt ? raft::sqrt(outVal * (outVal > 0)) : outVal;
   }
 
   __device__ AccT operator()(DataT aData) const noexcept { return aData; }
@@ -86,10 +104,16 @@ struct l2_exp_distance_op {
     for (int i = 0; i < Policy::AccRowsPerTh; ++i) {
 #pragma unroll
       for (int j = 0; j < Policy::AccColsPerTh; ++j) {
-        DataT val = regxn[i] + regyn[j] - (DataT)2.0 * acc[i][j];
-        // val could be negative due to numerical instability, especially when
-        // calculating self distance. Clamp to 0 to avoid potential NaN in sqrt
-        acc[i][j] = val * (raft::abs(val) >= DataT(0.0001));
+        DataT accVal = acc[i][j];
+        DataT val    = regxn[i] + regyn[j] - (DataT)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 * (val > 0) * !((val * val < get_clamp_precision<DataT>()) * (regxn[i] == regyn[j]));
       }
     }
     if (sqrt) {

diff --git a/cpp/include/raft/neighbors/detail/knn_brute_force.cuh b/cpp/include/raft/neighbors/detail/knn_brute_force.cuh
@@ -29,6 +29,7 @@
 #include <cstdint>
 #include <iostream>
 #include <raft/core/resources.hpp>
+#include <raft/distance/detail/distance_ops/l2_exp.cuh>
 #include <raft/distance/distance.cuh>
 #include <raft/distance/distance_types.hpp>
 #include <raft/linalg/map.cuh>
@@ -186,6 +187,7 @@ void tiled_brute_force_knn(const raft::resources& handle,
         auto row_norms = search_norms.data();
         auto col_norms = precomputed_index_norms ? precomputed_index_norms : index_norms.data();
         auto dist      = temp_distances.data();
+        bool sqrt      = metric == raft::distance::DistanceType::L2SqrtExpanded;
 
         raft::linalg::map_offset(
           handle,
@@ -194,15 +196,9 @@ void tiled_brute_force_knn(const raft::resources& handle,
             IndexType row = i + (idx / current_centroid_size);
             IndexType col = j + (idx % current_centroid_size);
 
-            auto val = row_norms[row] + col_norms[col] - 2.0 * dist[idx];
-
-            // due to numerical instability (especially around self-distance)
-            // the distances here could be slightly negative, which will
-            // cause NaN values in the subsequent sqrt. Clamp to 0
-            val = val * (val >= 0.0001);
-            if (metric == raft::distance::DistanceType::L2SqrtExpanded) { val = sqrt(val); }
-            val = distance_epilogue(val, row, col);
-            return val;
+            raft::distance::detail::ops::l2_exp_cutlass_op<ElementType, ElementType> l2_op(sqrt);
+            auto val = l2_op(row_norms[row], col_norms[col], dist[idx]);
+            return distance_epilogue(val, row, col);
           });
       } else if (metric == raft::distance::DistanceType::CosineExpanded) {
         auto row_norms = search_norms.data();

diff --git a/cpp/test/distance/fused_l2_nn.cu b/cpp/test/distance/fused_l2_nn.cu
@@ -57,6 +57,7 @@ RAFT_KERNEL naiveKernel(raft::KeyValuePair<int, DataT>* min,
     auto diff = midx >= m || nidx >= n ? DataT(0) : x[xidx] - y[yidx];
     acc += diff * diff;
   }
+
   if (Sqrt) { acc = raft::sqrt(acc); }
   ReduceOpT redOp;
   typedef cub::WarpReduce<raft::KeyValuePair<int, DataT>> WarpReduce;
@@ -343,7 +344,7 @@ const std::vector<Inputs<double>> inputsd = {
   {0.00001, 128, 32, 33, 1234ULL},  {0.00001, 128, 64, 33, 1234ULL},
   {0.00001, 128, 128, 65, 1234ULL}, {0.00001, 64, 128, 129, 1234ULL},
 
-  {0.00001, 1805, 134, 2, 1234ULL}, {0.00001, 8192, 1024, 25, 1234ULL},
+  {0.00001, 1805, 134, 2, 1234ULL},  //{0.00001, 8192, 1024, 25, 1234ULL},
 };
 typedef FusedL2NNTest<double, false> FusedL2NNTestD_Sq;
 TEST_P(FusedL2NNTestD_Sq, Result)

diff --git a/cpp/test/neighbors/ann_cagra.cuh b/cpp/test/neighbors/ann_cagra.cuh
@@ -274,7 +274,7 @@ class AnnCagraTest : public ::testing::TestWithParam<AnnCagraInputs> {
                                   distances_Cagra,
                                   ps.n_queries,
                                   ps.k,
-                                  0.001,
+                                  0.003,
                                   min_recall));
       EXPECT_TRUE(eval_distances(handle_,
                                  database.data(),
@@ -515,7 +515,7 @@ class AnnCagraFilterTest : public ::testing::TestWithParam<AnnCagraInputs> {
                                   distances_Cagra,
                                   ps.n_queries,
                                   ps.k,
-                                  0.001,
+                                  0.003,
                                   min_recall));
       EXPECT_TRUE(eval_distances(handle_,
                                  database.data(),
@@ -628,7 +628,7 @@ class AnnCagraFilterTest : public ::testing::TestWithParam<AnnCagraInputs> {
                                   distances_Cagra,
                                   ps.n_queries,
                                   ps.k,
-                                  0.001,
+                                  0.003,
                                   min_recall));
       EXPECT_TRUE(eval_distances(handle_,
                                  database.data(),

diff --git a/cpp/test/neighbors/ann_ivf_pq.cuh b/cpp/test/neighbors/ann_ivf_pq.cuh
@@ -312,7 +312,7 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
     // Hence, encoding-decoding chain often leads to altering both the PQ codes and the
     // reconstructed data.
     compare_vectors_l2(
-      handle_, vectors_1.view(), vectors_2.view(), label, compression_ratio, 0.025);
+      handle_, vectors_1.view(), vectors_2.view(), label, compression_ratio, 0.04);  // 0.025);
   }
 
   void check_packing(index<IdxT>* index, uint32_t label)