Надеюсь, я пишу в нужном месте. Я пытаюсь переместить данные из одного графического процессора в другой. (Я должен уточнить, что я не хочу использовать никакой cuMemcpy. В некотором смысле, я воспроизводлю это, что происходит с cuMemcpyPToP.)
Используя унифицированную адресацию, можно написать простое ядро
[…]
if ( i < m ) w[i] = v[i];
[…]
, где v и w - массивы размера m, а i - соответствующий индекс.
Теперь рассмотрим два устройства dev_0 и dev_1, массивы v и w размещены в dev_0, dev_1 соответственно. В этом случае при соответствующей настройке (см. Ниже) это ядро работает. Но проблема в том, как уведомить цель, то есть dev_1, о том, что данные на самом деле находятся в ее памяти? Для этого мы использовали процедуры события и ожидания. К сожалению, этот механизм не работает: если вы печатаете данные на dev_1, иногда все считанные данные являются правильными, иногда только их часть, как показано в следующем выводе,
b[0] = 0.00000
b[1] = 0.00000
b[2] = 0.00000
b[3] = 0.00000
b[4] = 0.86748
b[5] = 0.83402
b[6] = 0.81273
[…]
Из этой печати мы Можно предположить, что когда dev_1 читает элементы с 0 по 3, данные не видны, не зафиксированы или просто еще не получены. Затем данные становятся видимыми и могут быть доступны, так что остальные элементы не равны нулю.
Теперь мои вопросы: 1 / Мы записываем событие, связанное с движущимся ядром данных. Когда ядро завершило работу, означает ли это, что данные отправлены dev_0 или данные получены dev_1? 2 / Как мы можем полагаться на драйвер CUDA для получения события / уведомления о том, что данные были записаны в dev_1? 3 / Любое другое предложение?
Спасибо,
Вот код настройки, который я использую, чтобы разрешить удаленный доступ к памяти:
//On dev_0
cudaEventCreateWithFlags
cudaIpcGetEventHandle
Exchange info with device_1
cudaIpcOpenMemHandle
//on dev_1:
cudaIpcGetMemHandle
Exchange info with device_0
cudaIpcOpenEventHandle
Рабочий код является следующим:
//On dev_0
call cuda kernel
call cudaRecordEvent
//On dev_1
cudaStreamWaitEvent
print the data on the stream
РЕДАКТИРОВАТЬ: я нашел в документации https://docs.nvidia.com/cuda/gpudirect-rdma/#sync -поведение
Пример ситуации для сетевого взаимодействия Сценарий - это когда операция записи в сеть RDMA завершается сторонним сетевым устройством, и данные записываются в отображение BAR GPU. Хотя при чтении записанных данных с помощью GPU BAR или операции копирования в память CUDA будут возвращены вновь записанные данные, при одновременном запуске ядра GPU в эту сетевую запись могут наблюдаться устаревшие данные, частично записанные данные или данные, записанные вне -order.
Короче говоря, ядро графического процессора полностью несовместимо с параллельной RDMA для операций GPUDirect, и доступ к памяти, перезаписанной сторонним устройством в такой ситуации, будет считаться гонкой данных. Чтобы устранить это несоответствие и удалить гонку данных, операция записи DMA должна завершиться относительно потока ЦП, который запустит зависимое ядро графического процессора.
Поэтому я предполагаю, что наблюдаю устаревшие данные. Теперь я предполагаю, что записанное событие соответствует отправке данных из источника, а не записи данных по цели. Я ожидал бы, что проверенное событие на цели соответствует записи, а не чтению. Я прав ? Если да, то как я могу положиться на драйвер cuda?
РЕДАКТИРОВАТЬ 2: По просьбе Роберта Кровеллы я копирую код своей игрушки.
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <mpi.h>
// CUDA kernels
template <typename T>
__global__
void copyDevKer( int n, const T *dA, T *dB)
{
for (int tid = threadIdx.x + blockIdx.x * blockDim.x;
tid < n;
tid += blockDim.x * gridDim.x )
dB[tid] = dA[tid];
}
__global__
static void printfDevKer(int n, const double *buf){
if ( ! threadIdx.x )
for ( int i = 0 ; i < n; ++i)
printf ( "buf[%i] = %f\n", i, buf[i] );
}
static void initArray (int n, double *array)
{
double x = 0.0;
for ( int i = 0; i < n; ++i ) {
x = 2.0 * i / n;
array[i] = exp(-(x * x));
}
}
int cudaSetupP2PComm( void *d_array,
int source,
int dest,
cudaEvent_t *memEvent,
cudaIpcEventHandle_t *eventHandle,
cudaIpcMemHandle_t *memHandle,
void **d_rArray,//remote array
MPI_Comm comm)
{
int rank = 0;
int tag = 0;
MPI_Status mpi_status;
cudaError_t cuerr;
MPI_Comm_rank ( comm, &rank );
if ( rank == source ){
cuerr = cudaEventCreateWithFlags ( memEvent,
cudaEventInterprocess | cudaEventDisableTiming );
if ( cuerr != cudaSuccess ){
fprintf ( stderr, "Error %d\n", __LINE__ );
return -1;
}
cuerr = cudaIpcGetEventHandle ( eventHandle, *memEvent );
if ( cuerr != cudaSuccess ){
fprintf ( stderr, "Error %d\n", __LINE__ );
return -1;
}
//Exchange handles
MPI_Recv ( memHandle, 64, MPI_CHAR, dest, tag, comm, &mpi_status);
MPI_Send ( eventHandle, 64, MPI_CHAR, dest, tag, comm );
//Get a pointer of the remove memory
cuerr = cudaIpcOpenMemHandle ( d_rArray, *memHandle, cudaIpcMemLazyEnablePeerAccess );
if ( cuerr != cudaSuccess ){
fprintf ( stderr, "Error %d\n", __LINE__ );
return -1;
}
}else{
cuerr = cudaIpcGetMemHandle ( memHandle, d_array );
if ( cuerr != cudaSuccess ){
fprintf ( stderr, "Cannot get IPC_handle from d_array\n" );
return -1;
}
//Exchange handles
MPI_Send ( memHandle, 64, MPI_CHAR, source, tag, comm );
MPI_Recv ( eventHandle, 64, MPI_CHAR, source, tag, comm, &mpi_status );
cuerr = cudaIpcOpenEventHandle ( memEvent, *eventHandle );
if ( cuerr != cudaSuccess ){
fprintf ( stderr, "Error %d\n", __LINE__ );
return -1;
}
}
MPI_Barrier ( comm );
return 0;
}
int cudaDestroyP2PComm( void *d_array,
int source,
int dest,
cudaEvent_t memEvent,
cudaIpcEventHandle_t eventHandle,
cudaIpcMemHandle_t memHandle,
void *d_rArray,
MPI_Comm comm)
{
int rank = 0;
int tag = 0;
MPI_Status mpi_status;
cudaError_t cuerr;
MPI_Comm_rank ( comm, &rank );
if ( rank == source ) {
//Has to be done before the cudaFree of d_rArray on the target
cudaIpcCloseMemHandle ( d_rArray );
}
cudaEventDestroy ( memEvent );
MPI_Barrier ( comm );
return 0;
}
void cuda_send( void *d_array,
int n,
int dest,
int tag,
MPI_Comm comm,
cudaEvent_t copyEvent,
void *d_rArray,
cudaStream_t stream)
{
size_t arraysize = n * sizeof(double);
cudaError_t cuerr;
int nthread = 1024;
int nblock = ( n + nthread - 1 ) / nthread;
copyDevKer
<<< nblock, nthread, 0, stream >>>
( n, (double*)d_array, (double*)d_rArray );
cudaEventRecord ( copyEvent, stream );
#ifdef FORCE_ORDER
MPI_Barrier ( comm );
#endif
}
void cuda_recv( void *d_array, //unused in this example
int n, //unused in this example
int source, //unused in this example
int tag, //unused in this example
MPI_Comm comm,
cudaEvent_t copyEvent,
cudaStream_t stream)
{
#ifdef FORCE_ORDER
MPI_Barrier ( comm );
#endif
cudaStreamWaitEvent ( stream, copyEvent, 0 );
}
void bench(int n, MPI_Comm comm)
{
int rank = 0;
int tag = 0;
int source = 0;
int dest = 1;
size_t arraysize = n * sizeof(double);
double *array = (double*) malloc ( arraysize );
void *d_array = NULL;
void *d_rArray = NULL; // Do not free it
cudaError_t cuerr;
cudaStream_t stream;
cudaIpcMemHandle_t ipcMemHandle;
cudaEvent_t copyEvent;
cudaIpcEventHandle_t ipcEventHandle;
initArray ( n, array );
cuerr = cudaMalloc ( &d_array, arraysize );
if ( cuerr != cudaSuccess ) {
fprintf ( stderr, "Error, cannot allocate d_array of size %zu on DEVICE\n",
arraysize );
free ( array );
return ;
}
//Setup the handles and events vars
int err = cudaSetupP2PComm ( d_array, source, dest, ©Event, &ipcEventHandle,
&ipcMemHandle, &d_rArray, comm);
if ( err ) {
fprintf ( stderr, "Error, cannot create env for cuda p2p comm\n" );
cudaFree ( d_array );
free ( array );
return ;
}
MPI_Comm_rank ( comm, &rank );
cudaStreamCreate ( &stream );
//Actual communication
if ( rank == 0 ) {
cudaMemcpy ( d_array, array, arraysize, cudaMemcpyHostToDevice );
cuda_send ( d_array, n, dest, tag, comm, copyEvent, d_rArray, stream );
} else {
cuda_recv ( d_array, n, source, tag, comm, copyEvent, stream );
//Print the first 10 elements
printfDevKer <<< 1, 1, 0, stream >>> ( 10, (double*)d_array );
}
cudaStreamSynchronize ( stream );
// Clean-up
cudaDestroyP2PComm ( d_array, source, dest, copyEvent, ipcEventHandle,
ipcMemHandle, d_rArray, comm);
cudaStreamDestroy ( stream );
cudaFree ( d_array );
free(array);
}
#define USAGE "Usage %s :\n--n <int>\tsize of the array\n"
int main(int argc, char* argv[])
{
MPI_Comm comm;
int n = 10;
int rank = 0;
int size = 1;
int devId = 0;
int ndevice = 0;
for (int i = 1; i < argc; ++i){
if (!strcmp(argv[i], "-h")){
printf(USAGE, argv[0]);
return 0;
}
if (!strcmp(argv[i], "--n")){
n = atoi(argv[i + 1]);
continue;
}
}
MPI_Init ( &argc, &argv );
MPI_Comm_dup(MPI_COMM_WORLD, &comm);
MPI_Comm_rank ( comm, &rank );
MPI_Comm_size ( comm, &size );
cudaGetDeviceCount ( &ndevice );
cudaSetDevice ( ( ndevice == size ) ? rank : (rank) % ndevice );
cudaGetDevice ( &devId );
printf ( "MPI rank %d/%d linked to DEVICE %d/%d\n",
rank, size, devId, ndevice );
bench ( n, comm );
MPI_Finalize();
return 0;
}
После компиляции с CUDA и MPI, я выполняю код дважды и получаю следующий вывод
$mpirun -n 2 ./a.out --n 10000
MPI rank 0/2 linked to DEVICE 0/8
MPI rank 1/2 linked to DEVICE 1/8
buf[0] = 1.000000
buf[1] = 1.000000
buf[2] = 1.000000
buf[3] = 1.000000
buf[4] = 0.999999
buf[5] = 0.999999
buf[6] = 0.999999
buf[7] = 0.999998
buf[8] = 0.999997
buf[9] = 0.999997
$mpirun -n 2 ./a.out --n 10000
MPI rank 0/2 linked to DEVICE 0/8
MPI rank 1/2 linked to DEVICE 1/8
buf[0] = 0.000000
buf[1] = 0.000000
buf[2] = 0.000000
buf[3] = 0.000000
buf[4] = 0.000000
buf[5] = 0.000000
buf[6] = 0.000000
buf[7] = 0.000000
buf[8] = 0.000000
buf[9] = 0.000000
Теперь, если я принудительно вызову cudaEventRecord перед cudaStreamWaitEvent путем компиляции с -DFORCE_ORDER, я получу правильный контент / вывод.
Однако я не должен использовать этот дорогостоящий барьер здесь.
Я забыл упомянуть, что я использую машину DGX-1 с установленной на ней CUDA 10.1 / 10.2.