Вычисление Cuda на нескольких GPU

До CUDA 4.0 программирование на нескольких графических процессорах требовало многопоточного программирования на центральном процессоре. Это может быть сложно, особенно когда вам нужно синхронизировать и / или обмениваться данными между потоками или графическими процессорами. И если весь ваш параллелизм заключен в коде графического процессора, то наличие нескольких потоков процессора может усложнить ваше программное обеспечение без повышения производительности, выходящего за рамки того, что делает графический процессор.

Итак, начиная с CUDA 4.0, вы можете легко запрограммировать несколько графических процессоров из однопоточной хост-программы. Вот несколько слайдов об этом, которые я представил в прошлом году.

Программирование нескольких графических процессоров может быть таким простым:

int numDevs = 0;
cudaGetNumDevices(&numDevs);
...
for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    kernel<<<blocks, threads>>>(args);
}

Для вашего конкретного примера скалярных произведений вы можете использовать thrust::inner_product в качестве отправной точки. Я бы сделал это для прототипа. Но посмотрите мои комментарии в конце об узких местах с пропускной способностью.

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

// assume the deviceIDs of the two 2050s are dev0 and dev1.
// assume that the whole vector for the dot product is on the host in h_data
// assume that n is the number of elements in h_vecA and h_vecB.

int numDevs = 0;
cudaGetNumDevices(&numDevs);
...
float result = 0.f;
for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    device_vector<float> vecA(h_vecA + d*(n/d), h_vecA + (d+1)*(n/d)-1);
    device_vector<float> vecB(h_vecB + d*(n/d), h_vecB + (d+1)*(n/d)-1);
    result += thrust::inner_product(vecA.begin(), vecA.end(), vecB.begin(), 0.f);
}

(Я признаю, что указанная выше индексация неверна, если n не является четным числом, кратным numDevs, но я оставлю это исправление в качестве упражнения для читателя. :)

Это просто, и это отличное начало. Сначала заставьте его работать, а затем оптимизируйте.

После того, как он заработает, если все, что вы делаете на устройствах, - это точечные продукты, вы обнаружите, что ваша пропускная способность ограничена - в основном PCI-e, и вы также не получите параллелизма между устройствами, потому что thust :: _ 4_ синхронный из-за обратного чтения для возврата результата. Итак, вы можете использовать cudaMemcpyAsync (конструктор device_vector будет использовать cudaMemcpy). Но более простым и, вероятно, более эффективным подходом было бы использование «нулевой копии» - прямой доступ к памяти хоста (это также обсуждается в презентации программирования с несколькими графическими процессорами, ссылка на которую приведена выше). Поскольку все, что вы делаете, - это чтение каждого значения один раз и добавление его к сумме (параллельное повторное использование происходит в копии общей памяти), вы также можете читать его напрямую с хоста, а не копировать его с хоста на устройство, а затем читать это из памяти устройства в ядре. Кроме того, вам может потребоваться асинхронный запуск ядра на каждом графическом процессоре, чтобы обеспечить максимальный параллелизм.

Вы можете сделать что-то вроде этого:

int bytes = sizeof(float) * n;
cudaHostAlloc(h_vecA, bytes, cudaHostAllocMapped | cudaHostAllocPortable);
cudaHostAlloc(h_vecB, bytes, cudaHostAllocMapped | cudaHostAllocPortable);
cudaHostAlloc(results, numDevs * sizeof(float), cudaHostAllocMapped | cudaHostAllocPortable);
// ... then fill your input arrays h_vecA and h_vecB


for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    cudaEventCreate(event[d]));
    cudaHostGetDevicePointer(&dptrsA[d], h_vecA, 0);
    cudaHostGetDevicePointer(&dptrsB[d], h_vecB, 0);
    cudaHostGetDevicePointer(&dresults[d], results, 0);
}

...

for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    int first = d * (n/d);
    int last   = (d+1)*(n/d)-1;
    my_inner_product<<<grid, block>>>(&dresults[d], 
                                      vecA+first, 
                                      vecA+last, 
                                      vecB+first, 0.f);
    cudaEventRecord(event[d], 0);
}

// wait for all devices
float total = 0.0f;
for (int d = 0; d < devs; d++) {
    cudaEventSynchronize(event[d]);
    total += results[numDevs];
}

Чтобы создать несколько потоков, вы можете использовать OpenMP или pthreads. Чтобы сделать то, о чем вы говорите, кажется, что вам нужно создать и запустить два потока (параллельная секция omp или pthread_create), чтобы каждый выполнял свою часть вычислений и сохранял свой промежуточный результат в отдельных переменных для всего процесса. (напомним, глобальные переменные автоматически распределяются между потоками процесса, поэтому исходный поток сможет видеть изменения, сделанные двумя порожденными потоками). Чтобы исходные потоки ожидали завершения других, выполните синхронизацию (используя глобальный барьер или операцию соединения потоков) и объедините результаты в исходном потоке после завершения двух порожденных потоков (если вы разделяете массивы пополам и вычисляя скалярное произведение путем умножения соответствующих элементов и выполнения глобального суммирования половин, необходимо только добавить два промежуточных результата из двух порожденных потоков).

Вы также можете использовать MPI или fork, и в этом случае связь может осуществляться аналогично сетевому программированию ... каналы / сокеты или связь и синхронизация через (блокировку) отправку и получение.