Как лучше всего минимизировать сложность кода при сохранении данных стека для использования параллелизма?

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

Я нашел и распараллелил функцию, скрытую в нескольких функциях/циклах. На эту функцию приходится ~ 98% времени обработки, но она не использует достаточно параллелизма, чтобы быть полезной (порядка пары блоков..). При одновременном выполнении код работает намного быстрее. Однако в результате я вынужден поддерживать большой список объектов стека, которые я должен перебирать несколько раз, см. код ниже:

void do_work(int i, ...) {
    // computationally expensive stuff...
}

void prereq_stuff(int i) {
    int foo;
    double bar;

    // lots of big divergent control structures...

    do_work(i); // maybe arrive here..

    // output and what not that needs foo/bar...
}

int main() {

    for (int i = 0; i < BIG_NUMBER; i++) {
        prereq_stuff(i);
    }

    return 0;
}

Превратился в...

// a struct that contains all the stack data..
struct StackMem {
    int foo;
    double bar;
};

void do_work_on_gpu(List<StackMem> contexts) {
    // launch a kernel to handle to expensive stuff..
}

void prereq_stuff(StackMem* context, int i) {
    // maybe queue up data for do_work_on_gpu()...
}

void cleanup_stuff(StackMem* context, int i) {
    // output and what not that needs foo/bar...
}

int main() {

   List<StackMem> contexts; // some container of stack objects

   for (int i = 0; i < BIG_NUMBER; i++) {
        StackMem* context = contexts.add();
        prereq_stuff(context, i);
    }

    do_work_on_gpu(contexts); // calls the CUDA kernel

    for (int i = 0; i < contexts.size(); i++) {
        cleanup_stuff(contexts.get(i), i);
    }

    return 0;
}

Есть ли какая-то конструкция/шаблон дизайна, которую я могу здесь использовать? Или это настолько просто, насколько это возможно, когда все данные для вызова do_work() доступны одновременно?


person Brandon Wilson    schedule 17.07.2017    source источник
comment
Это зависит от кода/алгоритма, лежащего в основе множества больших расходящихся управляющих структур. Вы написали, что процессорная версия требовала около 98% времени обработки. Каково соотношение (частично) версии GPU? И как долго работает prereq_stuff по сравнению с do_work on_gpu? Я не уверен, что есть хороший общий ответ, и этот вопрос лучше задать на Code Review.   -  person BlameTheBits    schedule 18.07.2017
comment
@Shadow Обязательная ссылка на справочный центр проверки кода, чтобы избежать неправильно размещенных вопросов.   -  person Mast    schedule 18.07.2017
comment
@Mast: я не думаю, что это настоящий код, это больше похоже на набросок кода.   -  person einpoklum    schedule 18.07.2017
comment
@einpoklum Точно.   -  person Mast    schedule 18.07.2017


Ответы (1)


Вот как бы я подошел к этому сценарию:

  • Ваша структура данных контекста будет либо

    struct { int* foos; double* bars; } contexts;
    

    or

    struct context_t { int foos; double bar; };
    struct context_t* contexts;
    

    (или что-то подобное, но более красивое, например, используйте gsl::span_t<context_t>, если вы знаете, что это такое)

  • Вы будете предварительно выделять память для всех ваших контекстов, желательно закрепляя память с помощью cudaMallocHost() (поэтому я не предлагал std::vector).

  • Вместо создания контекста при каждом вызове prereq_stuff() просто передайте ему итератор конкретного контекста или индекс и массив контекстов, чтобы он не помещал свой контекст в какое-то произвольное место в памяти.

  • Передайте массив(ы) контекста, используя один cudaMemcpy().

  • Каждый поток, или блок, или что-то еще использует свой идентификатор, чтобы решить, с каким контекстом работать.

  • Будьте осторожны при объединении чтений, так как контекст составляет 12 байт, а чтение памяти на графическом процессоре NVIDIA осуществляется в единицах по 32 байта (через кеш текстур/__ldg()) или 128 байт.

PS. Что бы вы ни реализовали, убедитесь, что вы создали профиль свой код и посмотрите на время вычислений и ввода-вывода, а также на время, которое вы тратите на работу ЦП, чтобы увидеть, какой у вас есть потенциал для улучшения.

person einpoklum    schedule 18.07.2017