Changes

* __constant__ (CUDA) or __constant (OpenCL) memory -- Constants are not allowed to change during the execution of a particular kernel. Historically, this was used by Pixel Shaders as they read texture data. The texture-data could be computed and loaded onto the GPU, but the data was not allowed to change during the Pixel Shader's execution. Both NVidia and AMD GPUs have special caches (and in AMD's case: special registers called sGPRs) which accelerate constant-data. The caches associated with this memory space is sometimes called K$ (Konstant-cache), but have difficult caching issues and has to be independently flushed if the its data ever changes. The main benefit in both AMD and NVidia systems is that K$ values are broadcast extremely efficiently to all threads in a wavefront, but only if all threads in a wavefront are reading from the same memory location. Instead of haing 32-memory reads (NVidia) or 64-memory reads (AMD GCN), a read from K$ can be optimized into a single-read, broadcast to all 32 or 64-threads of a Warp or Wavefront.

* __shared__ (CUDA) or __local__ __local (OpenCL) memory -- This is highly-accelerated memory regions designed for threads to exchange data within a CUDA Block or OpenCL Workgroup. On AMD Systems, there is more Local "LDS" memory than even L1 Cache (GCN) or L0 Cache (RDNA).

* Default (CUDA) or Private __private (OpenCL) Memory -- Private memory typically maps to a GPU-register, and is inaccessible to other threads. If a kernel requires more memory than what can exist in GPU-registers, the data will automatically spill over into global VRAM (with an associated performance penalty). In practice, this spillover is well interleaved, well-optimized, and reduced to as small a subset as possible through compiler optimizations.