Latenza di accesso alla memoria della CPU dei dati allocati con malloc() rispetto a cudaHostAlloc() su Tegra TK1

Question

Sto eseguendo un semplice test che confronta la latenza di accesso dei dati allocati con malloc() e dati allocati con cudaHostAlloc() da l'host (la cpu sta eseguendo gli accessi). Ho notato che l'accesso ai dati assegnati con cudaHostAlloc() è molto più lento dell'accesso ai dati allocati con malloc() sul Jetson Tk1.

Questo non è il caso per le GPU discrete e sembra essere applicabile solo a TK1. Dopo alcune indagini, ho scoperto che i dati allocati con cudaHostAlloc() sono mappati in memoria (mmap) nelle aree/dev/nvmap dello spazio di indirizzamento del processo. Questo non è il caso per i normali dati malloc'd che sono mappati sull'heap del processo. Comprendo che questa mappatura potrebbe essere necessaria per consentire alla GPU di accedere ai dati poiché i dati di cudaHostAlloc devono essere visibili sia dall'host sia dal dispositivo.

La mia domanda è la seguente: Da dove proviene il sovraccarico di accesso ai dati di cudaHostAlloc'd dall'host? I dati sono mappati su/dev/nvmap senza cache sulle cache della CPU?

Answer 1

Credo di aver trovato il motivo di questo comportamento. Dopo ulteriori indagini (utilizzando e guardando lo nvmap driver code) ho scoperto che la fonte del sovraccarico deriva dal fatto che i dati allocati con cudaHostAlloc() sono contrassegnati come "non accessibili" utilizzando il flag NVMAP_HANDLE_UNCACHEABLE. Una chiamata a pgprot_noncached() viene effettuata per assicurare che le PTE pertinenti siano contrassegnate come non memorizzabili.

Il comportamento degli accessi dell'host ai dati assegnati utilizzando cudaMallocManaged() è diverso. I dati verranno memorizzati nella cache (utilizzando la flag NVMAP_HANDLE_CACHEABLE). Pertanto l'accesso a questi dati dall'host è equivalente ai dati malloc()'d. È inoltre importante notare che il runtime CUDA non consente l'accesso di dispositivi (GPU) a qualsiasi dato allocato con cudaMallocManaged() contemporaneamente all'host e tale azione genererebbe un segfault. Il runtime, tuttavia, consente accessi simultanei ai dati cudaHostAlloc()'d sia sul dispositivo che sull'host e credo che questo sia uno dei motivi per cui i dati cudaHostAlloc()'d non sono accessibili.