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

Question

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

__device__ __constant__ CaseParams deviceCaseParams;

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

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

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

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

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

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

Answer 1

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

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

Answer 2

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

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

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