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（150 Gb/s，理論最大帶寬約爲224Gb/s）。

NVIDIA發佈了一個測量傳輸帶寬的示例，您可以閱讀有關here的信息。您可以使用它來爲自己探索硬件的性能。