Ảnh hưởng của việc sử dụng bộ nhớ trang có thể cho bản sao bộ nhớ không đồng bộ?

Question

Trong CUDA C Best Practices Guide Phiên bản 5.0, Mục 6.1.2, nó được viết rằng:

Ngược lại với cudaMemcpy(), phiên bản asynchronous transfer đòi hỏi bộ nhớ lưu trữ gắn (xem Memory Pinned), và nó chứa một đối số bổ sung , một ID luồng.

Có nghĩa là chức năng cudaMemcpyAsync sẽ bị lỗi nếu tôi sử dụng bộ nhớ đơn giản.

Nhưng đây không phải là những gì đã xảy ra.

Chỉ cần cho mục đích thử nghiệm, tôi đã thử chương trình sau đây:

Kernel:

__global__ void kernel_increment(float* src, float* dst, int n) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    if(tid<n) 
     dst[tid] = src[tid] + 1.0f; 
} 

chính:

int main() 
{ 
    float *hPtr1, *hPtr2, *dPtr1, *dPtr2; 

    const int n = 1000; 

    size_t bytes = n * sizeof(float); 

    cudaStream_t str1, str2; 

    hPtr1 = new float[n]; 
    hPtr2 = new float[n]; 

    for(int i=0; i<n; i++) 
     hPtr1[i] = static_cast<float>(i); 

    cudaMalloc<float>(&dPtr1,bytes); 
    cudaMalloc<float>(&dPtr2,bytes); 

    dim3 block(16); 
    dim3 grid((n + block.x - 1)/block.x); 

    cudaStreamCreate(&str1); 
    cudaStreamCreate(&str2); 

    cudaMemcpyAsync(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice,str1); 
    kernel_increment<<<grid,block,0,str2>>>(dPtr1,dPtr2,n); 
    cudaMemcpyAsync(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost,str1); 

    printf("Status: %s\n",cudaGetErrorString(cudaGetLastError())); 

    cudaDeviceSynchronize(); 

    printf("Status: %s\n",cudaGetErrorString(cudaGetLastError())); 

    cudaStreamDestroy(str1); 
    cudaStreamDestroy(str2); 

    cudaFree(dPtr1); 
    cudaFree(dPtr2); 

    for(int i=0; i<n; i++) 
     std::cout<<hPtr2[i]<<std::endl; 

    delete[] hPtr1; 
    delete[] hPtr2; 

    return 0; 
} 

Chương trình đưa ra chính xác. Mảng tăng lên thành công.

Cách thực hiện cudaMemcpyAsync mà không có bộ nhớ bị khóa trang? Tôi có thiếu gì đó ở đây không?

Answer 1

cudaMemcpyAsync về cơ bản là phiên bản không đồng bộ của cudaMemcpy. Điều này có nghĩa là nó không chặn chuỗi máy chủ gọi khi cuộc gọi sao chép được phát hành. Đó là hành vi cơ bản của cuộc gọi.

Tùy chọn, nếu cuộc gọi được khởi chạy vào luồng không mặc định và nếu bộ nhớ máy chủ được cấp phát và thiết bị có công cụ sao chép DMA miễn phí, hoạt động sao chép có thể xảy ra trong khi GPU thực hiện đồng thời hoạt động: thực thi hạt nhân hoặc bản sao khác (trong trường hợp GPU có hai công cụ sao chép DMA). Nếu tất cả các điều kiện này không thỏa mãn, thao tác trên GPU có chức năng giống với cuộc gọi chuẩn cudaMemcpy, nghĩa là. nó nối tiếp các hoạt động trên GPU, và không có đồng thời sao chép hạt nhân thực hiện hoặc đồng thời nhiều bản sao có thể xảy ra. Sự khác biệt duy nhất là thao tác không chặn luồng máy chủ lưu trữ cuộc gọi.

Trong mã ví dụ của bạn, nguồn máy chủ và bộ nhớ đích không được ghim. Vì vậy, việc chuyển giao bộ nhớ không thể chồng lấp với việc thực thi hạt nhân (nghĩa là chúng hoạt động theo chuỗi trên GPU). Các cuộc gọi vẫn không đồng bộ trên máy chủ lưu trữ. Vì vậy, những gì bạn có là chức năng tương đương với:

cudaMemcpy(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice); 
kernel_increment<<<grid,block>>>(dPtr1,dPtr2,n); 
cudaMemcpy(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost); 

với ngoại lệ mà tất cả các cuộc gọi không đồng bộ trên máy chủ, do đó khối thread host tại cudaDeviceSynchronize() gọi chứ không phải tại mỗi cuộc gọi chuyển bộ nhớ.

Đây là hành vi tuyệt đối được mong đợi.