Я хочу разделить и скопировать массив данных с хоста на память устройства нескольких gpu. Кроме того, я хочу выполнять все эти операции копирования одновременно.
Для этого я использую cudaMemcpyAsync, который запускаю в приватном потоке каждого GPU.
Вот что я делаю (Сомнения в коде отмечены комментариями, начинающимися с ?? )
#define SIZE 1000
#define GPUCOUNT 2
int* hostData = nullptr;
int *devData[GPUCOUNT];
cudaStream_t stream[GPUCOUNT];
// Create one stream per GPU
for ( int i=0; i != GPUCOUNT ; ++i )
{
// DO I need to call cudaSetDevice before creating stream for each GPU ??
cudaStreamCreate(&stream[i]));
}
// Allocate pinned data on host
cudaMallocHost (&hostData, SIZE );
// Allocate data on each device and copy part of host data to it
for( int i=0; i != GPUCOUNT ; ++i )
{
cudaSetDevice(i);
cudaMalloc( (void**) &devData[i], sizeof(int) * SIZE/GPUCOUNT ); // ?? Does blocking behavior of cudamalloc prevents asynch memcpy invoked in stream of other GPUs from running concurrently
cudaMemcpyAsync( (void*) devData[i], hostData + i*SIZE/GPUCOUNT, SIZE/GPUCOUNT, cudaMemcpyHostToDevice, stream[i] );
}
// Some CPU code while copy is happening
// ....
// Wait for copy on all streams to finish
cudaDeviceSynchronize();
// Do something else
Когда я читаю руководство по программированию на C, я вижу, что описанные выше операции копирования памяти не будут выполняться асинхронно, потому что между двумя последовательными запусками асинхронного копирования памяти я вызываю операцию хоста, которая выделяет память устройства (блокирующий вызов).
3.2.5.5.4. Неявная синхронизация
Две команды из разных потоков не могут выполняться одновременно, если какая-либо из следующих операций выполняется между ними хост-потоком:
‣ выделение памяти хоста с блокировкой страницы,
‣ выделение памяти устройства,
‣ набор памяти устройства,
‣ копирование памяти между двумя адресами в одну и ту же память устройства,
‣ любую команду CUDA в поток по умолчанию,
Если приведенная выше причина кажется правдой, мне нужно разделить выделение памяти и операцию копирования.
// Allocate data on each device
for( int i=0; i != GPUCOUNT ; ++i )
{
cudaSetDevice(i);
cudaMalloc( (void**) &devData[i], sizeof(int) * SIZE/GPUCOUNT );
}
// Copy part of host data to each device
for( int i=0; i != GPUCOUNT ; ++i )
{
// ?? DO I need to call cudaSetDevice before memory copy ??
// CUDA guide says:"A memory copy will succeed even if it is issued to a stream that is not associated to the current device."
cudaMemcpyAsync( (void*) devData[i], hostData + i*SIZE/GPUCOUNT, SIZE/GPUCOUNT, cudaMemcpyHostToDevice, stream[i] );
}
Действителен ли мой приведенный выше анализ?
Кроме того, нельзя ли сделать это без создания явного потока для каждого графического процессора путем запуска операции cudaMemcpyAsync в потоке по умолчанию (идентификатор потока 0) каждого графического процессора?. Я основываю это на следующих утверждениях, сделанных в руководстве по программированию CUDA C:
Каждое устройство имеет свой собственный поток по умолчанию (см. Поток по умолчанию), поэтому команды, отправленные в поток по умолчанию устройства, могут выполняться не по порядку или одновременно с командами, отправленными в поток по умолчанию любого другого устройства.
Тогда код будет выглядеть так
#define SIZE 1000
#define GPUCOUNT 2
int* hostData = nullptr;
int *devData[GPUCOUNT];
// Allocate pinned data on host
cudaMallocHost (&hostData, SIZE );
// Allocate data on each device
for( int i=0; i != GPUCOUNT ; ++i )
{
cudaSetDevice(i);
cudaMalloc( (void**) &devData[i], sizeof(int) * SIZE/GPUCOUNT );
}
// Copy part of host data to each device
for( int i=0; i != GPUCOUNT ; ++i )
{
// ?? DO I need to call cudaSetDevice before memory copy ??
// CUDA guide says:"A memory copy will succeed even if it is issued to a stream that is not associated to the current device."
cudaMemcpyAsync( (void*) devData[i], hostData + i*SIZE/GPUCOUNT, SIZE/GPUCOUNT, cudaMemcpyHostToDevice, 0 );
}
// Some CPU code while copy is happening
// ....
// Wait for copy on all streams to finish
cudaDeviceSynchronize();
// Do something else
