From 411c1721da774099e3bfd62b5392cec20c01729a Mon Sep 17 00:00:00 2001
From: Robert Maynard <rmaynard@nvidia.com>
Date: Wed, 15 Nov 2023 02:29:21 -0800
Subject: [PATCH] Add linker script to support large cuda cubin files (#3115)

Summary:
nvcc starting with CUDA 11.5 offers a `-hls` option to generate host side linker scripts to support large cubin file support.
Since faiss supports CUDA 11.4 we replicate that behavior but injecting the same linker script into the link line manually.

Pull Request resolved: https://github.com/facebookresearch/faiss/pull/3115

Reviewed By: mdouze

Differential Revision: D51308908

Pulled By: algoriddle

fbshipit-source-id: c6dd073cd3f44dbc99d2e2da97f79b9ebc843b59
---
 faiss/gpu/CMakeLists.txt | 16 ++++++++++++++++
 1 file changed, 16 insertions(+)

diff --git a/faiss/gpu/CMakeLists.txt b/faiss/gpu/CMakeLists.txt
index ad7d2103fa..0ab6ff3cea 100644
--- a/faiss/gpu/CMakeLists.txt
+++ b/faiss/gpu/CMakeLists.txt
@@ -261,6 +261,22 @@ foreach(header ${FAISS_GPU_HEADERS})
   )
 endforeach()
 
+# Prepares a host linker script and enables host linker to support
+# very large device object files.
+# This is what CUDA 11.5+ `nvcc -hls=gen-lcs -aug-hls` would generate
+file(WRITE "${CMAKE_CURRENT_BINARY_DIR}/fatbin.ld"
+[=[
+SECTIONS
+{
+.nvFatBinSegment : { *(.nvFatBinSegment) }
+__nv_relfatbin : { *(__nv_relfatbin) }
+.nv_fatbin : { *(.nv_fatbin) }
+}
+]=]
+)
+target_link_options(faiss PRIVATE "${CMAKE_CURRENT_BINARY_DIR}/fatbin.ld")
+target_link_options(faiss_avx2 PRIVATE "${CMAKE_CURRENT_BINARY_DIR}/fatbin.ld")
+
 find_package(CUDAToolkit REQUIRED)
 target_link_libraries(faiss PRIVATE CUDA::cudart CUDA::cublas $<$<BOOL:${FAISS_ENABLE_RAFT}>:raft::raft> $<$<BOOL:${FAISS_ENABLE_RAFT}>:raft::compiled> $<$<BOOL:${FAISS_ENABLE_RAFT}>:nvidia::cutlass::cutlass>)
 target_link_libraries(faiss_avx2 PRIVATE CUDA::cudart CUDA::cublas $<$<BOOL:${FAISS_ENABLE_RAFT}>:raft::raft> $<$<BOOL:${FAISS_ENABLE_RAFT}>:raft::compiled>  $<$<BOOL:${FAISS_ENABLE_RAFT}>:nvidia::cutlass::cutlass>)