[QST] Custom Distance function & what is the DI directive/keyword? about raft HOT 9 CLOSED

Hi Daniel, adding a custom distance function requires some plumbing:

• Add the distance type to the enum here:

  raft/cpp/include/raft/distance/distance_types.hpp

  Line 25 in affd29b

  /** evaluate as dist_ij = sum(x_ik^2) + sum(y_ij)^2 - 2*sum(x_ik * y_jk) */

• Add to the switch case here:

  raft/cpp/include/raft/distance/distance-inl.cuh

  Line 259 in affd29b

  switch (metric) {

• Add an overload of distance_impl here:

  raft/cpp/include/raft/distance/detail/distance.cuh

  Line 115 in affd29b

  void distance_impl(raft::resources const& handle,

• Add a distance op, implementing the distance using the template that you already found here:

  raft/cpp/include/raft/distance/detail/distance_ops/template.cuh

  Line 27 in affd29b

  template <typename DataType, typename AccType, typename IdxType>

I might be missing some things, like adding some includes, but this should mostly cover it for the header-only use case. This way, you can call the function using raft::distance::pairwise_distance(...).

If you want the distance operation to be built as part of the shared library, you may need to add an instance here:

Hope this helps you get started!

from raft.

The DI is a macro that expands to __device__ inline. It is defined here:

from raft.

It worked! Fantastic! I will keep the question open for now in case there are any other issues while setting up a custom distance function. Thank you!

from raft.

Thanks for the answers! I am stuck at a much simpler step when trying to compile the template.

I installed raft via mamba:

mamba install -c rapidsai -c conda-forge -c nvidia raft-dask pylibraft cuda-version=11.4

which worked.

Then I copied the template code from the repository into a seperate directory and set-up the following cmake file

cmake_minimum_required(VERSION 3.23.1 FATAL_ERROR)
project(test_raft LANGUAGES CXX CUDA)
enable_language(CUDA)
find_package(raft REQUIRED)
add_executable(basic_example  test.cu)
target_link_libraries(basic_example PRIVATE raft::raft)

which runs without any issues.

However, I then run the make command and receive the following error

[ 50%] Building CUDA object CMakeFiles/basic_example.dir/test.cu.o
[100%] Linking CUDA executable basic_example
CMakeFiles/basic_example.dir/test.cu.o: In function `raft::resource::cublas_resource::cublas_resource(rmm::cuda_stream_view)':
tmpxft_00004343_00000000-6_test.cudafe1.cpp:(.text._ZN4raft8resource15cublas_resourceC2EN3rmm16cuda_stream_viewE[_ZN4raft8resource15cublas_resourceC5EN3rmm16cuda_stream_viewE]+0x35): undefined reference to `cublasCreate_v2'
tmpxft_00004343_00000000-6_test.cudafe1.cpp:(.text._ZN4raft8resource15cublas_resourceC2EN3rmm16cuda_stream_viewE[_ZN4raft8resource15cublas_resourceC5EN3rmm16cuda_stream_viewE]+0x8b): undefined reference to `cublasSetStream_v2'
CMakeFiles/basic_example.dir/test.cu.o: In function `raft::resource::cublas_resource::~cublas_resource()':
tmpxft_00004343_00000000-6_test.cudafe1.cpp:(.text._ZN4raft8resource15cublas_resourceD2Ev[_ZN4raft8resource15cublas_resourceD5Ev]+0x24): undefined reference to `cublasDestroy_v2'
CMakeFiles/basic_example.dir/test.cu.o: In function `raft::resource::get_cublas_handle(raft::resources const&)':
tmpxft_00004343_00000000-6_test.cudafe1.cpp:(.text._ZN4raft8resource17get_cublas_handleERKNS_9resourcesE[_ZN4raft8resource17get_cublas_handleERKNS_9resourcesE]+0x119): undefined reference to `cublasSetStream_v2'
CMakeFiles/basic_example.dir/test.cu.o: In function `cublasStatus_t raft::linalg::detail::cublasgemm<float>(cublasContext*, cublasOperation_t, cublasOperation_t, int, int, int, float const*, float const*, int, float const*, int, float const*, float*, int, CUstream_st*)':
tmpxft_00004343_00000000-6_test.cudafe1.cpp:(.text._ZN4raft6linalg6detail10cublasgemmIfEE14cublasStatus_tP13cublasContext17cublasOperation_tS6_iiiPKT_S9_iS9_iS9_PS7_iP11CUstream_st[_ZN4raft6linalg6detail10cublasgemmIfEE14cublasStatus_tP13cublasContext17cublasOperation_tS6_iiiPKT_S9_iS9_iS9_PS7_iP11CUstream_st]+0x49): undefined reference to `cublasSetStream_v2'
tmpxft_00004343_00000000-6_test.cudafe1.cpp:(.text._ZN4raft6linalg6detail10cublasgemmIfEE14cublasStatus_tP13cublasContext17cublasOperation_tS6_iiiPKT_S9_iS9_iS9_PS7_iP11CUstream_st[_ZN4raft6linalg6detail10cublasgemmIfEE14cublasStatus_tP13cublasContext17cublasOperation_tS6_iiiPKT_S9_iS9_iS9_PS7_iP11CUstream_st]+0x38a): undefined reference to `cublasSgemm_v2'
collect2: error: ld returned 1 exit status
make[2]: *** [basic_example] Error 1
make[1]: *** [CMakeFiles/basic_example.dir/all] Error 2
make: *** [all] Error 2

The reason I used mamba was because I thought I may not be installing from source correctly. I had ran the 'build.sh' file with a combination of flags and tags but I kept getting the same error when I tried to compile. Then I installed the latest cuBLAS version, but that does not seem to help. Do you have any advice?

nvcc --version: 11.4
cmake --version: 3.27.1

from raft.

Add something like this to the CMake:

target_link_libraries (name_of_your_program PRIVATE 
  raft::raft 
  other_deps 
  CUDA::cublas // <- this one might be missing
)

from raft.

Great! I am curious which distance function you are trying to implement?

from raft.

I would like to have a different distance metric for various features and change their weighting as well. For example, when using the correlation metric, the correlation measure normalizes the data implicitly, so I can not simply scale the feature to change its contribution to the overall distance.

I have some cuda code for my own kNN implementation, but it seems to be quite a bit slower than the implementation of raft. Seems like it may be a combination of RMM and other optimization tricks?

from raft.

Interesting! Weighting indeed seems like a useful feature.

Seems like it may be a combination of RMM and other optimization tricks?

I cannot speak for your code, but the meat of the algorithm can be found here:

We use the "standard" tricks like tiling in shared memory, avoiding shared memory bank conflicts, and loading data using vectorized loads.

from raft.

I am closing this issue. Feel free to open a new one if any questions arise!

from raft.