Странное поведение cudaMemcpyAsync: 1. cudaMemcpyKind не имеет значения. 2. Копировать не удается, но молча

Question

Я знакомлюсь с новым кластером, оснащенным графическими процессорами Pascal P100 +Nvlink. Я написал программу для пинг-понга для проверки пропускной способности gpu<->gpu и gpu<->cpu и однорангового доступа. (Я знаю, что примеры cuda содержат такую ​​программу, но я хотел бы сделать это сам для лучшего понимания.) Пропускная способность Nvlink кажется разумной (~35 ГБ / с, двунаправленная, с теоретическим максимумом 40). Однако при отладке пинг-понга я обнаружил странное поведение.

Прежде всего, cudaMemcpyAsync успешно выполняется независимо от того, какой параметр cudaMemcpyKind я указываю, например, если cudaMemcpyAsync копирует память с хоста на устройство, он будет успешным, даже если я передам cudaMemcpyDeviceToHost в качестве вида.

Во-вторых, когда память хоста не заблокирована, cudaMemcpyAsync выполняет следующие действия:

• Копирование памяти с хоста на устройство кажется успешным (нет ошибок segfaults или cuda во время выполнения, и данные отображаются для правильной передачи).
• Копирование памяти с устройства на хост происходит автоматически: segfault не происходит, и cudaDeviceSynchronize после того, как memcpy возвращает cudaSuccess, но проверка данных показывает, что данные на gpu не были правильно переданы хосту.

Можно ли ожидать такого поведения? Я включил минимальный рабочий пример кода, который демонстрирует его в моей системе (образец не является приложением для пинг-понга, все, что он делает, это проверяет cudaMemcpyAsync с различными параметрами).

У P100s включен UVA, поэтому мне кажется, что cudaMemcpyAsync просто выводит расположение указателей src и dst и игнорирует аргумент cudaMemcpyKind. Однако я не уверен, почему cudaMemcpyAsync не может выдать ошибку для хост-памяти без блокировки страницы. Я был под впечатлением, что было строго нет-нет.

#include <stdio.h>
#include <cuda_runtime.h>
#include <stdlib.h>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void checkDataDevice( int* current, int* next, int expected_current_val, int n )
{
  int tid = threadIdx.x + blockIdx.x*blockDim.x;
  for( int i = tid; i < n; i += blockDim.x*gridDim.x )
  {
    if( current[i] != expected_current_val )
      printf( "Error on device:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

void checkDataHost( int* current, int* next, int expected_current_val, int n )
{
  for( int i = 0; i < n; i++ )
  {
    if( current[i] != expected_current_val )
      printf( "Error on host:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

int main( int argc, char** argv )
{
  bool pagelocked = true;
  // invoking the executable with any additional argument(s) will turn off page locked memory, i.e.,
  // Run with pagelocked memory:  ./a.out
  // Run with ordinary malloc'd memory: ./a.out jkfdlsja
  if( argc > 1 )
    pagelocked = false;

  int copybytes = 1e8; // Ok to use int instead of size_t for 1e8.

  cudaStream_t* stream = (cudaStream_t*)malloc( sizeof(cudaStream_t) );
  cudaStreamCreate( stream );

  int* srcHost;
  int* dstHost;
  int* srcDevice;
  int* dstDevice;

  cudaMalloc( (void**)&srcDevice, copybytes );
  cudaMalloc( (void**)&dstDevice, copybytes );
  if( pagelocked )
  {
    printf( "Using page locked memory\n" );
    cudaMallocHost( (void**)&srcHost, copybytes );
    cudaMallocHost( (void**)&dstHost, copybytes );
  }
  else
  {
    printf( "Using non page locked memory\n" );
    srcHost = (int*)malloc( copybytes );
    dstHost = (int*)malloc( copybytes );
  }

  for( int i = 0; i < copybytes/sizeof(int); i++ )
    srcHost[i] = 1;

  cudaMemcpyKind kinds[4];
  kinds[0] = cudaMemcpyHostToDevice;
  kinds[1] = cudaMemcpyDeviceToHost;
  kinds[2] = cudaMemcpyHostToHost;
  kinds[3] = cudaMemcpyDeviceToDevice;

  // Test cudaMemcpyAsync in both directions,
  // iterating through all "cudaMemcpyKinds" to verify
  // that they don't matter.
  int expected_current_val = 1;
  for( int kind = 0; kind<4; kind++ )
  {
    // Host to device copy 
    cudaMemcpyAsync( dstDevice
        , srcHost
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataDevice<<<56*8,256>>>( dstDevice
        , srcDevice
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;

    // Device to host copy
    cudaMemcpyAsync( dstHost
        , srcDevice
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataHost( dstHost
        , srcHost
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;
  }

  free( stream );

  cudaFree( srcDevice );
  cudaFree( dstDevice );
  if( pagelocked )
  {
    cudaFreeHost( srcHost );
    cudaFreeHost( dstHost );
  }
  else
  {
    free( srcHost );
    free( dstHost );
  }

  return 0;
}

Answer 1

Когда возникают проблемы с кодом CUDA, я настоятельно рекомендую использовать строгую (== проверяется каждый код возврата вызова) правильную проверку ошибок CUDA.

Ваша проверка ошибок имеет недостатки, и недостатки приводят к некоторой путанице.

Прежде всего, в случае блокировки страницы данный (сопоставленный) указатель доступен / действителен как на хосте, так и на устройстве. Поэтому любое возможное перечисление направления (H2D, D2H, D2D, H2H) является законным и действительным. В результате ошибки не будут возвращены, и операция копирования будет успешной.

В случае отсутствия блокировки страницы вышеприведенное неверно, поэтому, вообще говоря, указанное направление передачи лучше соответствовало предполагаемому направлению передачи, как проверено указателями. Если это не так, cudaMemcpyAsync вернет код ошибки (cudaErrorInvalidValue == 11). В вашем случае вы игнорируете этот результат ошибки. Вы можете доказать это себе, если у вас достаточно терпения (было бы лучше, если бы вы просто отметили первую ошибку, а не распечатывали каждое несоответствие в элементах 10M+), запустив код с cuda-memcheck (еще одна хорошая вещь, которую нужно делать всякий раз, когда у вас возникают проблемы с кодом CUDA) или просто выполнять тщательную тщательную проверку ошибок.

Когда cudaMemcpyAsync операция указывает на сбой, операция не завершается успешно, поэтому данные не копируются, а проверка данных указывает на несоответствия. Надеемся, что теперь это не удивительно, поскольку ожидаемая операция копирования фактически не произошла (и при этом она не завершилась "тихо").

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

Это не правильно для cudaMemcpyAsync, Ошибка, которая может быть обнаружена при вызове cudaMemcpyAsync операция будет немедленно возвращена самим вызовом и не будет возвращена в результате последующих вызовов CUDA (ясно), так как этот тип ошибки не является липким.

Мораль этой истории:

1. Сделайте правильную проверку ошибок CUDA. Неукоснительно.
2. Запустите свой код с cuda-memcheck,

Вот полностью проработанный пример с тривиальной модификацией вашего кода, чтобы сделать вывод "нормальным" в случае сбоя, демонстрирующий, что в случае сбоя указана ошибка:

$ cat t153.cu
#include <stdio.h>
#include <stdlib.h>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void checkDataDevice( int* current, int* next, int expected_current_val, int n )
{
  int tid = threadIdx.x + blockIdx.x*blockDim.x;
  for( int i = tid; i < n; i += blockDim.x*gridDim.x )
  {
    if( current[i] != expected_current_val )
      printf( "Error on device:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

void checkDataHost( int* current, int* next, int expected_current_val, int n )
{
  for( int i = 0; i < n; i++ )
  {
    if( current[i] != expected_current_val ){
      printf( "Error on host:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
      exit(0);}
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

int main( int argc, char** argv )
{
  bool pagelocked = true;
  // invoking the executable with any additional argument(s) will turn off page locked memory, i.e.,
  // Run with pagelocked memory:  ./a.out
  // Run with ordinary malloc'd memory: ./a.out jkfdlsja
  if( argc > 1 )
    pagelocked = false;

  int copybytes = 1e8; // Ok to use int instead of size_t for 1e8.

  cudaStream_t* stream = (cudaStream_t*)malloc( sizeof(cudaStream_t) );
  cudaStreamCreate( stream );

  int* srcHost;
  int* dstHost;
  int* srcDevice;
  int* dstDevice;

  cudaMalloc( (void**)&srcDevice, copybytes );
  cudaMalloc( (void**)&dstDevice, copybytes );
  if( pagelocked )
  {
    printf( "Using page locked memory\n" );
    cudaMallocHost( (void**)&srcHost, copybytes );
    cudaMallocHost( (void**)&dstHost, copybytes );
  }
  else
  {
    printf( "Using non page locked memory\n" );
    srcHost = (int*)malloc( copybytes );
    dstHost = (int*)malloc( copybytes );
  }

  for( int i = 0; i < copybytes/sizeof(int); i++ )
    srcHost[i] = 1;

  cudaMemcpyKind kinds[4];
  kinds[0] = cudaMemcpyHostToDevice;
  kinds[1] = cudaMemcpyDeviceToHost;
  kinds[2] = cudaMemcpyHostToHost;
  kinds[3] = cudaMemcpyDeviceToDevice;

  // Test cudaMemcpyAsync in both directions,
  // iterating through all "cudaMemcpyKinds" to verify
  // that they don't matter.
  int expected_current_val = 1;
  for( int kind = 0; kind<4; kind++ )
  {
    // Host to device copy
    cudaMemcpyAsync( dstDevice
        , srcHost
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataDevice<<<56*8,256>>>( dstDevice
        , srcDevice
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;

    // Device to host copy
    cudaMemcpyAsync( dstHost
        , srcDevice
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataHost( dstHost
        , srcHost
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;
  }

  free( stream );

  cudaFree( srcDevice );
  cudaFree( dstDevice );
  if( pagelocked )
  {
    cudaFreeHost( srcHost );
    cudaFreeHost( dstHost );
  }
  else
  {
    free( srcHost );
    free( dstHost );
  }

  return 0;
}
$ nvcc -arch=sm_61 -o t153 t153.cu
$ cuda-memcheck ./t153 a
========= CUDA-MEMCHECK
Using non page locked memory
========= Program hit cudaErrorInvalidValue (error 11) due to "invalid argument" on CUDA API call to cudaMemcpyAsync.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0x2ef423]
=========     Host Frame:./t153 [0x489a3]
=========     Host Frame:./t153 [0x2e11]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf5) [0x21ec5]
=========     Host Frame:./t153 [0x2a49]
=========
Error on host:  expected = 2, current[0] = 0
========= ERROR SUMMARY: 1 error
$