Question about the cacheConfig value in nsight systems

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello!

From this guide, it seems that the cacheConfig value reported by NSight Systems should correspond to one of the ENUM_CUDA_FUNC_CACHE_CONFIG values. When I use sqlite from the nsys reports, I see that ENUM_CUDA_FUNC_CACHE_CONFIG has the following values:

[(0, ‘CU_FUNC_CACHE_PREFER_NONE’, ‘None’), (1, ‘CU_FUNC_CACHE_PREFER_SHARED’, ‘Shared’), (2, ‘CU_FUNC_CACHE_PREFER_L1’, ‘L1’), (3, ‘CU_FUNC_CACHE_PREFER_EQUAL’, ‘Equal’)]

However, when I use sqlite to get the cacheConfig, some kernels show a value of 4.Could you please explain why this is happening, or if I am missing something?

Thank you!

2

I suggest questions like this be asked on the nsight systems forum.

1 Like
3

Have you access to those kernels and can you programmatically read their func cache config with the Cuda API instead of Nsight Compute?

4

Yes, I have access to these kernels and I can set the shared memory preference with cudaFuncSetCacheConfig or cudaFuncSetAttribute/cudaFuncAttributePreferredSharedMemoryCarveout
Ideally, I want to check if NSYS can help with also kernels that I don’t have access to, this is why I am checking with some custom setups first.

I am not sure how to get the func cache config that was actually used for the kernel using the CUDA API, do you have any recommendations?

5

There is only a cudaDeviceGetCacheConfig

6

Thank you! For the nsight systems value, do you know if this is the ‘preferred’ or the ‘actual’ cache config?
Also, I see that nsight systems has a sharedMemoryExecuted metric, which seems to be: Shared memory size set by the driver (from the guide). Does this mean that this is the value that the driver carved out for shared memory in each SM? (so the remaining will be L1?)

7

I don’t know (for both questions), perhaps somebody knowledgeable could answer?