Errors running raft on custom kernel[QST]

[MikeChenfu]

What is your question?
Hello, I am trying add raft into my custom function but got some error while compiling the code raft::random::make_blobs. Appreciate it if anyone has an idea about this.

Here is my sample source code and cMakeLists.txt
.

#include <stdio.h>
#include <raft/core/handle.hpp>
#include <raft/core/device_mdarray.hpp>
#include <raft/random/make_blobs.cuh>
// #include <raft/distance/distance.cuh>
// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>

__global__ void vectorAdd(const float *A, const float *B, float *C,
                          int numElements) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;

  if (i < numElements) {
    C[i] = A[i] + B[i] + 0.0f;
  }
}

int main(void) {
  raft::handle_t handle;

  int n_samples = 5000;
  int n_features = 50;

  // works good
  auto input = raft::make_device_matrix<float>(handle, n_samples, n_features);
  auto labels = raft::make_device_vector<int>(handle, n_samples);
  auto output = raft::make_device_matrix<float>(handle, n_samples, n_samples);
  // has a problem
  raft::random::make_blobs(handle, input.view(), labels.view());

cMakeLists.txt

cmake_minimum_required(VERSION 3.18)

project(basic_example VERSION 0.0.1 LANGUAGES CXX CUDA)

set(CPM_DOWNLOAD_VERSION v0.32.2)
file(DOWNLOAD https://github.com/cpm-cmake/CPM.cmake/releases/download/${CPM_DOWNLOAD_VERSION}/get_cpm.cmake ${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake)
include(${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake)

find_package(raft)
# Configure your project here
add_executable(basic_example  src/test.cu)
target_link_libraries(basic_example PRIVATE raft::raft)
target_compile_features(basic_example PRIVATE cxx_std_17)

It works well without raft::random::make_blobs(handle, input.view(), labels.view());

(raft_dev_env_v1) cuopt_user@8e1b03d71dff:~/raft/c_pq_io$ bash build.sh  && ./basic/build/basic_example
-- Found raft: /home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/lib/cmake/raft/raft-config.cmake (found version "22.10.0")  
-- Configuring done
-- Generating done
-- Build files have been written to: /home/cuopt_user/raft/c_pq_io/basic/build
Consolidate compiler generated dependencies of target basic_example
[ 50%] Building CUDA object CMakeFiles/basic_example.dir/src/test.cu.o
[100%] Linking CUDA executable basic_example
[100%] Built target basic_example
[Vector addition of 50000 elements]
Copy input data from the host memory to the CUDA device
CUDA kernel launch with 196 blocks of 256 threads
Copy output data from the CUDA device to the host memory
Test PASSED
Done

Errors

(raft_dev_env_v1) cuopt_user@8e1b03d71dff:~/raft/c_pq_io$ bash build.sh  && ./basic/build/basic_example
-- Configuring done
-- Generating done
-- Build files have been written to: /home/cuopt_user/raft/c_pq_io/basic/build
Consolidate compiler generated dependencies of target basic_example
[ 50%] Building CUDA object CMakeFiles/basic_example.dir/src/test.cu.o
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/linalg/detail/unary_op.cuh(85): error: calling a constexpr __host__ function("operator()") from a __global__ function("writeOnlyUnaryOpKernel") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.
          detected during:
            instantiation of "void raft::linalg::detail::writeOnlyUnaryOpKernel(OutType *, IdxType, Lambda) [with OutType=int, Lambda=lambda [](int *, int)->void, IdxType=int]" 
(93): here
            instantiation of "void raft::linalg::detail::writeOnlyUnaryOpCaller(OutType *, IdxType, Lambda, cudaStream_t) [with OutType=int, Lambda=lambda [](int *, int)->void, IdxType=int, TPB=256]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/linalg/unary_op.cuh(71): here
            instantiation of "void raft::linalg::writeOnlyUnaryOp(OutType *, IdxType, Lambda, cudaStream_t) [with OutType=int, Lambda=lambda [](int *, int)->void, IdxType=int, TPB=256]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/random/detail/make_blobs.cuh(51): here
            instantiation of "void raft::random::detail::generate_labels(IdxT *, IdxT, IdxT, __nv_bool, raft::random::RngState &, cudaStream_t) [with IdxT=int]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/random/detail/make_blobs.cuh(239): here
            instantiation of "void raft::random::detail::make_blobs_caller(DataT *, IdxT *, IdxT, IdxT, IdxT, cudaStream_t, __nv_bool, const DataT *, const DataT *, DataT, __nv_bool, DataT, DataT, uint64_t, raft::random::GeneratorType) [with DataT=float, IdxT=int]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/random/make_blobs.cuh(174): here
            instantiation of "void raft::random::make_blobs(const raft::handle_t &, raft::device_matrix_view<DataT, IdxT, layout>, raft::device_vector_view<IdxT, IdxT, raft::layout_c_contiguous>, IdxT, std::optional<raft::device_matrix_view<DataT, IdxT, layout>>, std::optional<raft::device_vector_view<DataT, IdxT, raft::layout_c_contiguous>>, DataT, __nv_bool, DataT, DataT, uint64_t, raft::random::GeneratorType) [with DataT=float, IdxT=int, layout=raft::layout_c_contiguous]" 
/home/cuopt_user/raft/c_pq_io/basic/src/test.cu(38): here

/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/linalg/detail/unary_op.cuh(85): error: calling a constexpr __host__ function("operator()") from a __global__ function("writeOnlyUnaryOpKernel") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.
          detected during:
            instantiation of "void raft::linalg::detail::writeOnlyUnaryOpKernel(OutType *, IdxType, Lambda) [with OutType=int, Lambda=lambda [](int *, int)->void, IdxType=int]" 
(93): here
            instantiation of "void raft::linalg::detail::writeOnlyUnaryOpCaller(OutType *, IdxType, Lambda, cudaStream_t) [with OutType=int, Lambda=lambda [](int *, int)->void, IdxType=int, TPB=256]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/linalg/unary_op.cuh(71): here
            instantiation of "void raft::linalg::writeOnlyUnaryOp(OutType *, IdxType, Lambda, cudaStream_t) [with OutType=int, Lambda=lambda [](int *, int)->void, IdxType=int, TPB=256]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/random/detail/make_blobs.cuh(51): here
            instantiation of "void raft::random::detail::generate_labels(IdxT *, IdxT, IdxT, __nv_bool, raft::random::RngState &, cudaStream_t) [with IdxT=int]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/random/detail/make_blobs.cuh(239): here
            instantiation of "void raft::random::detail::make_blobs_caller(DataT *, IdxT *, IdxT, IdxT, IdxT, cudaStream_t, __nv_bool, const DataT *, const DataT *, DataT, __nv_bool, DataT, DataT, uint64_t, raft::random::GeneratorType) [with DataT=float, IdxT=int]" 
/home/cuopt_user/miniconda/conda/envs/raft_dev_env_v1/include/raft/random/make_blobs.cuh(174): here
            instantiation of "void raft::random::make_blobs(const raft::handle_t &, raft::device_matrix_view<DataT, IdxT, layout>, raft::device_vector_view<IdxT, IdxT, raft::layout_c_contiguous>, IdxT, std::optional<raft::device_matrix_view<DataT, IdxT, layout>>, std::optional<raft::device_vector_view<DataT, IdxT, raft::layout_c_contiguous>>, DataT, __nv_bool, DataT, DataT, uint64_t, raft::random::GeneratorType) [with DataT=float, IdxT=int, layout=raft::layout_c_contiguous]" 
/home/cuopt_user/raft/c_pq_io/basic/src/test.cu(38): here

2 errors detected in the compilation of "/home/cuopt_user/raft/c_pq_io/basic/src/test.cu".
make[2]: *** [CMakeFiles/basic_example.dir/build.make:76: CMakeFiles/basic_example.dir/src/test.cu.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:83: CMakeFiles/basic_example.dir/all] Error 2
make: *** [Makefile:111: all] Error 2

[cjnolet]

Hi @MikeChenfu, thanks for opening an issue about this.

It looks like --expt-relaxed-constexpr needs to be added to your CUDA Compile flags. You should be able to do this using the target_compile_options in cmake. Can you try this and let us know if it still doesn't compile?

For example:

target_compile_options(basic_example PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda --expt-relaxed-constexpr>)

[MikeChenfu]

@cjnolet Thanks for the quick reply. That works. Appreciate it!

[cjnolet]

And thanks again for opening the issue. It's an indication that this line should probably be added to our build docs.