Increasing utilization - gdrcopy_copybw

Aperta
#288 3 commenti 0 reazioni 0 assegnatari Vedi su GitHub

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
4/5
Tempo stimato
3-5 giorni
Idoneità per principianti
25/100
Tipo di issue
Bug
Chiarezza
Da chiarire
Stato di attività
Ferma
Stack tecnologico
c, linux
Ambito
performance

Direzione di ricerca

Start with the gdrcopy_copybw benchmark and the memcpy_uncached_store_avx implementation mentioned in the issue. Compare the existing _mm256_load_pd and _mm256_stream_pd path with the _mm512 equivalents on the stated NVIDIA A100 80GB PCIe and Gen 4 PCIe setup. Done means documenting the reason for the current width and whether 512-bit operations improve utilization for 32kB-8MB buffers.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

Hi,

I am running the gdrcopy_copybw benchmark on NVIDIA A100 80GB PCIe with Gen 4 PCIe.
It appears that the utilization doesn't reach its maximal possible value, getting about 20 GB/s out of the possible 32 GB/s, for buffers of sizes 32kB-8MB.
Upon looking into the code, it appears that in your implementation of memcpy_uncached_store_avx you are using 256 bit functions _mm256_load_pd and _mm256_stream_pd.
What could be the reason for that? Is there a reason for not using 512 bit functions _mm512_load_pd and _mm512_stream_pd instead? Could using the 512 bit functions increase the utilization?

Thanks!

Lingua principale
C
Stelle
1.4k
Fork
198
Metriche di merge delle PR
Nessuna PR unita negli ultimi 30g

Preparare l'ambiente

Questo progetto non fornisce container di sviluppo, Dockerfile né guida per i contributori, quindi l'ambiente è a tuo carico: parti dal suo README e consulta la nostra guida al primo contributo per i passaggi generali.

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di NVIDIA/gdrcopy

Tutte le issue di NVIDIA/gdrcopy

Issue simili

Altre issue su C

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.