ТЕХНИЧЕСКОЕ ЗАДАНИЕ: Написать pinned memory аллокатор

1. Цель задачи

Научиться использовать закреплённую (pinned) память для ускорения передачи данных между CPU и GPU. Реализовать собственный аллокатор pinned memory на основе cudaHostAlloc, организовать асинхронное копирование host-to-device (H2D) с перекрытием (overlap) вычислений и передачи данных.

Ключевой результат Рабочий аллокатор pinned памяти + демонстрационный скрипт, который измерит ускорение при overlap вычислений и передачи по сравнению с синхронным подходом.

2. Исходные данные

Перед началом необходимо иметь:

Если нет реального GPU — симулируем:

1. Установить WSL2 с CUDA (если GPU нет физически, но есть поддержка WSL) — или использовать облачный GPU (Google Colab, AWS) с Jupyter.
2. В крайнем случае — написать заглушку, которая эмулирует pinned память через mmap с флагом MAP_LOCKED, и асинхронную передачу через потоки Python. Но для получения реальных метрик GPU обязателен.

3. Технологический стек

4. Этапы выполнения

Этап 1: Настройка окружения и проверка GPU (30 минут)

Действия

1. Установить CUDA Toolkit
   Скачать и установить с официального сайта. Проверить версию:

   nvcc --version
   

2. Установить PyCUDA

   pip install pycuda
   

3. Проверить доступность GPU и базовые свойства
   Написать скрипт:

   import pycuda.driver as cuda
   import pycuda.autoinit
   print(f"Device: {cuda.Device(0).name()}")
   print(f"Total memory: {cuda.Device(0).total_memory() / 1e9:.2f} GB")
   

   Убедиться, что код выполняется без ошибок.

Ожидаемый результат этапа Рабочее окружение, определён GPU, установлены библиотеки.

Этап 2: Реализация аллокатора pinned memory (1 час)

Действия

1. Написать класс PinnedAllocator, который использует cudaHostAlloc
   В PyCUDA это делается через cuda.host_alloc() с флагами:

   import pycuda.driver as cuda
   import numpy as np
   
   class PinnedAllocator:
       def __init__(self, size_bytes):
           self.size = size_bytes
           # Выделяем pinned память (флаг 0 — стандартный, можно также cudaHostAllocMapped, но используем default)
           self.ptr = cuda.host_alloc(size_bytes, flags=0)  # возвращает указатель
           self.host_arr = np.frombuffer(self.ptr, dtype=np.float32)
       
       def __del__(self):
           cuda.host_free(self.ptr)
       
       def get_array(self, shape):
           return np.frombuffer(self.ptr, dtype=np.float32).reshape(shape)
   

   Альтернатива (более низкоуровневая): написать на C++ с cudaHostAlloc и cudaFreeHost.

2. Добавить тестовую аллокацию
   Выделить 1 ГБ, записать данные, проверить, что память закреплена (сравнить скорость передачи в/из GPU с обычной malloc).

3. Сравнить с обычной памятью (pageable)
   Выделить через numpy.zeros и передать на GPU — замерить время cuda.memcpy_htod. То же самое сделать с pinned — убедиться, что pinned быстрее.

Ожидаемый результат этапа Класс PinnedAllocator, который выделяет память, доступную для асинхронных копий. Измерение пропускной способности (как минимум 2x-3x быстрее pageable).

Этап 3: Асинхронная копия H2D с использованием streams (1–1.5 часа)

Действия

1. Создать stream (поток CUDA)

   stream = cuda.Stream()
   

2. Реализовать асинхронное копирование

   # Предположим, есть device_array (cuda.mem_alloc(size_bytes))
   cuda.memcpy_htod_async(device_array, host_ptr, stream=stream)
   

3. Синхронизация
   Поток синхронизируется через stream.synchronize().

4. Проверить коррекность
   Скопировать данные обратно (синхронно) и сравнить с исходным массивом.

Ожидаемый результат этапа Функция, копирующая данные на GPU асинхронно, без блокировки основного потока CPU.

Этап 4: Демонстрация overlap вычислений и передачи (1.5–2 часа)

Действия

1. Разработать сценарий с двумя потоками CUDA

   • Stream 1: копирование данных (H2D)
   • Stream 2: вычисления на GPU (например, SAXPY или матричное умножение) на уже переданных данных.

   Важно Вычисления должны запускаться на данных, которые уже на устройстве (или копироваться в один stream, а вычисляться в другом — с корректной синхронизацией через события).

2. Структура кода

   import numpy as np
   import pycuda.autoinit
   import pycuda.driver as cuda
   from pycuda.compiler import SourceModule
   
   # Компиляция простого ядра: умножение массива на скаляр
   mod = SourceModule("""
       __global__ void scale(float *a, float s, int n) {
           int idx = threadIdx.x + blockIdx.x * blockDim.x;
           if (idx < n) a[idx] *= s;
       }
   """)
   scale_kernel = mod.get_function("scale")
   
   # Выделение памяти
   N = 10**7
   pinned_alloc = PinnedAllocator(N * 4)
   host_arr = pinned_alloc.get_array((N,))
   host_arr[:] = 1.0
   
   device_arr = cuda.mem_alloc(N * 4)
   
   # Синхронный случай (для сравнения)
   start = cuda.Event()
   end = cuda.Event()
   start.record()
   cuda.memcpy_htod(device_arr, host_arr)
   scale_kernel(device_arr, np.float32(2.0), np.int32(N), block=(256,1,1), grid=(N//256+1,1))
   cuda.memcpy_dtoh(host_arr, device_arr)
   end.record()
   end.synchronize()
   sync_time = cuda.event_elapsed_time(start, end)
   
   # Асинхронный с overlap
   stream1 = cuda.Stream()
   stream2 = cuda.Stream()
   host_arr[:] = 1.0   # сброс
   
   start.record()
   cuda.memcpy_htod_async(device_arr, host_arr, stream=stream1)
   # Можно сразу запустить ядро на stream2 — оно будет ждать окончания копии через события?
   # Правильнее: поставить событие после копии в stream1, затем ядро в stream2 ждёт это событие.
   event = cuda.Event()
   stream1.record_event(event)
   stream2.wait_for_event(event)
   scale_kernel(device_arr, np.float32(2.0), np.int32(N), block=(256,1,1), grid=(N//256+1,1), stream=stream2)
   stream2.synchronize()
   end.record()
   end.synchronize()
   async_time = cuda.event_elapsed_time(start, end)
   
   print(f"Sync time: {sync_time:.3f} ms")
   print(f"Async time (overlap): {async_time:.3f} ms")
   # При корректном overlap async_time должно быть меньше sync_time, т.к. передача и вычисления частично перекрываются.
   

3. Анализ

   • Если async_time > sync_time — возможно, нет реального перекрытия (проверить размеры, количество потоков, использовать nsys profile).
   • Убедиться, что для overlap необходима pinned память (с pageable асинхронные копии работают как синхронные).
   • Построить график времени от размера данных.

Ожидаемый результат этапа Скрипт, который демонстрирует измеримое снижение общего времени выполнения при overlap (например, для массива 100 МБ async быстрее на 20-40%).

Этап 5: Сравнение производительности и оформление (30 минут)

Действия

1. Зафиксировать метрики

   • Пропускная способность pinned vs pageable
   • Синхронное vs асинхронное время
   • Доля overlap (можно оценить через профилировщик)

2. Оформить результаты
   Написать краткий отчёт в README.md: цель, метод, результаты в виде таблицы, выводы.

Ожидаемый результат этапа Файл отчёта с цифрами и сравнительными графиками (опционально).

5. Критерии приемки (Definition of Done)

• Написан класс PinnedAllocator, корректно выделяющий и освобождающий закреплённую память.
• Реализована асинхронная копия H2D с использованием cudaMemcpyAsync и cudaStream_t.
• Демонстрационный скрипт запускается без ошибок на целевой системе.
• Скрипт выводит время синхронной и асинхронной версий для размера данных ≥ 10⁷ float32.
• Время асинхронной версии (с overlap) меньше времени синхронной как минимум на 15% (из-за перекрытия передачи и вычислений).
• Код содержит комментарии, объясняющие ключевые моменты (выделение pinned памяти, создание stream, синхронизация событий).
• Приложен файл README.md с результатами тестирования и инструкцией по запуску.

6. Ожидаемый результат

Основной артефакт Файл pinned_allocator_demo.py (или аналогичный), содержащий:

• Класс PinnedAllocator с методами alloc, free, get_array.
• Функцию compare_sync_async(size), которая выводит времена.
• Блок if __name__ == "__main__": для запуска демонстрации.

Дополнительно (опционально):

• Jupyter Notebook с визуализацией зависимости времени от размера.
• Makefile для компиляции (если используется C++).
• Профилировочный лог (файл .nsys-rep) с видимым overlap.

7. Возможные сложности и их решение

8. Бюджет времени (оценка)

Примечание Для первого раза опытному инженеру потребуется около 5 часов. Новичку — до 8 часов с учётом чтения документации.

9. Связанные вопросы из базы знаний

10. Чек-лист самопроверки

• Я проверил, что cudaHostAlloc действительно возвращает указатель в закреплённой памяти (через numpy.frombuffer не падает с ошибкой).
• Я убедился, что асинхронная копия не блокирует CPU-поток (можно вставить time.sleep и наблюдать, что копия завершается без ожидания).
• Я сравнил время с pageable памятью и получил значительное ускорение для больших массивов.
• Я проверил, что при overlap общее время меньше суммы времени копии и вычислений (т.е. перекрытие есть).
• Я протестировал освобождение памяти (деструктор __del__ / host_free) — нет утечек.