В предыдущих разделах были описаны основные циклы задач вычислительной механики, решаемых сеточными методами. В текущем разделе рассмотрим методику распараллеливания данных циклов.
На общей памятиРаспараллеливание циклов, витки которых независимы, на системе с общей памятью выполнить относительно просто.
Система OpenMP предназначена для такого вида распараллеливания и поэтому достаточно добавить директиву:
#pragma omp parallel for
for (n = 0; n < N; n++) {
A_new[n] = f(A_old[n]);
}
При исполнении система OpenMP породит некоторое количество нитей (которое зависит как от аппаратной конфигурации вычислительной системы, так и от настроек запуска), которые будут в параллель выполнять части данного цикла.
На CUDAРаспараллеливание на графических ускорителях такого вида циклов тоже относительно просто. Для переноса на CUDA необходимо выделить тело цикла в виде вспомогательной функции (назовем ее iter) и объявить ее с модификатором __global__. Цикл в ней будет неявным — точнее, превратится в параллельный CUDA-вызов.
__global__ void iter() {
int n = blockIdx.x * blockDim.x + threadIdx.x;
if (n < N) { // проверка необходима, поскольку количество реально запускаемых нитей кратно числу блоков и поэтому может превосходить N
A_new[n] = f(A_old[n]);
}
}
Тогда вызов цикла, исполняемого параллельно на множестве CUDA-ядер, будет выглядеть так:
iter<<<nBlocks, nThreads>>>();
Параметры nBlocks и nThreads выбираются таким образом, чтобы устройство было максимально загружено, и при этом чтобы хватило ресурсов. А именно: устройство имеет ограничение на максимальное количество тредов в одном блоке, и nThreads не может его превосходить; кроме того, есть ограничение по количеству регистровой памяти, в которую помещаются стековые переменные, — если тредов в блоке слишком много, то на всех может не хватить регистров, и тогда придется еще уменьшить nThreads.
Для удобства рекомендуется выделить виток цикла в одну функцию в отдельный заголовочный файл (например, iter_n.h):
DEVICE void iter_n(int n) {
A_new[n] = f(A_old[n]);
}
Тогда один и тот же код можно использовать для генерации кода как для ускорителя, так и для центрального процессора. Достаточно лишь подставить нужное определение препроцессорной переменной DEVICE — оно должно быть пустым для CPU-версии и определяться как __device__ для CUDA-версии. Дополнительное удобство такого подхода состоит в том, что обе версии кода хранятся в виде общего источника, что помогает поддерживать эквивалентность обеих версий.
Теперь вызывающий код выглядит следующим образом:
// CPU-версия
#define DEVICE #include "iter_n.h" // в этом файле код «витка»
#pragma omp parallel for
for (n = 0; n < N; n++) {
iter_n(n);
} // CUDA-версия: декларация витка цикла как функции #define DEVICE __device__ #include "iter_n.h" // в этом файле код «витка»
__global__ void iter() {
int n = blockIdx.x * blockDim.x + threadIdx.x;
if (n < N) {
iter_n(n); // вызов «витка»
}
}
// ...и ее CUDA-вызов iter<<<nBlocks, nThreads>>>();
С помощью средств препроцессора можно сделать такой вызывающий код универсальным для всех функций того же типа, что iter().
На распределенной памятиВ случае распределенной памяти недостаточно просто распараллелить работу — витки циклов — между процессами. Необходимо распределить данные и организовать работу по перемещению данных между узлами вычислительной системы.
Описанная организация данных и циклов позволяет сделать такое разделение с наименьшими затратами.
Для разделения данных удобно использовать способ распараллеливания циклов методом «Преобразования массивов», описанный в разделе Способы распараллеливания циклов. В данном способе изменяется смысл массивов, описывающих структуру сетки, и массивов, хранящих данные. Вместо того, чтобы хранить данные для всей сетки, эти массивы хранят данные только для части сетки, которая расположена на данном узле.
Использование массивов и описанных выше методов позволяет сделать такую подмену без изменения кода циклов, “незаметно” для программиста. В результате, цикл интерпретируется по-новому: вместо последовательного кода, который обрабатывает все ячейки сетки, данный цикл исполняется несколькими процессами (как принято в MPI или SHMEM), каждый процесс исполняет данный цикл для своей части сетки.
Следует отметить, что такого преобразования программы недостаточно, для корректного распараллеливания. Например, при вычислении нового значения в ячейке сетки могут потребоваться старые значения в соседних ячейках, которые могут быть расположены на других узлах вычислительной системы.
В общем случае автоматически понять, какие данные необходимо передать, трудно или невозможно. Однако, пользователь зачастую знает, какие данные нужны для вычислений. Поэтому, необходимо попросить пользователя передать эту информацию библиотеке: записать эту информацию в виде специальных вызовов или псевдокомментариев. Причем, может быть, с некоторым запасом: корректность не будет нарушена, если пользователь укажет большее число данных, чем нужно на самом деле. Однако, если пользователь укажет слишком много данных (например, все данные - вся сетка), то это будет корректно, но не эффективно.
ЗаключениеСледующим этапом распараллеливания является выделение коммуникационных шаблонов каждого цикла: какие данные и в каких соседних ячейках необходимы для вычисления данных в заданной ячейке. Определение способа задания шаблонов, необходимых для задач из данной области. И реализация передачи данных согласно указанным шаблонам. |