Запуск операций копирования асинхронной памяти на нескольких графических процессорах

Я хочу разделить & скопируйте массив данных на хосте в память устройства с несколькими графическими процессорами. Кроме того, я хочу сделать все эти операции копирования одновременно.

Для этого я использую cudaMemcpyAsync, который я запускаю в приватном потоке каждого графического процессора.

Вот что я делаю (сомнения в коде отмечены комментариями, начинающимися с ??)

#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. Неявная синхронизация

Две команды из разных потоков не могут выполняться одновременно, если они есть
из следующих операций выдается хостом между ними
нить:

Allocation выделение памяти хоста с блокировкой страницы,

Allocation распределение памяти устройства,

Memory набор памяти устройства,

Copy копия памяти между двумя адресами в одну и ту же память устройства,

‣ любая команда 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