Memory transaction size

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I am trying to understand the effects of L1 and L2 caching in CUDA and have some trouble matching some profiling results with the statements in the programming guide. Specifically, I cannot seem to produce results consistent with the 128 byte memory transactions that should be used with L1 caching of global memory accesses enabled (using ‘-Xptxas -dlcm=ca’).

For profiling, I use the following (obviously inefficient) code:

#define NUM_ELEMENTS (1000000)

__device__ int array[NUM_ELEMENTS];
__device__ int sum = 0;

__global__ void kernel_forward()
{
  int s = 0;
  for(int i = 0; i < NUM_ELEMENTS; i++)
    s += array[i];
  sum = s;
}

int main(int argc, char **argv)
{
  kernel_forward<<<1, 1>>>();
  cudaDeviceSynchronize();

  return 0;
}

The results with L2-only caching (‘-Xptxas -dlcm=cg’ or no such argument) are as one might expect, i.e., a cache miss for every 8th 4-byte word:

1                      l2_tex_read_hit_rate                   L2 Hit Rate (Texture Reads)      87.50%      87.50%      87.50%
1                        tex_cache_hit_rate                        Unified Cache Hit Rate       0.00%       0.00%       0.00%
1                           global_hit_rate                               Global Hit Rate       0.00%       0.00%       0.00%
1                    dram_read_transactions               Device Memory Read Transactions      645084      645084      645084
1                      l2_read_transactions                          L2 Read Transactions     1535211     1535211     1535211
1                  l2_tex_read_transactions               L2 Transactions (Texture Reads)     1000000     1000000     1000000

With ‘-Xptxas -dlcm=ca’, I would have expected an L1 hit rate of (1 - 1/32) = 96.875%, due to this statement in the programming guide, referring to CC 2.x but referenced in descriptions of the subsequent CCs as well:

“A cache line is 128 bytes and maps to a 128 byte aligned segment in device memory. Memory accesses that are cached in both L1 and L2 are serviced with 128-byte memory transactions whereas memory accesses that are cached in L2 only are serviced with 32-byte memory transactions.”

However, the results indicate that although cache hits moved to L1, there is still a cache miss for every 8th element:

1                      l2_tex_read_hit_rate                   L2 Hit Rate (Texture Reads)       0.00%       0.00%       0.00%
1                        tex_cache_hit_rate                        Unified Cache Hit Rate      87.50%      87.50%      87.50%
1                           global_hit_rate                               Global Hit Rate      87.50%      87.50%      87.50%
1                    dram_read_transactions               Device Memory Read Transactions      603701      603701      603701
1                      l2_read_transactions                          L2 Read Transactions      608022      608022      608022
1                  l2_tex_read_transactions               L2 Transactions (Texture Reads)      125000      125000      125000

The above results were generated on a GTX 950. However, a GTX 850 Ti and a GTX 1050 showed the same cache hit rates. Am I misinterpreting the results or conducting the wrong experiment, or is the memory transaction size actually always 32 bytes, regardless of the cache configuration?

Thanks for any hints,
PA

2

The actual behavior here varies from one GPU architecture to another, according to my observations.

GPUs of Maxwell and Pascal families have a unified texture/L1. AFAIK texture has always had a 32 byte granularity, effectively (if for no other reason) because it is still flowing through L2, which always has a 32 byte granularity.

Prior architectures (Kepler, Fermi) had a separate dedicated L1 (and/or unified with shared memory), and for these architectures an L1 miss always triggered 128 bytes of downstream traffic.

But an L1 miss on Maxwell/Pascal may only trigger 32 bytes of downstream traffic, depending on the actual global access pattern.