Paměťový model CUDA

Hierarchie pamětí CUDA

registry vlákna
- 65k registrů na SM → ~ 32 reg na vlákno
lok. paměť : zásobník (int x[5];), uint3
- je cachovaná L1 - existuje vhodný přístupový vzor pro přístup
sdílená paměť bloku
- __shared float[10]; - vidí všechna vlákna
- změna přístupové vzoru
  - správný přístupový vzor, vedlejší vlákna přistupují do vedlejší paměti
  - dlaždice 32x32 do shared proměnné
  - nejhorší - latence 400 taktů na prvek vs 32 taktů
- v 1 taktu až 32 hodnot
globální paměť
- životnost po celou dobu kernelu
texturní paměť
- jiná indexace, než přes load/store
- jiná/nekoherentní cache než u ld/st - nesmím sahat zároveň přes texture a ld/st
- cachuje prostorově - natáhne se cacheline 1 nad a 1 pod ní
- aproximace bodu při rozdílném vzorkováním
konstantní paměť
- cachování spec. cache

Minimalizovat přístupy do globální paměti
Udržovat data pro výpočet co nejblíže procesoru – sdílená paměť/registry
Provádět nad daty uvnitř procesoru co nejvíce výpočtů pro překrytí latence způsobené čtením dat z globální paměti

malloc()
- alokovaný objekt je označen jako pageable
- musí přes mezi buffer, které má pinned stránky - jsou zamknuté proti odswapování = 2x kopie dat
cudaHostAlloc() - předejde se mezi alokaci bufferu

cudaHostAllocMapped (adresa cudaHostGetDevicePointer())
typicky data, které čteme 1x nebo zapisujeme 1x
..
ukládání dat
- kolečko dat naplní do bufferu a pak zapíšeme do host RAM
- místo toho přímo z kolečka naplníme buffer do host RAM
cudaHostAllocManaged
- bin strom místo serializace do pole

blockDim.x * blockDim.x + threadIdx.x i += blockDim.x * gridDim.x