CUDA 8.0 - cudaMemcpy（） - 線形または一定時間の操作ですか？

Question

CUDA 8.0 cudaMemcpy()は、同時に全メモリブロックにコピーするか、バイト単位でコピーしますか？

私は、コピー時間を制限したいがcudaMemcpy()は線形または一定時間操作である場合、私は指定のドキュメントに何かを見つけることができません。

Answer 1

同期メモリ転送は時定数ではなく、むしろ固定レイテンシ成分と転送サイズに比例する成分の両方を有します。小さなサイズでは、レイテンシが大規模であり、転送速度の制限はメモリまたはバスの帯域幅によって制限されます。

は、以下の些細なベンチマークを検討：異なるデータサイズでこれを実行

#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(); 

} 

は、以下の挙動を示す：デバイスホストおよびホスト - デバイス漸近両方

問題のマシンのPCI-eバスの帯域幅の約60％に近づきます（固定ホストメモリを使用して約6.5Gb/s以上に達することができます）。一方、カーネルはメインメモリ帯域幅の約70％に達しますGPU（理論上の最大帯域幅約224Gb/sの150Gb/s）。

NVIDIAの船あなたがについてhereを読むことができる転送帯域幅を測定するためのサンプル。これを使用して、ハードウェアのパフォーマンスを自分で調べることができます。