Возможен ли одновременный cudaMemcpyAsync? - PullRequest
Новая
       10

Возможен ли одновременный cudaMemcpyAsync?

2 голосов
/ 02 апреля 2019

Я пишу тестовый код, чтобы ознакомиться с параллельными атрибутами cudaMemcpyAsync.

Когда я пытался выполнить одновременное cudaMemcpyAsync в единственном контексте , операции копирования ставятся в очередь и выполняются одна за другой с пропускной способностью 12.4 GB/s, что соответствует ответу здесь :

Но когда я попытался сделать одновременные cudaMemcpyAsync в различных контекстах (разделив их на 4 процесса), кажется, что первый и последний работают одновременно:

Первые 2 последовательных cudaMemcpyAsync работают с пропускной способностью 12.4 GB/s, в то время как последние 2 одновременных работают с пропускной способностью 5.3 GB/s.

Как я могу сделать одновременные cudaMemcpyAsync в одном контексте?

Я использую CUDA9.0 на TITAN Xp, который имеет 2 механизма копирования.

EDIT

Код для сценария 1: 

#include <stdio.h>
#include <pthread.h>
#include <stdlib.h>
#include <assert.h>
#include <time.h>

inline
cudaError_t checkCuda(cudaError_t result)
{
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
  return result;
}

const int nStreams = 8;
const int N = 100000000;
const int bytes = N * sizeof(int);
int* arr_H;
int* arr_D[nStreams];
cudaStream_t stream[nStreams];
int args[nStreams];
pthread_t threads[nStreams];

void* worker(void *arg)
{
  int i = *((int *)arg);
  checkCuda(cudaMemcpyAsync(arr_D[i], arr_H, bytes, cudaMemcpyHostToDevice, stream[i]));

  return NULL;
}

int main()
{
  for(int i = 0; i < nStreams; i++)
    checkCuda(cudaStreamCreate(&stream[i]));

  checkCuda(cudaMallocHost((void**)&arr_H, bytes));
  for (int i = 0; i < N; i++)
    arr_H[i] = random();

  for (int i = 0; i < nStreams; i++)
    checkCuda(cudaMalloc((void**)&arr_D[i], bytes));

  for (int i = 0; i < nStreams; i++) {
    args[i] = i;
    pthread_create(&threads[i], NULL, worker, &args[i]);
  }

  for (int i = 0; i < nStreams; i++)
    pthread_join(threads[i], NULL);

  cudaFreeHost(arr_H);
  for (int i = 0; i < nStreams; i++) {
    checkCuda(cudaStreamDestroy(stream[i]));
    cudaFree(arr_D[i]);
  }

  return 0;

Код для сценария 2: 

#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <time.h>

inline
cudaError_t checkCuda(cudaError_t result)
{
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
  return result;
}

int main()
{
  const int nStreams = 1;
  const int N = 100000000;
  const int bytes = N * sizeof(int);
  int* arr_H;
  int* arr_D[nStreams];
  cudaStream_t stream[nStreams];

  for(int i = 0; i < nStreams; i++)
    checkCuda(cudaStreamCreate(&stream[i]));

  checkCuda(cudaMallocHost((void**)&arr_H, bytes));
  for (int i = 0; i < N; i++)
    arr_H[i] = random();

  for (int i = 0; i < nStreams; i++)
    checkCuda(cudaMalloc((void**)&arr_D[i], bytes));

  for (int i = 0; i < nStreams; i++)
    checkCuda(cudaMemcpyAsync(arr_D[i], arr_H, bytes, cudaMemcpyHostToDevice, stream[i]));

  cudaFreeHost(arr_H);
  for (int i = 0; i < nStreams; i++) {
    checkCuda(cudaStreamDestroy(stream[i]));
    cudaFree(arr_D[i]);
  }

  return 0;
}

Код 2 в основном скопирован из кода 1. Я использовал скрипт Python для одновременного запуска нескольких процессов: 

#!/usr/bin/env python3
import subprocess

N = 4

processes = [subprocess.Popen('./a.out', shell=True) for _ in range(N)]

for process in processes:
    process.wait()
...