Bump CUDA -> 11.4 and cuDNN -> 8.0.5

~~большой~~ PR по переключению дефолтной версии CUDA с cuDNN в Аркадии

Обновляем CUDA: 10.1 -> 11.4
Обновляем cuDNN: 7.6.5 -> 8.0.5

Помимо простого обновления версий, данный PR содержит следующее:

* От перехода на CUDA 11.4 честно сломался только [один проект](https://a.yandex-team.ru/arcadia/cv/imgclassifiers/danet/backend/gpu_cuda?rev=rXXXXXX), поэтому просто там правим
* Где-то поменялись тесты на объём потребляемой памяти, поэтому их просто переканонизируем
* По поводу удаления таргетов из clang_prev_targets:
  - `ld.lld` из поставки clang14 не справляется с линковкой - не может найти символы в попруненных либах, хотя они там есть
  - но можно просто переключиться на автосборку, потому что те проблемы [совместимости cuDNN и clang](https://st.yandex-team.ru/#655b977f30316b33e3a5ec87), для которых gpu-шные таргеты добавляли в clang_prev_targets, в автосборке уже не существуют
* По поводу удаления таргетов из cuda11_targets - теперь cuda11_targets отличается от автосборки версией TensorRT (5 в автосборке vs 7 в cuda11) и режимом сборки (relwithdebinfo в автосборке vs release в cuda11), но есть два момента
  - те, таргеты, которые я удаляю, не зависят от TensorRT, поэтому их не имеет смысла их тестить и в автосборке, и в cuda11 (про таргет `ads/quality/sis/tests/cuda11_arch80` мне сейчас ничего не известно; если что - выпилим его отдельным таргетом)
  - на самом деле дублирование даже вредно - в `dict/mt/daemon/tests/gpu` есть тест на потребление памяти gpu, который имеет разные результаты в зависимости от release vs relwithdebinfo, поэтому для этого таргета мы в принципе не можем собираться одновременно и в автосборке, и в cuda11

По поводу дефолтной автосборки в принципе пришлось сделать четыре приседания:
* пишем кастомный скрипт для линкера:
  - идея скрипта: переставляем большую секцию с gpu-шным кодом (`.nv_fatbin`) после `.bss` (самая дальняя секция, куда можно ожидать ссылки из кода бинаря), чтобы было меньше шансов нарваться на проблемы с relocation overflow в (`.bss`)
  - замечание: в самом скрипте используем в нём только `INSERT AFTER`, чтобы [ld.ldd и дальше применял свои дефолтные правила для остальной программы](https://releases.llvm.org/16.0.0/tools/lld/docs/ELF/linker_script.html#sections-command)
  - сделал замеры перфа генезисного инференса через `ml/zeliboba/libs/ynmt_lm/translate/bin` на том сценарии, который оказался под рукой  - не заметил каких-то изменений
* добавляем nvcc-флаг `-Xfatbin=-compress-all`, чтобы наш fatbin сжимался и занимал меньше пространства в бинаре (далее он будет разжиматься на старте программы,  так что оверхеда быть не должно)
  - сделал замеры перфа генезисного инференса через `ml/zeliboba/libs/ynmt_lm/translate/bin` на том сценарии, который оказался под рукой  - не заметил каких-то изменений
* добавляем флаг линкера `--no-relax`, так как опция "relaxation of relocatable symbols" (когда мы load в регистр + jump по адресу в регистре заменяем на относительный jump) уменьшает допустимый диапазон оффсетов, которые мы можем кодировать, что также приводит к relocation overflow
  - больше информации можно найти в https://maskray.me/blog/2023-05-14-relocation-overflow-and-code-models
  - сделал замеры перфа генезисного инференса через `ml/zeliboba/libs/ynmt_lm/translate/bin` на том сценарии, который оказался под рукой  - не заметил каких-то изменений
* ещё чуть более гранулярно выпиливаем лобзиком архитектуры из либ в поставках CUDA / cuDNN - в категории "ассемблерный код" (`compute_XX`) оставляем только для последней версии поддерживаемой архитектуры (то есть `compute_86`), так как на новых GPU-шках мы будем JIT-компилироваться из него, а на старых GPU-шках мы возьмём готовый машинный код для них

Эти идеи были в основном взяты из следующих источников:
- https://maskray.me/blog/2021-07-04-sections-and-overwrite-sections#insert-before-and-insert-after и https://discourse.llvm.org/t/lld-relocation-overflows-and-nv-fatbin/58889 (про перемещение `.nv_fatbin`)
- https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html#nvprune (про `nvprune`)
- https://maskray.me/blog/2023-05-14-relocation-overflow-and-code-models#relocation-overflow (`-Wl,--no-relax`)
- https://discourse.llvm.org/t/lld-relocation-overflows-and-nv-fatbin/58889/6 + https://github.com/pytorch/pytorch/pull/43074/files#diff-1e7de1ae2d059d21e1dd75d5812d5a34b0222cef273b7c3a2af62eb747f9d20aR360 (про `-Xfatbin=-compress-all`)

Дополнительные комментарии:
- можно дополнительно не выпиливать ненужные архитектуры, но этот флаг дополнительно уменьшает размеры бинарей (например, таргет `ml/zeliboba/libs/ynmt_lm/translate/bin`уменьшается на 120 MiB, что составляет примерно 5% от размера бинаря)
- кастомный скрипт для линкера позволяет не развлекаться с ещё более гранулярным выпиливанием архитектур (например, для таргетов `TENSORFLOW_WITH_CUDA`)

Оставшийся технический долг, который будет делать в других PRах
* донести функциональность из `link_exe.py` в `link_dyn_lib.py`
* перейти на TensorRT 7
* выпилить cuda11_targets
* выпилить флаг `TENSORFLOW_WITH_CUDA` в пользу `CUDA_REQUIRED` и удалить tensorflow_with_cuda_targets (??)
* поднять версию стандарта c++ для гпушного кода с 14 до 17

2 files changed, 93 insertions, 41 deletions

diff --git a/build/scripts/link_exe.py b/build/scripts/link_exe.py
index 4650315b89..c20875c77a 100644
--- a/build/scripts/link_exe.py
+++ b/build/scripts/link_exe.py
@@ -4,6 +4,7 @@ import os.path
 import sys
 import subprocess
 import optparse
+import textwrap
 
 import process_command_files as pcf
 
@@ -43,31 +44,74 @@ CUDA_LIBRARIES = {
 }
 
 
-def prune_cuda_libraries(cmd, prune_arches, nvprune_exe, build_root):
-    def name_generator(prefix):
-        for idx in itertools.count():
-            yield prefix + '_' + str(idx)
+class CUDAManager:
+    def __init__(self, known_arches, nvprune_exe):
+        self.fatbin_libs = self._known_fatbin_libs(set(CUDA_LIBRARIES))
 
-    def compute_arch(arch):
-        _, ver = arch.split('_', 1)
-        return 'compute_{}'.format(ver)
+        self.prune_args = []
+        if known_arches:
+            for arch in known_arches.split(':'):
+                self.prune_args.append('-gencode')
+                self.prune_args.append(self._arch_flag(arch))
+
+        self.nvprune_exe = nvprune_exe
 
-    libs_to_prune = set(CUDA_LIBRARIES)
+    def has_cuda_fatbins(self, cmd):
+        return bool(set(cmd) & self.fatbin_libs)
 
-    # does not contain device code, nothing to prune
-    libs_to_prune.remove('-lcudart_static')
+    def _known_fatbin_libs(self, libs):
+        libs_wo_device_code = {
+            '-lcudart_static'
+        }
+        return set(libs) - libs_wo_device_code
 
-    tmp_names_gen = name_generator('cuda_pruned_libs')
+    def _arch_flag(self, arch):
+        _, ver = arch.split('_', 1)
+        return 'arch=compute_{},code={}'.format(ver, arch)
+
+    def prune_lib(self, inp_fname, out_fname):
+        if self.prune_args:
+            prune_command = [self.nvprune_exe] + self.prune_args + ['--output-file', out_fname, inp_fname]
+            subprocess.check_call(prune_command)
+
+    def write_linker_script(self, f):
+        # This script simply says:
+        # * Place all `.nv_fatbin` input sections from all input files into one `.nv_fatbin` output section of output file
+        # * Place it after `.bss` section
+        #
+        # Motivation can be found here: https://maskray.me/blog/2021-07-04-sections-and-overwrite-sections#insert-before-and-insert-after
+        # TL;DR - we put section with a lot of GPU code directly after the last meaningful section in the binary
+        # (which turns out to be .bss)
+        # In that case, we decrease chances of relocation overflows from .text to .bss,
+        # because now these sections are close to each other
+        script = textwrap.dedent("""
+            SECTIONS {
+                .nv_fatbin : { *(.nv_fatbin) }
+            } INSERT AFTER .bss
+        """).strip()
+
+        f.write(script)
+
+
+def process_cuda_libraries(cmd, cuda_manager, build_root):
+    if not cuda_manager.has_cuda_fatbins(cmd):
+        return cmd
+
+    def tmpdir_generator(prefix):
+        for idx in itertools.count():
+            path = os.path.abspath(os.path.join(build_root, prefix + '_' + str(idx)))
+            os.makedirs(path)
+            yield path
 
-    arch_args = []
-    for arch in prune_arches.split(':'):
-        arch_args.append('-gencode')
-        arch_args.append('arch={},code={}'.format(compute_arch(arch), arch))
+    tmpdir_gen = tmpdir_generator('cuda_pruned_libs')
 
     flags = []
     cuda_deps = set()
+
+    # Because each directory flag only affects flags that follow it,
+    # for correct pruning we need to process that in reversed order
     for flag in reversed(cmd):
-        if flag in libs_to_prune:
+        if flag in cuda_manager.fatbin_libs:
             cuda_deps.add('lib' + flag[2:] + '.a')
             flag += '_pruned'
         elif flag.startswith('-L') and os.path.exists(flag[2:]) and os.path.isdir(flag[2:]) and any(f in cuda_deps for f in os.listdir(flag[2:])):
@@ -75,14 +119,12 @@ def prune_cuda_libraries(cmd, prune_arches, nvprune_exe, build_root):
             from_deps = list(cuda_deps & set(os.listdir(from_dirpath)))
 
             if from_deps:
-                to_dirpath = os.path.abspath(os.path.join(build_root, next(tmp_names_gen)))
-                os.makedirs(to_dirpath)
+                to_dirpath = next(tmpdir_gen)
 
                 for f in from_deps:
-                    # prune lib
                     from_path = os.path.join(from_dirpath, f)
                     to_path = os.path.join(to_dirpath, f[:-2] + '_pruned.a')
-                    subprocess.check_call([nvprune_exe] + arch_args + ['--output-file', to_path, from_path])
+                    cuda_manager.prune_lib(from_path, to_path)
                     cuda_deps.remove(f)
 
                 # do not remove current directory
@@ -93,7 +135,16 @@ def prune_cuda_libraries(cmd, prune_arches, nvprune_exe, build_root):
         flags.append(flag)
 
     assert not cuda_deps, ('Unresolved CUDA deps: ' + ','.join(cuda_deps))
-    return reversed(flags)
+    flags = list(reversed(flags))
+
+    # add custom linker script
+    to_dirpath = next(tmpdir_generator('cuda_linker_script'))
+    script_path = os.path.join(to_dirpath, 'script')
+    with open(script_path, 'w') as f:
+        cuda_manager.write_linker_script(f)
+    flags.append('-Wl,--script={}'.format(script_path))
+
+    return flags
 
 
 def remove_excessive_flags(cmd):
@@ -234,10 +285,12 @@ if __name__ == '__main__':
         else:
             cmd.append('-Wl,-no-pie')
 
+
     if opts.dynamic_cuda:
         cmd = fix_cmd_for_dynamic_cuda(cmd)
-    elif opts.cuda_architectures:
-        cmd = prune_cuda_libraries(cmd, opts.cuda_architectures, opts.nvprune_exe, opts.build_root)
+    else:
+        cuda_manager = CUDAManager(opts.cuda_architectures, opts.nvprune_exe)
+        cmd = process_cuda_libraries(cmd, cuda_manager, opts.build_root)
     cmd = ProcessWholeArchiveOption(opts.arch, opts.whole_archive_peers, opts.whole_archive_libs).construct_cmd(cmd)
 
     if opts.custom_step:
diff --git a/build/ymake_conf.py b/build/ymake_conf.py
index 33c55da3d5..8004d5618c 100755
--- a/build/ymake_conf.py
+++ b/build/ymake_conf.py
@@ -2272,6 +2272,13 @@ class Cuda(object):
         self.peerdirs = ['build/platform/cuda']
 
         self.nvcc_flags = [
+            # Compress fatbinary to reduce size of .nv_fatbin and prevent problems with linking
+            #
+            # Idea comes from many resources, one of them is https://discourse.llvm.org/t/lld-relocation-overflows-and-nv-fatbin/58889/6
+            # Some sources suggest using `-Xfatbin=-compress-all`, other suggest using `-Xcuda-fatbinary --compress-all`
+            # We will use the same flag as in nixpkgs
+            # (https://github.com/NixOS/nixpkgs/pull/220402/files#diff-a38e6c4e8421c03dc6c2a60c9a172ceb4059048b65798e5d4a400a7a4a5720ffR167)
+            "-Xfatbin=-compress-all",
             # Allow __host__, __device__ annotations in lambda declaration.
             "--expt-extended-lambda",
             # Allow host code to invoke __device__ constexpr functions and vice versa
@@ -2366,7 +2373,7 @@ class Cuda(object):
 
     def auto_cuda_version(self):
         if self.use_arcadia_cuda.value:
-            return '10.1'
+            return '11.4'
 
         if not self.have_cuda.value:
             return None
@@ -2398,10 +2405,6 @@ class Cuda(object):
             # do not impose any restrictions, when build not for "linux 64-bit"
             return ''
 
-        # do not include 'lto' type,
-        # because we already perform static linking
-        supported_types = ['compute', 'sm']
-
         # Equality to CUDA 11.4 is rather strict comparison
         # TODO: find out how we can relax check (e.g. to include more version of CUDA toolkit)
         if self.cuda_version.value == '11.4':
@@ -2410,19 +2413,15 @@ class Cuda(object):
             #   (these devices run only on arm64)
             # * drop support for '37'
             #   the single place it's used in Arcadia is https://a.yandex-team.ru/arcadia/sdg/sdc/third_party/cub/common.mk?rev=r13268523#L69
-            supported_vers = ['35',
-                              '50', '52',
-                              '60', '61',
-                              '70', '75',
-                              '80', '86']
+            return ':'.join(
+                ['sm_35',
+                 'sm_50', 'sm_52',
+                 'sm_60', 'sm_61',
+                 'sm_70', 'sm_75',
+                 'sm_80', 'sm_86',
+                 'compute_86'])
         else:
-            supported_vers = []
-
-        cuda_architectures = ['{typ}_{ver}'.format(typ=typ, ver=ver)
-                              for typ in supported_types
-                              for ver in supported_vers]
-
-        return ':'.join(cuda_architectures)
+            return []
 
     def auto_use_arcadia_cuda(self):
         return not self.cuda_root.from_user
@@ -2497,7 +2496,7 @@ class CuDNN(object):
         return self.cudnn_version.value in ('7.6.5', '8.0.5')
 
     def auto_cudnn_version(self):
-        return '7.6.5'
+        return '8.0.5'
 
     def print_(self):
         if self.cuda.have_cuda.value and self.have_cudnn():