Текущая реализация описания некоторой части данных, с которой осуществляются операции чтения и/или записи, выражена в виде абстрактного интерфейсного класса Си++ SharedDatum.
class SharedDatum {
public:
virtual void copyToDevice() = 0;
virtual void copyFromDevice() = 0;
bool validOnCuda, validOnHost;
};
Подразумевается, что каждый вид разделяемого куска данных должен быть описан в виде класса, наследующего SharedDatum. Структурные особенности типа shared datum заключаются в описании того, какие именно данные и каким образом копируются из системной памяти на CUDA-устройство (метод copyToDevice()) и обратно (метод copyFromDevice()). Например, в случае сплошного массива должна быть один раз вызвана функция cudaMemcpy() библиотеки CUDA с необходимыми параметрами: адреса указателей на массивы в системной памяти и памяти устройства, размер массива и направление копирования (в устройство или наоборот). Все эти параметры зафиксированы в момент создания конкретного экземпляра shared datum, благодаря чему копирующие методы класса SharedDatum не имеют дополнительных параметров. Это же верно и в общем случае — всякий экземпляр shared datum (не только представляющий массив, как в данном примере) должен содержать всю необходимую информацию о том, как копировать свои данные. Важно, что библиотечная функция-обработчик «не знает» конкретных реализаций экземпляров shared datum, с которыми имеет дело. Обработчик может вызвать методы копирования и использует лишь поля validOnCuda и validOnHost для того, чтобы определить необходимость копирования данных в CUDA-устройство или обратно. Чтобы было возможным реализовать такой обработчик, введено следующее ограничение: все экземпляры shared datum должны быть созданы на стадии инициализации и в дальнейшем оставаться неизменными. Ниже приведены два библиотечных класса, наследующих SharedDatum: SharedArray, описывающий сплошной массив, и SharedVar, описывающий одну переменную.
class SharedArray : public SharedDatum {
// указатели на массив в системной памяти и в памяти CUDA соответственно
double *arr_onHost, *arr_onCuda;
// размер каждого из массивов
size_t size;
public:
SharedArray(double *arr_onHost, double *arr_onCuda, size_t size): arr_onHost(arr_onHost), arr_onCuda(arr_onCuda), size(size) {}
void copyToDevice() {
// CUDA-вызов копирования
cudaMemcpy(arr_onCuda, arr_onHost, size, cudaMemcpyHostToDevice);
}
void copyFromDevice() {
// CUDA-вызов копирования
cudaMemcpy(arr_onHost, arr_onCuda, size, cudaMemcpyDeviceToHost);
}
};
class SharedVar : public SharedDatum {
// так называемый символ CUDA, который связан с переменной в памяти устройства
const char *sym;
// ссылка на переменную в системной памяти
double& var;
public:
SharedVar(const char *sym, double& var): sym(sym), var(var) {}
void copyToDevice() {
// CUDA-вызов, копирующий данные в устройство
cudaMemcpyToSymbol(sym, &var, sizeof(double));
}
void copyFromDevice() {
// CUDA-вызов, копирующий данные из устройства
cudaMemcpyFromSymbol(&var, sym, sizeof(double));
}
};
Пользователю предоставлена возможность описания других, более сложных, структур данных в виде классов, наследующих SharedDatum, а вышеупомянутые библиотечные классы могут служить в качестве шаблона. Описание стадии представляется интерфейсным классом Stage, куда входят указатели runOnHost и runOnDevice на функции исполнения данной стадии на центральном процессоре и CUDA-процессоре соответственно, а также два списка указателей на shared datum: тех, что считываются, и тех, что модифицируются в процессе исполнения.
class Stage {
public:
void (* const runOnHost)();
void (* const runOnDevice)();
unsigned const target; // целевая архитектура (CUDA или CPU)
SharedDatum **reads; // набор считываемых shared datum
SharedDatum **writes; // набор модифицируемых shared datum
virtual void run(); Stage(unsigned target, void (*runOnHost)(), void (*runOnDevice)(), // далее следуют два NULL-terminated списка: reads и writes ...);
};
Метод Stage::run() служит для исполнения стадии и определяется довольно просто — он вызывает либо функцию runOnHost(), либо runOnDevice() в зависимости от типа устройства исполнения (CPU или CUDA), на который указывает поле target. Однако есть возможность переопределить поведение, реализовав какой-либо иной механизм выбора устройства исполнения. Например, может потребоваться сделать несколько CUDA-вызовов в рамках одной стадии. Такое бывает при выполнении редукции, когда множество элементов, по которым происходит редукция, больше количества тредов в устройстве. Такого определения Stage достаточно для реализации общего механизма, обеспечивающего исполнение всех стадий на тех типах устройств, которые задал пользователь, и экономичное (то есть минимально необходимое) копирование данных при смене типа устройства в рамках каждого потока исполнения. Это реализовано в библиотечной функции, имеющий следующий прототип:
void runStage(Stage *);
Теперь пользовательский код может выглядеть так:
// массив с описанием всех стадий программы
Stage **stages = { ..., NULL };
// Тело главного цикла
{
...
Stage **s = stages;
while(*s)
runStage(*s++); // исполнение очередной стадии (включая копирование)
...
}
Здесь важно, что при вызове исполнения стадии не передается дополнительных параметров — все особенности каждого запуска «спрятаны» в описании стадии и установлены с помощью конструктора, то есть функция runStage() «не заглядывает» в особенности реализации каждой стадии. В текущей реализации библиотеки при создании экземпляра Stage необходимо указывать целевую архитектуру (явно установить поле target). После чего данная стадия всегда будет исполняться на заданном устройстве. Тем не менее, Stage сконструирована так, что она потенциально умеет исполняться на любом (CUDA или CPU) устройстве. Это можно использовать для реализации динамической балансировки, что заложено в планах по дальнейшему развитию библиотеки Centaur. Динамическая балансировка допускает исполнение стадии на любом устройстве или даже на устройствах разных типов одновременно с целью минимизировать полное время исполнения стадии. Разумеется, для этого должны быть реализованы механизмы, позволяющие определить, сколько тредов и на каких устройства целесообразно запускать. В простых случаях пользователь может явно указать долю тредов, которые следует запустить на CUDA-устройстве. В этом случае конструктор может выглядеть следующим образом:
StageLB::StageLB(float targetCUDA, void (*runOnHost)(), void (*runOnDevice)(), ...);
Здесь параметр targetCUDA (вместо target) указывает долю тредов, которые следует исполнять на CUDA-устройстве, оставшиеся же будут исполнены на CPU. Поскольку, по определению стадии, зависимости по данным между тредами отсутствуют, такой подход прост в реализации и достаточно эффективен, если параметр targetCUDA подобран правильно. |