Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Use fusedL2NN in ANN kmeans #821

Merged
merged 20 commits into from
Oct 4, 2022
Merged

Use fusedL2NN in ANN kmeans #821

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
90730e3
Initial integration of fusedL2NN in ANN kmeans
Nyrio Sep 6, 2022
e69d532
Precompute X norm for L2 metrics
Nyrio Sep 7, 2022
c42c0b7
Fix integer overflow and pass index type to predict_float_core
Nyrio Sep 8, 2022
eb2ee66
Merge remote-tracking branch 'origin/branch-22.10' into enh-ann-kmeans
Nyrio Sep 12, 2022
76d4908
Fix for integer types
Nyrio Sep 12, 2022
4b497ff
Cleanup debug function
Nyrio Sep 12, 2022
fd290fa
Use int64_t indices in IVF-Flat test to enable large test cases
Nyrio Sep 13, 2022
8cfee2c
Replace if by switch-case, C-style casts by C++ style, use int8 share…
Nyrio Sep 13, 2022
4736f93
Do not batch when using fusedL2NN
Nyrio Sep 22, 2022
2f8500d
Support for unsigned label type
Nyrio Sep 22, 2022
0a1ce14
Make pre-computed norm an optional argument
Nyrio Sep 23, 2022
ab2890a
Rewrite batch heuristic + change hard-coded integer types to templated
Nyrio Sep 23, 2022
ef9b281
Merge remote-tracking branch 'origin/branch-22.10' into enh-ann-kmeans
Nyrio Sep 26, 2022
1d0ab4c
Index type capacity assertion + use make_device_mdarray with memory r…
Nyrio Sep 29, 2022
3dcac56
Use thrust::gather instead of copy_selected for norm
Nyrio Sep 30, 2022
0bb179d
Solve warnings due to printf qualifier / type mismatch
Nyrio Sep 30, 2022
0e4ad82
Merge remote-tracking branch 'origin/branch-22.10' into enh-ann-kmeans
Nyrio Sep 30, 2022
5cfb67e
Use fusedL2NN wrapper compiled in distance library
Nyrio Sep 30, 2022
4e35df5
Merge remote-tracking branch 'origin/branch-22.10' into enh-ann-kmeans
Nyrio Oct 4, 2022
fa946c5
Fix ivf-pq / ann utils merge
Nyrio Oct 4, 2022
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
3 changes: 3 additions & 0 deletions cpp/bench/spatial/knn.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,9 @@

#if defined RAFT_NN_COMPILED
#include <raft/spatial/knn/specializations.cuh>
#if defined RAFT_DISTANCE_COMPILED
#include <raft/cluster/specializations.cuh>
#endif
#endif

#include <rmm/mr/device/managed_memory_resource.hpp>
Expand Down
6 changes: 3 additions & 3 deletions cpp/include/raft/distance/detail/fused_l2_nn.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ struct MinAndDistanceReduceOpImpl {
DI void init(DataT* out, DataT maxVal) { *out = maxVal; }
DI void init(KVP* out, DataT maxVal)
{
out->key = -1;
out->key = 0;
out->value = maxVal;
}
};
Expand Down Expand Up @@ -150,7 +150,7 @@ __global__ __launch_bounds__(P::Nthreads, 2) void fusedL2NNkernel(OutT* min,
KVPair val[P::AccRowsPerTh];
#pragma unroll
for (int i = 0; i < P::AccRowsPerTh; ++i) {
val[i] = {-1, maxVal};
val[i] = {0, maxVal};
}

// epilogue operation lambda for final value calculation
Expand Down Expand Up @@ -222,7 +222,7 @@ __global__ __launch_bounds__(P::Nthreads, 2) void fusedL2NNkernel(OutT* min,
// reset the val array.
#pragma unroll
for (int i = 0; i < P::AccRowsPerTh; ++i) {
val[i] = {-1, maxVal};
val[i] = {0, maxVal};
}
};

Expand Down
4 changes: 0 additions & 4 deletions cpp/include/raft/distance/fused_l2_nn.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,10 +101,6 @@ void fusedL2NN(OutT* min,
bool initOutBuffer,
cudaStream_t stream)
{
// Assigning -1 to unsigned integers results in a compiler error.
// Enforce a signed IdxT here with a clear error message.
static_assert(std::is_signed_v<IdxT>, "fusedL2NN only supports signed index types.");

// When k is smaller than 32, the Policy4x4 results in redundant calculations
// as it uses tiles that have k=32. Therefore, use a "skinny" policy instead
// that uses tiles with a smaller value of k.
Expand Down
3 changes: 2 additions & 1 deletion cpp/include/raft/linalg/coalesced_reduction.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,8 @@ void coalescedReduction(OutType* dots,
ReduceLambda reduce_op = raft::Sum<OutType>(),
FinalLambda final_op = raft::Nop<OutType>())
{
detail::coalescedReduction(dots, data, D, N, init, stream, inplace, main_op, reduce_op, final_op);
detail::coalescedReduction<InType, OutType, IdxType>(
dots, data, D, N, init, stream, inplace, main_op, reduce_op, final_op);
}

/**
Expand Down
96 changes: 48 additions & 48 deletions cpp/include/raft/linalg/detail/norm.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,32 +37,32 @@ void rowNormCaller(Type* dots,
{
switch (type) {
case L1Norm:
raft::linalg::reduce(dots,
data,
D,
N,
(Type)0,
rowMajor,
true,
stream,
false,
raft::L1Op<Type, IdxType>(),
raft::Sum<Type>(),
fin_op);
raft::linalg::reduce<Type, Type, IdxType>(dots,
data,
D,
N,
(Type)0,
rowMajor,
true,
stream,
false,
raft::L1Op<Type, IdxType>(),
raft::Sum<Type>(),
fin_op);
break;
case L2Norm:
raft::linalg::reduce(dots,
data,
D,
N,
(Type)0,
rowMajor,
true,
stream,
false,
raft::L2Op<Type>(),
raft::Sum<Type>(),
fin_op);
raft::linalg::reduce<Type, Type, IdxType>(dots,
data,
D,
N,
(Type)0,
rowMajor,
true,
stream,
false,
raft::L2Op<Type>(),
raft::Sum<Type>(),
fin_op);
break;
default: ASSERT(false, "Invalid norm type passed! [%d]", type);
};
Expand All @@ -80,32 +80,32 @@ void colNormCaller(Type* dots,
{
switch (type) {
case L1Norm:
raft::linalg::reduce(dots,
data,
D,
N,
(Type)0,
rowMajor,
false,
stream,
false,
raft::L1Op<Type, IdxType>(),
raft::Sum<Type>(),
fin_op);
raft::linalg::reduce<Type, Type, IdxType>(dots,
data,
D,
N,
(Type)0,
rowMajor,
false,
stream,
false,
raft::L1Op<Type, IdxType>(),
raft::Sum<Type>(),
fin_op);
break;
case L2Norm:
raft::linalg::reduce(dots,
data,
D,
N,
(Type)0,
rowMajor,
false,
stream,
false,
raft::L2Op<Type, IdxType>(),
raft::Sum<Type>(),
fin_op);
raft::linalg::reduce<Type, Type, IdxType>(dots,
data,
D,
N,
(Type)0,
rowMajor,
false,
stream,
false,
raft::L2Op<Type, IdxType>(),
raft::Sum<Type>(),
fin_op);
break;
default: ASSERT(false, "Invalid norm type passed! [%d]", type);
};
Expand Down
8 changes: 4 additions & 4 deletions cpp/include/raft/linalg/detail/reduce.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,16 +44,16 @@ void reduce(OutType* dots,
FinalLambda final_op = raft::Nop<OutType>())
{
if (rowMajor && alongRows) {
raft::linalg::coalescedReduction(
raft::linalg::coalescedReduction<InType, OutType, IdxType>(
dots, data, D, N, init, stream, inplace, main_op, reduce_op, final_op);
} else if (rowMajor && !alongRows) {
raft::linalg::stridedReduction(
raft::linalg::stridedReduction<InType, OutType, IdxType>(
dots, data, D, N, init, stream, inplace, main_op, reduce_op, final_op);
} else if (!rowMajor && alongRows) {
raft::linalg::stridedReduction(
raft::linalg::stridedReduction<InType, OutType, IdxType>(
dots, data, N, D, init, stream, inplace, main_op, reduce_op, final_op);
} else {
raft::linalg::coalescedReduction(
raft::linalg::coalescedReduction<InType, OutType, IdxType>(
dots, data, N, D, init, stream, inplace, main_op, reduce_op, final_op);
}
}
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/raft/linalg/reduce.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,7 @@ void reduce(OutType* dots,
ReduceLambda reduce_op = raft::Sum<OutType>(),
FinalLambda final_op = raft::Nop<OutType>())
{
detail::reduce(
detail::reduce<InType, OutType, IdxType>(
dots, data, D, N, init, rowMajor, alongRows, stream, inplace, main_op, reduce_op, final_op);
}

Expand Down
3 changes: 2 additions & 1 deletion cpp/include/raft/linalg/strided_reduction.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,8 @@ void stridedReduction(OutType* dots,
ReduceLambda reduce_op = raft::Sum<OutType>(),
FinalLambda final_op = raft::Nop<OutType>())
{
detail::stridedReduction(dots, data, D, N, init, stream, inplace, main_op, reduce_op, final_op);
detail::stridedReduction<InType, OutType, IdxType>(
dots, data, D, N, init, stream, inplace, main_op, reduce_op, final_op);
}

/**
Expand Down
Loading