From 7dab7e2fd17c268d17df3a879a3f2f18291b4dd7 Mon Sep 17 00:00:00 2001 From: Allard Hendriksen Date: Thu, 7 Jul 2022 14:56:17 +0200 Subject: [PATCH] dbscan: remove naive versions of algorithms This fixes issue #3414. --- cpp/src/dbscan/adjgraph/algo.cuh | 1 - cpp/src/dbscan/adjgraph/naive.cuh | 65 -------------------- cpp/src/dbscan/adjgraph/runner.cuh | 4 +- cpp/src/dbscan/common.cuh | 43 ------------- cpp/src/dbscan/dbscan.cuh | 2 + cpp/src/dbscan/vertexdeg/naive.cuh | 93 ----------------------------- cpp/src/dbscan/vertexdeg/runner.cuh | 4 +- 7 files changed, 4 insertions(+), 208 deletions(-) delete mode 100644 cpp/src/dbscan/adjgraph/naive.cuh delete mode 100644 cpp/src/dbscan/common.cuh delete mode 100644 cpp/src/dbscan/vertexdeg/naive.cuh diff --git a/cpp/src/dbscan/adjgraph/algo.cuh b/cpp/src/dbscan/adjgraph/algo.cuh index 3918db5b23..759297d0e8 100644 --- a/cpp/src/dbscan/adjgraph/algo.cuh +++ b/cpp/src/dbscan/adjgraph/algo.cuh @@ -20,7 +20,6 @@ #include #include -#include "../common.cuh" #include "pack.h" #include diff --git a/cpp/src/dbscan/adjgraph/naive.cuh b/cpp/src/dbscan/adjgraph/naive.cuh deleted file mode 100644 index 6ce3994610..0000000000 --- a/cpp/src/dbscan/adjgraph/naive.cuh +++ /dev/null @@ -1,65 +0,0 @@ -/* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "../common.cuh" -#include "pack.h" -#include -#include -#include -#include - -namespace ML { -namespace Dbscan { -namespace AdjGraph { -namespace Naive { - -template -void launcher(const raft::handle_t& handle, - Pack data, - Index_ batch_size, - cudaStream_t stream) -{ - Index_ k = 0; - Index_ N = data.N; - ML::pinned_host_vector host_vd(batch_size + 1); - ML::pinned_host_vector host_adj(((batch_size * N) / 8) + 1); - ML::pinned_host_vector host_ex_scan(batch_size); - raft::update_host((bool*)host_adj.data(), data.adj, batch_size * N, stream); - raft::update_host(host_vd.data(), data.vd, batch_size + 1, stream); - handle.sync_stream(stream); - size_t adjgraph_size = size_t(host_vd[batch_size]); - ML::pinned_host_vector host_adj_graph(adjgraph_size); - for (Index_ i = 0; i < batch_size; i++) { - for (Index_ j = 0; j < N; j++) { - /// TODO: change layout or remove; cf #3414 - if (host_adj[i * N + j]) { - host_adj_graph[k] = j; - k = k + 1; - } - } - } - host_ex_scan[0] = Index_(0); - for (Index_ i = 1; i < batch_size; i++) - host_ex_scan[i] = host_ex_scan[i - 1] + host_vd[i - 1]; - raft::update_device(data.adj_graph, host_adj_graph.data(), adjgraph_size, stream); - raft::update_device(data.ex_scan, host_ex_scan.data(), batch_size, stream); -} -} // namespace Naive -} // namespace AdjGraph -} // namespace Dbscan -} // namespace ML diff --git a/cpp/src/dbscan/adjgraph/runner.cuh b/cpp/src/dbscan/adjgraph/runner.cuh index ebc26b3bd4..4af3b21739 100644 --- a/cpp/src/dbscan/adjgraph/runner.cuh +++ b/cpp/src/dbscan/adjgraph/runner.cuh @@ -17,7 +17,6 @@ #pragma once #include "algo.cuh" -#include "naive.cuh" #include "pack.h" namespace ML { @@ -41,8 +40,7 @@ void run(const raft::handle_t& handle, { Pack data = {vd, adj, adj_graph, adjnnz, ex_scan, N}; switch (algo) { - // TODO: deprecate naive runner. cf #3414 - case 0: Naive::launcher(handle, data, batch_size, stream); break; + case 0: ASSERT(false, "Incorrect algo '%d' passed! Naive version of adjgraph has been removed.", algo); case 1: Algo::launcher(handle, data, batch_size, row_counters, stream); break; default: ASSERT(false, "Incorrect algo passed! '%d'", algo); } diff --git a/cpp/src/dbscan/common.cuh b/cpp/src/dbscan/common.cuh deleted file mode 100644 index 02da48ed1b..0000000000 --- a/cpp/src/dbscan/common.cuh +++ /dev/null @@ -1,43 +0,0 @@ -/* - * Copyright (c) 2018-2021, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -namespace Dbscan { - -/// Default "ds1" diff-squared-accumulate traits specialization for value_t->accum_t -/// Currently only supported for float/double -template -struct ds_accummulate { - /// Single-component "ds1" diff-squared vector type - typedef value_t dp_vector_t; - - /// Compute "ds1" float->float - inline __device__ static void mad(float& d, const float& a, const float& b, const float& c) - { - float diff = a - b; - asm volatile("fma.rn.f32 %0, %1, %1, %2;\n" : "=f"(d) : "f"(diff), "f"(c)); - } - - /// Compute "ds1" double->double - inline __device__ static void mad(double& d, const double& a, const double& b, const double& c) - { - double diff = a - b; - asm volatile("fma.rn.f64 %0, %1, %1, %2;\n" : "=d"(d) : "d"(diff), "d"(c)); - } -}; - -} // namespace Dbscan diff --git a/cpp/src/dbscan/dbscan.cuh b/cpp/src/dbscan/dbscan.cuh index 24595ff931..6b8c38435f 100644 --- a/cpp/src/dbscan/dbscan.cuh +++ b/cpp/src/dbscan/dbscan.cuh @@ -110,6 +110,8 @@ void dbscanFitImpl(const raft::handle_t& handle, { raft::common::nvtx::range fun_scope("ML::Dbscan::Fit"); ML::Logger::get().setLevel(verbosity); + // XXX: for algo_vd and algo_adj, 0 (naive) is no longer an option and has + // been removed. int algo_vd = (metric == raft::distance::Precomputed) ? 2 : 1; int algo_adj = 1; int algo_ccl = 2; diff --git a/cpp/src/dbscan/vertexdeg/naive.cuh b/cpp/src/dbscan/vertexdeg/naive.cuh deleted file mode 100644 index f332367ca9..0000000000 --- a/cpp/src/dbscan/vertexdeg/naive.cuh +++ /dev/null @@ -1,93 +0,0 @@ -/* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "pack.h" -#include - -namespace ML { -namespace Dbscan { -namespace VertexDeg { -namespace Naive { - -using namespace MLCommon; - -/** number of threads in a CTA along X dim */ -static const int TPB_X = 32; -/** number of threads in a CTA along Y dim */ -static const int TPB_Y = 8; - -/** - * @brief Naive distance matrix evaluation and epsilon neighborhood construction - * @param data input struct containing vertex degree computation params - * @param start_vertex_id which vertex to begin the computations from - * @param batch_size number of vertices in this batch - */ -template -__global__ void vertex_degree_kernel(Pack data, - Index_ start_vertex_id, - Index_ batch_size) -{ - const Type Zero = (Type)0; - Index_ row = (blockIdx.y * TPB_Y) + threadIdx.y; - Index_ col = (blockIdx.x * TPB_X) + threadIdx.x; - Index_ N = data.N; - if ((row >= batch_size) || (col >= N)) return; - Type eps = data.eps; - Type eps2 = eps * eps; - Type sum = Zero; - Index_ D = data.D; - const Type* x = data.x; - bool* adj = data.adj; - Index_* vd = data.vd; - for (Index_ d = 0; d < D; ++d) { - Type a = __ldg(x + (row + start_vertex_id) * D + d); - Type b = __ldg(x + col * D + d); - Type diff = a - b; - sum += (diff * diff); - } - Index_ res = (sum <= eps2); - adj[row * N + col] = res; - /// TODO: change layout or remove; cf #3414 - - if (sizeof(Index_) == 4) { - raft::myAtomicAdd((int*)(vd + row), (int)res); - raft::myAtomicAdd((int*)(vd + batch_size), (int)res); - } else if (sizeof(Index_) == 8) { - raft::myAtomicAdd((unsigned long long*)(vd + row), res); - raft::myAtomicAdd((unsigned long long*)(vd + batch_size), res); - } -} - -template -void launcher(Pack data, - Index_ start_vertex_id, - Index_ batch_size, - cudaStream_t stream) -{ - ASSERT(sizeof(Index_) == 4 || sizeof(Index_) == 8, "index_t should be 4 or 8 bytes"); - - dim3 grid(raft::ceildiv(data.N, (Index_)TPB_X), raft::ceildiv(batch_size, (Index_)TPB_Y), 1); - dim3 blk(TPB_X, TPB_Y, 1); - data.resetArray(stream, batch_size + 1); - vertex_degree_kernel<<>>(data, start_vertex_id, batch_size); -} - -} // namespace Naive -} // namespace VertexDeg -} // namespace Dbscan -} // namespace ML diff --git a/cpp/src/dbscan/vertexdeg/runner.cuh b/cpp/src/dbscan/vertexdeg/runner.cuh index ea1da4f4e0..99804b365b 100644 --- a/cpp/src/dbscan/vertexdeg/runner.cuh +++ b/cpp/src/dbscan/vertexdeg/runner.cuh @@ -17,7 +17,6 @@ #pragma once #include "algo.cuh" -#include "naive.cuh" #include "pack.h" #include "precomputed.cuh" @@ -41,8 +40,7 @@ void run(const raft::handle_t& handle, { Pack data = {vd, adj, x, eps, N, D}; switch (algo) { - // TODO: deprecate naive runner. cf #3414 - case 0: Naive::launcher(data, start_vertex_id, batch_size, stream); break; + case 0: ASSERT(false, "Incorrect algo '%d' passed! Naive version of vertexdeg has been removed.", algo); case 1: Algo::launcher(handle, data, start_vertex_id, batch_size, stream, metric); break;