Wpływ wykorzystania pamięci stronicowanej do kopiowania pamięci asynchronicznej?

Question

W CUDA C Best Practices Guide Version 5.0, pkt 6.1.2, jest napisane, że:

W przeciwieństwie cudaMemcpy(), asynchronicznego transferu wersji wymagającej przypięte pamięć hosta (patrz podpięty Memory) oraz zawiera dodatkowy argument o nazwie ID strumienia.

Oznacza to, że funkcja cudaMemcpyAsync powinna zakończyć się niepowodzeniem, jeśli użyję zwykłej pamięci.

Ale tak się nie stało.

Tylko dla celów testowania, próbowałem następujący program:

jądra:

__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; 
} 

Main:

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; 
} 

Program dał poprawny wynik. Tablica inkrementowała się pomyślnie.

Jak wykonać cudaMemcpyAsync bez pamięci zablokowanej strony? Czy tu czegoś brakuje?

Answer 1

cudaMemcpyAsync jest zasadniczo asynchroniczną wersją cudaMemcpy. Oznacza to, że nie blokuje on wywołującego wątku hosta po wywołaniu kopiowania. To jest podstawowe zachowanie połączenia.

Opcjonalnie , jeśli połączenie jest uruchamiane w strumieniu innym niż domyślny, a pamięć hosta jest przypisem przypiętym, a urządzenie ma wolny mechanizm kopiowania DMA, operacja kopiowania może się zdarzyć, podczas gdy GPU wykonuje jednocześnie inne operacja: wykonanie jądra lub inna kopia (w przypadku GPU z dwoma silnikami kopiującymi DMA). Jeśli wszystkie te warunki nie są spełnione, operacja na GPU jest funkcjonalnie identyczna ze standardowym wywołaniem cudaMemcpy, tj. serializuje operacje na GPU i nie może wystąpić jednoczesne wykonanie kopii jądra lub jednoczesne wielokrotne kopiowanie. Jedyna różnica polega na tym, że operacja nie blokuje wywołującego wątku hosta.

W tym przykładowym kodzie źródło hosta i pamięć docelowa nie są przypięte. Tak więc transfer pamięci nie może nakładać się na wykonanie jądra (tzn. Serializuje operacje na GPU). dzwoni są nadal asynchroniczne na hoście. Więc co masz jest funkcjonalnym odpowiednikiem:

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

z tą różnicą, że wszystkie połączenia są asynchroniczne na komputerze, więc bloki nić gospodarza na wezwanie cudaDeviceSynchronize() zamiast na każdym z połączeń transferowych pamięci.

Jest to zachowanie absolutnie oczekiwane.