How to reduce the shared memory footprint of the low latency combine kernel?

I am encountering a runtime error when running the `internode_ll::combine kernel` with hidden=2048. The error is `cudaFuncSetAttribute` failed, which is caused by the high shared memory requirement of the kernel's send phase approx. (200KB).

```
RuntimeError: Failed: Assertion error DeepEP/csrc/kernels/internode_ll.cu:985 'cudaFuncSetAttribute(combine_func, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size) == cudaSuccess'
```
I identified three main parameters that control the shared memory size: `hidden`,`kNumStages`,`kNumMaxUnrolls`

To resolve the runtime error, I need to reduce the shared memory allocation. Given that the hidden size is fixed(2048), which of the other two parameters is the intended or safest one to adjust?