- Long latency
- Off-chip, read/write access
- Random access causes performance hit (sequential access is better)
- Host can read/write
- Reads are cached
- Writes invalidate the cache
Host can access global and constant variables with
cudaGetSymbolAddress()
cudaGetSymbolSize()
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()
Constants must be declared outside of a function body.