Функции __device__ и __host__

Ядра являются точками входа, но большая часть кода находится во вспомогательных функциях. CUDA предоставляет ещё два квалификатора, указывающих, где может выполняться обычная функция. Познакомимся с ними.

Функция __device__ выполняется на GPU и может вызываться только из другого кода GPU, например из ядра или другой функции устройства.

Считайте функции __device__ набором инструментов GPU. Они делают ядра более понятными, храня логику, которую повторно используют потоки устройства.

Квалификатор __host__ означает обычный код CPU. Он используется по умолчанию, поэтому обычные функции C++ уже являются функциями хоста.

Если не указывать квалификаторы, функция по умолчанию предназначена только для хоста. Поэтому обычный C++ по-прежнему компилируется внутри файла .cu.

Объедините два квалификатора, и функция будет скомпилирована для обеих сторон. nvcc создаст одну версию для CPU и одну для GPU.

Пометка вспомогательной функции как __host__ __device__ позволяет использовать одну и ту же математику в эталонном коде CPU и внутри ядер без дублирования. ✨

Код хоста не может вызывать функцию __device__, а код устройства не может вызывать функцию, предназначенную только для хоста. Компилятор контролирует эту границу.

Функцию __device__ нельзя запускать. Только ядра __global__ принимают запуск с тройными угловыми скобками от хоста.

Компилятор часто встраивает небольшие функции __device__, поэтому выделение логики во вспомогательные функции обычно ничего не стоит во время выполнения.

Выбирайте её по месту выполнения: __device__ — для вспомогательных функций GPU, __host__ — для кода CPU, а оба квалификатора — когда логика общая.

Вы узнали, что __device__ выполняется на GPU, __host__ — на CPU, значение по умолчанию — хост, а объединение обоих квалификаторов позволяет использовать одну функцию в обоих мирах. 🎯