So kopieren Sie mithilfe von Thrust- und CUDA-Streams Speicher asynchron vom Host auf das Gerät

Lesezeit: 7 Minuten

Ich möchte Speicher vom Host auf das Gerät kopieren, indem ich Schub wie in verwende

thrust::host_vector<float> h_vec(1 << 28);
thrust::device_vector<float> d_vec(1 << 28);
thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin());

Verwenden Sie CUDA-Streams analog dazu, wie Sie Speicher mithilfe von Streams vom Gerät auf das Gerät kopieren würden:

cudaStream_t s;
cudaStreamCreate(&s);

thrust::device_vector<float> d_vec1(1 << 28), d_vec2(1 << 28);
thrust::copy(thrust::cuda::par.on(s), d_vec1.begin(), d_vec1.end(), d_vec2.begin());

cudaStreamSynchronize(s);
cudaStreamDestroy(s);

Das Problem ist, dass ich die Ausführungsrichtlinie nicht auf CUDA setzen kann, um den Stream beim Kopieren vom Host auf das Gerät anzugeben, da Thrust in diesem Fall davon ausgehen würde, dass beide Vektoren auf dem Gerät gespeichert sind. Gibt es eine Möglichkeit, dieses Problem zu umgehen? Ich verwende die neueste Schubversion von github (in der Datei version.h steht 1.8).

  • Das Ankündigung, die ich gelesen habe ließ es klingen, als ob Streams für Underlying implementiert wurden Kernel-Aufrufe, nicht unbedingt pauschal in Schub. Wenn Sie Streams zum Kopieren vom Hostvektor zum Gerätevektor verwendet haben, möchten Sie wahrscheinlich einen verwenden gepinnte Zuweisung auf dem Host. Ich glaube daher, dass das, wonach Sie fragen, mit Schubvektoren und erreicht werden könnte cudaMemcpyAsync.

    – Robert Crovella

    31. Juli 2014 um 17:02 Uhr

  • Ja, sollten Sie verwenden cudaMemcpyAsync dafür direkt, wie Robert vorschlägt.

    – Jared Hoberock

    31. Juli 2014 um 17:26 Uhr

  • Ab heute (Mai 2016) finde ich den ersten Eintrag in der Dokumentation hier: schub.github.io/doc/… wirklich verstörend. Es besagt, dass wir throw::copy( throw::cuda::par.on(cudaStream), HostPtr, HostPtr+size, DevicePtr ); als gültige Syntax, ohne dass Ihre Kopie asynchron zu dem bestimmten Stream ausgegeben wird, den Sie im Parameter übergeben …

    – Tobey

    20. Mai 2016 um 16:32 Uhr


Benutzer-Avatar
Robert Crovella

Wie in den Kommentaren angegeben, glaube ich nicht, dass dies direkt mit möglich sein wird thrust::copy. Allerdings können wir verwenden cudaMemcpyAsync in einer Schubanwendung, um das Ziel von asynchronen Kopien und Überlappung von Kopie mit Berechnung zu erreichen.

Hier ist ein ausgearbeitetes Beispiel:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/system/cuda/experimental/pinned_allocator.h>
#include <thrust/system/cuda/execution_policy.h>
#include <thrust/fill.h>
#include <thrust/sequence.h>
#include <thrust/for_each.h>
#include <iostream>

// DSIZE determines duration of H2D and D2H transfers
#define DSIZE (1048576*8)
// SSIZE,LSIZE determine duration of kernel launched by thrust
#define SSIZE (1024*512)
#define LSIZE 1
// KSIZE determines size of thrust kernels (number of threads per block)
#define KSIZE 64
#define TV1 1
#define TV2 2

typedef int mytype;
typedef thrust::host_vector<mytype, thrust::cuda::experimental::pinned_allocator<mytype> > pinnedVector;

struct sum_functor
{
  mytype *dptr;
  sum_functor(mytype* _dptr) : dptr(_dptr) {};
  __host__ __device__ void operator()(mytype &data) const
    {
      mytype result = data;
      for (int j = 0; j < LSIZE; j++)
        for (int i = 0; i < SSIZE; i++)
          result += dptr[i];
      data = result;
    }
};

int main(){

  pinnedVector hi1(DSIZE);
  pinnedVector hi2(DSIZE);
  pinnedVector ho1(DSIZE);
  pinnedVector ho2(DSIZE);
  thrust::device_vector<mytype> di1(DSIZE);
  thrust::device_vector<mytype> di2(DSIZE);
  thrust::device_vector<mytype> do1(DSIZE);
  thrust::device_vector<mytype> do2(DSIZE);
  thrust::device_vector<mytype> dc1(KSIZE);
  thrust::device_vector<mytype> dc2(KSIZE);

  thrust::fill(hi1.begin(), hi1.end(),  TV1);
  thrust::fill(hi2.begin(), hi2.end(),  TV2);
  thrust::sequence(do1.begin(), do1.end());
  thrust::sequence(do2.begin(), do2.end());

  cudaStream_t s1, s2;
  cudaStreamCreate(&s1); cudaStreamCreate(&s2);

  cudaMemcpyAsync(thrust::raw_pointer_cast(di1.data()), thrust::raw_pointer_cast(hi1.data()), di1.size()*sizeof(mytype), cudaMemcpyHostToDevice, s1);
  cudaMemcpyAsync(thrust::raw_pointer_cast(di2.data()), thrust::raw_pointer_cast(hi2.data()), di2.size()*sizeof(mytype), cudaMemcpyHostToDevice, s2);

  thrust::for_each(thrust::cuda::par.on(s1), do1.begin(), do1.begin()+KSIZE, sum_functor(thrust::raw_pointer_cast(di1.data())));
  thrust::for_each(thrust::cuda::par.on(s2), do2.begin(), do2.begin()+KSIZE, sum_functor(thrust::raw_pointer_cast(di2.data())));

  cudaMemcpyAsync(thrust::raw_pointer_cast(ho1.data()), thrust::raw_pointer_cast(do1.data()), do1.size()*sizeof(mytype), cudaMemcpyDeviceToHost, s1);
  cudaMemcpyAsync(thrust::raw_pointer_cast(ho2.data()), thrust::raw_pointer_cast(do2.data()), do2.size()*sizeof(mytype), cudaMemcpyDeviceToHost, s2);

  cudaDeviceSynchronize();
  for (int i=0; i < KSIZE; i++){
    if (ho1[i] != ((LSIZE*SSIZE*TV1) + i)) { std::cout << "mismatch on stream 1 at " << i << " was: " << ho1[i] << " should be: " << ((DSIZE*TV1)+i) << std::endl; return 1;}
    if (ho2[i] != ((LSIZE*SSIZE*TV2) + i)) { std::cout << "mismatch on stream 2 at " << i << " was: " << ho2[i] << " should be: " << ((DSIZE*TV2)+i) << std::endl; return 1;}
    }
  std::cout << "Success!" << std::endl;
  return 0;
}

Für meinen Testfall habe ich RHEL5.5, Quadro5000 und cuda 6.5RC verwendet. Dieses Beispiel ist so konzipiert, dass Thrust sehr kleine Kernel erstellt (nur ein einzelner Threadblock, solange KSIZE klein ist, sagen wir 32 oder 64), so dass die Kerne, die der Schub erzeugt, aus thrust::for_each parallel laufen können.

Wenn ich diesen Code profiliere, sehe ich:

nvvp-Ausgang für die Anwendung von Schubströmen

Dies weist darauf hin, dass wir sowohl zwischen Schubkernen als auch zwischen Kopieroperationen und Schubkernen sowie einem asynchronen Datenkopieren bei der Vollendung der Kernel eine ordnungsgemäße Überlappung erreichen. Notiere dass der cudaDeviceSynchronize() Die Operation “füllt” die Zeitachse und zeigt an, dass alle asynchronen Operationen (Datenkopieren, Schubfunktionen) asynchron ausgegeben wurden und die Steuerung an den Host-Thread zurückgegeben wurde, bevor eine der Operationen im Gange war. All dies wird erwartet, richtiges Verhalten für vollständige Parallelität zwischen Host-, GPU- und Datenkopiervorgängen.

  • ist die Verwendung von pf pinned_allocator notwendig zu machen cudaMemcpyAsync richtig arbeiten auf a thrust::host_vector? Was würde passieren, wenn ich einen Standard verwenden würde thrust::host_vector?

    – Frau

    13. Juni 2015 um 11:46 Uhr

  • Standard host_vector verwendet einen nicht fixierten (dh nicht seitengesperrten) Zuordner. Das heißt, wenn Sie versuchen, a cudaMemcpyAsync, ist der Vorgang nicht asynchron. Und wenn Sie versuchen, diese Operation mit etwas anderem zu überlappen, wird sie sich nicht überschneiden. Schau mal hier und hier

    – Robert Crovella

    13. Juni 2015 um 13:55 Uhr


  • Beachten Sie, dass die Version von Thrust, die mit CUDA 7 geliefert wurde, eine Ausgabe Dies verhindert in einigen Fällen die ordnungsgemäße Ausgabe von Schubkernen an Streams. Die Problemumgehung wäre, 1. den Schub auf CUDA 7 auf die zu aktualisieren aktuelle Entwicklungsversion (einschließlich der Behebung des Problems) oder 2. zu CUDA 6.5 zurückkehren (oder zu einer zukünftigen CUDA-Toolkit-Version wechseln, sobald diese verfügbar ist.)

    – Robert Crovella

    13. Juni 2015 um 14:48 Uhr


  • Wurde dies in 7.5 behoben oder bleibt es bestehen?

    – Bar

    17. Dezember 2015 um 13:07 Uhr

  • Ja, die Schubversion, die mit CUDA 7.5 geliefert wird, hat die Lösung für dieses Problem übernommen.

    – Robert Crovella

    17. Dezember 2015 um 13:14 Uhr

Hier ist ein funktionierendes Beispiel mit thrust::cuda::experimental::pinned_allocator<T>:

// Compile with:
// nvcc --std=c++11 mem_async.cu -o mem_async

#include <cuda.h>
#include <cuda_runtime.h>
#include <cufft.h>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/fill.h>
#include <thrust/system/cuda/experimental/pinned_allocator.h>

#define LEN 1024

int main(int argc, char *argv[]) {
    thrust::host_vector<float, thrust::cuda::experimental::pinned_allocator<float>> h_vec(LEN);
    thrust::device_vector<float> d_vec(LEN);

    thrust::fill(d_vec.begin(), d_vec.end(), -1.0);

    cudaMemcpyAsync(thrust::raw_pointer_cast(h_vec.data()),
                    thrust::raw_pointer_cast(d_vec.data()),
                    d_vec.size()*sizeof(float),
                    cudaMemcpyDeviceToHost);

    // Comment out this line to see what happens.
    cudaDeviceSynchronize();

    std::cout << h_vec[0] << std::endl;
}

Kommentieren Sie den Synchronisierungsschritt aus und Sie sollten erhalten 0 aufgrund der asynchronen Speicherübertragung an die Konsole gedruckt.

1012190cookie-checkSo kopieren Sie mithilfe von Thrust- und CUDA-Streams Speicher asynchron vom Host auf das Gerät

This website is using cookies to improve the user-friendliness. You agree by using the website further.

Privacy policy