Как использовать постоянную память CUDA у программиста приятным способом?

Я работаю над хрустящим приложением, использующим CUDA. У меня есть некоторые статические данные, которые должны быть доступны для всех потоков, поэтому я поместил их в постоянную память следующим образом:

__device__ __constant__ CaseParams deviceCaseParams;

Я использую вызов cudaMemcpyToSymbol для переноса этих параметров с хоста на устройство:

void copyMetaData(CaseParams* caseParams)
{
    cudaMemcpyToSymbol("deviceCaseParams", caseParams, sizeof(CaseParams));
}

который работает.

В любом случае, кажется (по пробной версии и ошибкам, а также из чтения сообщений в сети), что по какой-то больной причине объявление deviceCaseParams и его копирование (вызов cudaMemcpyToSymbol) должны быть в одном файле, На данный момент у меня эти два файла .cu, но я действительно хочу иметь структуру параметров в файле .cuh, чтобы любая реализация могла видеть это, если захочет. Это означает, что я также должен иметь функцию copyMetaData в заголовочном файле, но это испортит связывание (уже определенный символ), так как оба файла .cpp и .cu включают этот заголовок (и, таким образом, компилятор MS С++ и nvcc компилирует его).

Есть ли у кого-нибудь советы по дизайну здесь?

Обновление: См. комментарии

Ответ 1

С обновленным CUDA (например, 3.2) вы должны иметь возможность делать memcpy из другой единицы перевода, если вы ищете символ во время выполнения (то есть, передавая строку в качестве первого аргумента arg cudaMemcpyToSymbol, как вы в своем примере).

Кроме того, с помощью устройств класса Fermi вы можете просто разделить память (cudaMalloc), скопировать в память устройства и передать аргумент как указатель const. Компилятор будет распознавать, если вы равномерно распределяете данные по искажениям, и если так будет использоваться постоянный кеш. Дополнительную информацию см. В Руководстве по программированию CUDA. Примечание: вам нужно скомпилировать с помощью -arch=sm_20.

Ответ 2

Если вы используете pre-Fermi CUDA, вы обнаружите, что эта проблема относится не только к постоянной памяти, она применима ко всему, что вы хотите на стороне CUDA. Единственные два пути, которые я нашел вокруг этого, - либо:

Напишите все CUDA в одном файле (.cu) или
Если вам нужно разбить код на отдельные файлы, ограничьте себя заголовками, которые затем включите ваш единственный файл .cu.

Если вам нужно разделить код между CUDA и C/С++ или иметь общий код, который вы разделяете между проектами, вариант 2 является единственным выбором. С самого начала кажется очень неестественным, но это решает проблему. Вы по-прежнему можете структурировать свой код, просто не в обычном режиме. Главными издержками являются то, что каждый раз, когда вы делаете сборку, вы скомпилируете все. Положительная сторона этого (по-моему, возможно, почему так работает) заключается в том, что компилятор CUDA имеет доступ ко всему исходному коду в одном ударе, который хорош для оптимизации.