Effetto dell'utilizzo di una memoria in grado di gestire la copia asincrona della memoria?

Question

In Guida CUDA C Best Practices versione 5.0, sezione 6.1.2, è scritto che:

In contrasto con cudaMemcpy(), la versione di trasferimento asincrono richiede memoria host appuntato (consultare Memoria Appuntato), e contiene un argomento aggiuntivo , un ID di flusso.

Significa che la funzione cudaMemcpyAsync non dovrebbe funzionare se utilizzo una memoria semplice.

Ma questo non è quello che è successo.

solo a scopo di test, ho provato il seguente programma:

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

principale:

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

Il programma ha dato uscita corretta. L'array è stato incrementato correttamente.

In che modo cudaMemcpyAsync è stato eseguito senza memoria di blocco pagina? Mi manca qualcosa qui?

Answer 1

cudaMemcpyAsync è fondamentalmente una versione asincrona di cudaMemcpy. Ciò significa che non blocca il thread dell'host chiamante quando viene emessa la chiamata di copia. Questo è il comportamento di base della chiamata.

Facoltativamente, se la chiamata è lanciato nel flusso non predefinito, e se la memoria host è un'allocazione bloccata, e il dispositivo ha un motore copia DMA libera, l'operazione di copia può avvenire mentre la GPU esegue simultaneamente un'altra operazione: esecuzione del kernel o un'altra copia (nel caso di una GPU con due motori di copia DMA). Se tutte queste condizioni non sono soddisfatte, l'operazione sulla GPU è funzionalmente identica a una chiamata standard cudaMemcpy, ad es. serializza le operazioni sulla GPU e non è possibile eseguire simultaneamente l'esecuzione del copia-kernel o più copie simultanee. L'unica differenza è che l'operazione non blocca il thread dell'host chiamante.

Nel codice di esempio, la memoria di origine e di destinazione dell'host non sono bloccate. Quindi il trasferimento di memoria non può sovrapporsi all'esecuzione del kernel (cioè serializzano le operazioni sulla GPU). Le chiamate sono ancora asincrone sull'host. Quindi, quello che hai è funzionalmente equivalente a:

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

con l'eccezione che tutte le chiamate sono asincrone sull'host, in modo che i blocchi di filettatura ospite al call cudaDeviceSynchronize() piuttosto che a ciascuna delle chiamate di trasferimento di memoria.

Questo comportamento è assolutamente previsto.