CUDA 8.0 - cudaMemcpy() - lineare oder konstante Zeitoperation?

Question

Kopiert CUDA 8.0 cudaMemcpy() gleichzeitig einen ganzen Speicherblock oder byteweise?

Ich möchte die Kopierzeit begrenzen, aber ich kann nichts in der Dokumentation finden, die angibt, ob cudaMemcpy() eine lineare oder konstante Zeitoperation ist.

Answer 1

Synchrone Speicherübertragungen sind keine konstante Zeit, sondern haben sowohl eine feste Latenzkomponente als auch eine Komponente, die proportional zur Übertragungsgröße ist. Bei kleinen Größen dominiert die Latenz, bei großen Größen ist die begrenzende Übertragungsgeschwindigkeit durch die Speicher- oder Busbandbreite begrenzt.

Betrachten Sie das folgende triviale Benchmark:

#include <iostream> 
#include <string> 
#include <algorithm> 

__global__ void memsetkernel(int *x, int n) 
{ 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    int stride = blockDim.x * gridDim.x; 
    for(; tid < n; tid += stride) { 
     x[tid] = threadIdx.x; 
    } 
} 

int main(int argc, char* argv[]) 
{ 
    // size 
    int n = 100; 
    int nreps = 10; 

    if (argc > 1) { 
     n = std::stoi(std::string(argv[1])); 
    } 

    size_t sz = sizeof(int) * size_t(n); 

    // host array 
    int* host = new int[n]; 

    // allocate size ints on device 
    int* device; 
    cudaMalloc(&device, sz); 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    { 
     int nthreads = 1024; 
     int nblocks = std::max(1, std::min(13*2, n/nthreads)); 
     memsetkernel<<<nblocks, nthreads>>>(device, n); 
     cudaDeviceSynchronize(); 
     cudaEventRecord(start); 
     for(int i=0; i<nreps; i++) { 
      memsetkernel<<<nblocks, nthreads>>>(device, n); 
     } 
     cudaEventRecord(stop); 
     cudaEventSynchronize(stop); 
     float milliseconds, kilobytes, bandwidth; 
     cudaEventElapsedTime(&milliseconds, start, stop); 
     milliseconds /= float(nreps); // Average of nreps 
     kilobytes = float(sz)/1e3f; 
     bandwidth = kilobytes/milliseconds;   
     std::cout << "kernel assignment: " << bandwidth << " Mb/s" << std::endl; 
    } 

    { 
     cudaMemcpy(host, device, sz, cudaMemcpyDeviceToHost); 
     cudaEventRecord(start); 
     for(int i=0; i<nreps; i++) { 
      cudaMemcpy(host, device, sz, cudaMemcpyDeviceToHost); 
     } 
     cudaEventRecord(stop); 
     cudaEventSynchronize(stop); 
     float milliseconds, kilobytes, bandwidth; 
     cudaEventElapsedTime(&milliseconds, start, stop); 
     milliseconds /= float(nreps); // Average of nreps 
     kilobytes = float(sz)/1e3f; 
     bandwidth = kilobytes/milliseconds;   
     std::cout << "DTOH: " << bandwidth << " Mb/s" << std::endl; 
    } 

    { 
     cudaMemcpy(device, host, sz, cudaMemcpyHostToDevice); 
     cudaEventRecord(start); 
     for(int i=0; i<nreps; i++) { 
      cudaMemcpy(device, host, sz, cudaMemcpyHostToDevice); 
     } 
     cudaEventRecord(stop); 
     cudaEventSynchronize(stop); 
     float milliseconds, kilobytes, bandwidth; 
     cudaEventElapsedTime(&milliseconds, start, stop); 
     milliseconds /= float(nreps); // Average of nreps 
     kilobytes = float(sz)/1e3f; 
     bandwidth = kilobytes/milliseconds; 
     std::cout << "HTOD: " << bandwidth << " Mb/s" << std::endl; 
    } 

    // reset device 
    cudaDeviceReset(); 

} 

des Lauf bei unterschiedlichen Datengrößen zeigt das folgende Verhalten:

Beiden Geräte-zu-Host und Host-to-device asymptotisch einen Wert von etwa 60% der Bandbreite des PCI-e-Busses der fraglichen Maschine erreichen (etwa 6,5 ​​Gb/s, höhere kann unter Verwendung eines gepinnten Host-Speichers erreicht werden), während der Kern etwa 70% der Hauptspeicherbandbreite erreicht die GPU (150 Gb/s mit einer theoretischen maximalen Bandbreite von etwa 224 Gb/s).

NVIDIA liefert ein Beispiel für die Messung der Übertragungsbandbreite, die Sie über here lesen können. Sie können dies verwenden, um die Leistung Ihrer Hardware für sich selbst zu erkunden.