I am writing a simple function that uses the occupancy calculation functions to maximize the occupancy and dynamic shared memory size (in order of priority) of a cooperative kernel launch, but I ran into some interesting issues. I wrote the following MRO: #include "helper_cuda.h" __global__ void f…

Hi , this responds to NVBUG ID 5209234 that we are in tracking with the CUDA engineering team . You are correct that what happens here is “Note that the maximum amount of shared memory per thread block is smaller than the maximum shared memory partition available per SM. The 1 KB of shared memory n…

One can argue about this being an unexpected value/sharp edge instead of an incorrect one/bug given that the user can query and subtract the per-block reserved memory, but at the very least the documentation should be clear about what this function returns.

First, I think it’s important to note that there isn’t always* such thing as “per-block reserved memory”, only per-SM reserved memory. (Update: I was incorrect about there being “reserved shared memory per SM”; there is indeed only “reserved shared memory per block”.) Second, maybe I missed them, b…

cudaGetDeviceProperties() gives you access to reservedSharedMemPerBlock. It being zero for certain architectures/GPUs would not be a problem. I just saw that this entry in the device properties is missing from the struct declaration in the documentation and only listed below that, i.e. another docu…

Upon further consideration, this is a documentation bug at the very least. The documentation for cudaOccupancyAvailableDynamicSMemPerBlockstates that it “returns in *dynamicSmemSize the maximum size of dynamic shared memory to allow numBlocks blocks per SM”, yet actually using dynamicSmemSize bytes …

I’d be curious to see if cudaFuncSetCacheConfig/cudaFuncAttributePreferredSharedMemoryCarveout changes the return value of cudaOccupancyAvailableDynamicSMemPerBlock on Kepler GPUs, or what the return value is if it doesn’t. Does it return the correct value according to the preferred carveout, or doe…

Volta and newer GPUs also have unified shared memory and L1 (+ Texture Cache), so you don’t need ancient Fermi/Kepler to check that. See e.g. 1. NVIDIA Ampere GPU Architecture Tuning Guide — NVIDIA Ampere Tuning Guide 12.8 documentation for your 3090.

It’s also worth noting that it isn’t as simple as “manually subtracting reservedSharedMemPerBlock from the output of cudaOccupancyAvailableDynamicSMemPerBlock()”; you also have to account for the possibility that the returned value is already less than or equal to the “true” maximum. This is the sim…

If func’s cudaFuncAttributeMaxDynamicSharedMemorySize is not set close to the value of cudaDevAttrMaxSharedMemoryPerBlockOptin (e.g. the default is 48KB) and/or numBlocks == 1, you’d be subtracting 1KB of shared memory for no reason.

Another interesting property of the bug is that once you use any static shared memory in the kernel (in a way that con’t be optimized away by the compiler), the problem seems to vanish. I.e. cudaOccupancyAvailableDynamicSMemPerBlock() does respect reserved shared memory, but only if there is also st…

cudaOccupancyAvailableDynamicSMemPerBlock returning incorrect value

Accelerated Computing CUDA CUDA Programming and Performance

paleonix April 8, 2025, 6:52pm 12

Honestly, after playing around with the code and reading the docs again, I’m not quite sure what I got hung up on earlier. This is clearly a functional bug. The docs and even the API name contain “dynamic shared memory” and there is no reason to believe reserved shared memory is part of that. Have you filed it as a bug already?

Topic		Replies	Views
How to use cudaOccupancyMaxPotentialBlockSizeVariableSMem CUDA Programming and Performance	7	694	October 17, 2023
Dynamic shared memory calculated by ncu larger than Max_shared_memory_per_block Nsight Compute cuda	3	869	September 7, 2023
Default value of max dynamic shared memory CUDA Programming and Performance cuda	8	466	December 23, 2024
shared memory and CUDA calculator CUDA Programming and Performance	6	4181	October 26, 2008
NCU dynamic shared memory display question Nsight Compute	1	568	April 3, 2024
Question about max shared memory in block and multiprocessor CUDA Programming and Performance	1	2397	February 20, 2024
[ON HOLD] Issue with cuda_occupancy and cudaDeviceSetCacheConfig(...) CUDA Programming and Performance	7	2447	June 26, 2018
Amount of Shared Memory CUDA Programming and Performance	10	4443	June 3, 2010
Find the limit of shared memory that can be used per block CUDA Programming and Performance	2	14309	March 17, 2017
question about dynamic shared memory CUDA Programming and Performance	3	2893	June 7, 2010

cudaOccupancyAvailableDynamicSMemPerBlock returning incorrect value

Related topics