До CUDA 4.0 для программирования с несколькими графическими процессорами требовалось программирование многопоточных процессоров. Это может быть сложной задачей, особенно когда вам нужно синхронизировать и / или общаться между потоками или графическими процессорами. И если весь ваш параллелизм заложен в коде вашего GPU, то наличие нескольких потоков CPU может увеличить сложность вашего программного обеспечения, не улучшая производительность дальше, чем делает GPU.
Итак, начиная с 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, и вы также не получите параллелизм между устройствами, потому что вот так: : inner_product является синхронным из-за обратного чтения для возврата результата. Таким образом, вы можете использовать cudaMemcpyAsync (конструктор device_vector будет использовать cudaMemcpy). Но более простым и, вероятно, более эффективным подходом будет использование «нулевого копирования» - прямой доступ к памяти хоста (также обсуждалось в презентации по программированию для нескольких графических процессоров, приведенной выше). Поскольку все, что вы делаете, - это читаете каждое значение один раз и добавляете его к сумме (параллельное повторное использование происходит в копии с общей памятью), вы можете также прочитать его непосредственно с хоста, а не копировать его с хоста на устройство, а затем прочитать это из памяти устройства в ядре. Кроме того, вы бы хотели асинхронный запуск ядра на каждом GPU, чтобы обеспечить максимальный параллелизм.
Вы могли бы сделать что-то вроде этого:
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];
}