no embedding headers

Author: divyegala

cpp/CMakeLists.txt

@@ -452,32 +452,7 @@ if(NOT BUILD_CPU_ONLY)
   )
   endblock()
 
-  find_file(
-    CUVS_BUILD_TIME_CUDA_FP16_H cuda_fp16.h
-    PATHS ${CUDAToolkit_INCLUDE_DIRS}
-    NO_DEFAULT_PATH REQUIRED
-    DOC "cuda_fp16.h from the CUDA toolkit (embedded for NVRTC UDF JIT)"
-  )
-
-  set(CUVS_EMBEDDED_CUDA_FP16_CPP
-      "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/embedded_cuda_fp16.cpp"
-  )
-  file(MAKE_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto")
-  add_custom_command(
-    OUTPUT ${CUVS_EMBEDDED_CUDA_FP16_CPP}
-    COMMAND
-      ${CMAKE_COMMAND} "-DINPUT=${CUVS_BUILD_TIME_CUDA_FP16_H}"
-      "-DOUTPUT=${CUVS_EMBEDDED_CUDA_FP16_CPP}" -P
-      "${CMAKE_CURRENT_SOURCE_DIR}/cmake/embed_cuda_fp16.cmake"
-    DEPENDS "${CUVS_BUILD_TIME_CUDA_FP16_H}"
-            "${CMAKE_CURRENT_SOURCE_DIR}/cmake/embed_cuda_fp16.cmake"
-    COMMENT
-      "Embedding cuda_fp16.h from CUDA toolkit for NVRTC (rebuilt when toolkit header changes)"
-    VERBATIM
-  )
-
   set(jit_lto_files
-      ${CUVS_EMBEDDED_CUDA_FP16_CPP}
       ${interleaved_scan_files}
       ${metric_files}
       ${filter_files}
@@ -767,10 +742,6 @@ if(NOT BUILD_CPU_ONLY)
     ${jit_lto_files}
   )
 
-  set_source_files_properties("${CUVS_EMBEDDED_CUDA_FP16_CPP}" PROPERTIES GENERATED TRUE)
-  add_custom_target(cuvs_embed_cuda_fp16_header DEPENDS "${CUVS_EMBEDDED_CUDA_FP16_CPP}")
-  add_dependencies(cuvs_objs cuvs_embed_cuda_fp16_header)
-
   set_target_properties(
     cuvs_objs
     PROPERTIES CXX_STANDARD 20

cpp/cmake/embed_cuda_fp16.cmake

@@ -3062,10 +3062,11 @@ namespace experimental::udf {
 /**
  * @brief Wrapper for vector elements that provides both packed and unpacked access.
  *
- * For float/half: trivial wrapper around scalar values
+ * For float: trivial wrapper around scalar values
  * For int8/uint8 with Veclen > 1: wraps packed bytes in a 32-bit word
  *
- * @tparam T Data type (float, __half, int8_t, uint8_t)
+ * @tparam T Data type (float, int8_t, uint8_t). Fp16 vector elements are not supported for UDFs
+ *             at this time (see `metric_interface` static_assert when `cuda_fp16.h` is available).
  * @tparam AccT Storage/accumulator type (float, __half, int32_t, uint32_t)
  * @tparam Veclen Vector length (1, 2, 4, 8, 16)
  */
@@ -3130,6 +3131,13 @@ template <typename T, typename AccT, int Veclen = 1>
 struct metric_interface {
   using point_type = point<T, AccT, Veclen>;
 
+#if CUVS_IVF_FLAT_UDF_HAS_CUDA_FP16
+  static_assert(
+    !(std::is_same_v<std::remove_cv_t<T>, __half> || std::is_same_v<std::remove_cv_t<T>, half>),
+    "IVF-Flat custom metric UDF does not support fp16 (__half / half) at this time; do not set "
+    "search_params.metric_udf for fp16 indices.");
+#endif
+
   virtual __device__ void operator()(AccT& acc, point_type x, point_type y) = 0;
   virtual ~metric_interface()                                               = default;
 };
@@ -3380,8 +3388,6 @@ __device__ __forceinline__ AccT max_elem(point<T, AccT, V> x, point<T, AccT, V>
  * the necessary types and utilities inline.
  */
 constexpr std::string_view jit_preamble_code = R"(
-#include <cuda_fp16.h>
-
 /* Fixed-width integer types for nvrtc */
 using int8_t = signed char;
 using uint8_t = unsigned char;

cpp/include/cuvs/detail/jit_lto/embedded_cuda_fp16.hpp

@@ -4,7 +4,6 @@
  */
 
 #include <cuvs/detail/jit_lto/NVRTCLTOFragmentCompiler.hpp>
-#include <cuvs/detail/jit_lto/embedded_cuda_fp16.hpp>
 
 #include <mutex>
 
@@ -37,8 +36,6 @@ NVRTCLTOFragmentCompiler::NVRTCLTOFragmentCompiler()
     std::string{"-rdc=true"},
     std::string{"--std=c++20"},
     std::string{"-default-device"},
-    // cuda_fp16.h skips stdlib / vector includes when compiling for NVRTC
-    std::string{"-D__CUDACC_RTC__"},
   };
 }
 
@@ -62,12 +59,9 @@ std::unique_ptr<UDFFatbinFragment> NVRTCLTOFragmentCompiler::compile(std::string
     return std::make_unique<UDFFatbinFragment>(key, it->second);
   }
 
-  char const* const udf_headers[]       = {cuvs::detail::jit_lto::k_nvrtc_embedded_cuda_fp16_h};
-  char const* const udf_include_names[] = {"cuda_fp16.h"};
-
   nvrtcProgram prog;
-  NVRTC_SAFE_CALL(nvrtcCreateProgram(
-    &prog, code.c_str(), "nvrtc_lto_fragment", 1, udf_headers, udf_include_names));
+  NVRTC_SAFE_CALL(
+    nvrtcCreateProgram(&prog, code.c_str(), "nvrtc_lto_fragment", 0, nullptr, nullptr));
 
   // Convert std::vector<std::string> to std::vector<const char*> for nvrtc API
   std::vector<const char*> opts;

cpp/include/cuvs/neighbors/ivf_flat.hpp

@@ -31,7 +31,7 @@ CUVS_METRIC(custom_l2, { acc += squared_diff(x, y); })
 // ============================================================================
 
 template <typename T>
-concept udf_test_fp_element = std::same_as<T, float> || std::same_as<T, __half>;
+concept udf_test_fp_element = std::same_as<T, float>;
 
 template <typename T>
 concept udf_test_int_byte_element = std::same_as<T, int8_t> || std::same_as<T, uint8_t>;
@@ -46,7 +46,7 @@ struct TestDataTraits<T> {
 
   static std::vector<T> database()
   {
-    // 4-dimensional float/__half dataset
+    // 4-dimensional float dataset
     // Vectors arranged for easy distance verification:
     //   db[0] = [0, 0, 0, 0]  - origin
     //   db[1] = [1, 0, 0, 0]  - unit along x
@@ -292,7 +292,7 @@ class IvfFlatUdfTest : public ::testing::Test {
   uint32_t n_probes_;
 };
 
-using TestTypes = ::testing::Types<float, __half, int8_t, uint8_t>;
+using TestTypes = ::testing::Types<float, int8_t, uint8_t>;
 TYPED_TEST_SUITE(IvfFlatUdfTest, TestTypes);
 
 // ============================================================================