Maybe I am confused.
The GPU’s cosntant cache is a small cache with a broadcast feature, generally used for data declared constant. I am not entirely sure whether this constant cache still exists in Maxwell, but I think it does. I am not aware that __ldg() can be used to read through the constant cache. Is there documentation that says LDG.CI performs a read through the constant cache?
LDG in Kepler was designed to read through the texture/read-only cache. “const restrict” pointers make it easier for the compiler to determine that it is safe to generate LDGs automatically without the explicit use of the __ldg() intrinsic. In Maxwell, the texture and L1 caches have been merged, so presumably LDG now reads through that combined cache.