Dram_sectors_read.sum cofusing in Nsight Compute

Accelerated Computing CUDA CUDA Programming and Performance
,
1

hi, I’m just learning the cuda. I have create a test case for testing L1 and L2 cache behavior and encounter some issue.

I try to read 64 bytes from global memory. Firstly it go to L1 cache and miss all two sectors, then it go to L2 cache for these 2 sectors, again it got cache miss, so it finally go to the dram. And then I found that it get 4 sectors from dram, which I assume it should be 2 sector as the granularity of L2 cache is 32-byte.

I use the RTX4000 to do the test.

the kernel is very simple

static __global__ void sumArraysGPU(short* a, short* b, short* res,  int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (k < n) {
        res[i] = a[i];
    }
        
}

image
image1611×431 29.9 KB

2

The GPU might be prefetching a full cache line.

3

Weird, If I make the block size to 1024, the result of dram_sectors_read.sum is always 4 sectors larger than the required sectors which is 64 sectors.

image
image1261×253 13.4 KB

4

I would not test those numbers with such low occupancy: There could be memory accesses like loading parameters or constants, which are just a few bytes overall, but distort your calculation. I am not sure, if this is an issue here, but for example the shared memory bank conflicts are known to have a slight variability for small numbers.

Use a loop and load from such memory addresses, that the data of each iteration is apart (not consecutive) in memory so that possible prefetching can still be seen as effect.

5

The tools and GPU HWPM system are not designed to be as accurate as you are requesting. Reasons you can run into issues.

  1. A metric may be collected over multiple passes. Given no method for 100% deterministic replay this can result in variance/error.
  2. GPU has many independent simultaneous engines that may increment a PM.
  3. On more recent GPUs tools have moved to _realtime metrics. These are not 100% accurate. The error is small for typical sample periods (e.g. ±32 for 10000 cycles) but on a small sample you may not see an increment.
  4. There are hardware features on hardware features. On 100 class HBM GPUs and Ada GPUs L2 has 64B promotion enabled by default. 100 class may also have ECC turned on which can increase traffic.

When testing PMs I generally launch equal warps per SM sub-partition (optimal launch) and produce sufficient work that any other increment reason will result in small variance that will be noise.

6

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.