# Configuring how much shared memory is available

**URL:** <https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080>\
**Category:** Support: How do I do ...?\
**Created:** [November 22, 2025, 4:56am UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080 "2025-11-22T04:56:39Z")\
**Posts on this page:** 9\
**Page:** 1

<div class="post-metadata">

**Author:** ![pnk](https://avatars.discourse-cdn.com/v4/letter/p/848f3c/32.png) [@pnk](https://numba.discourse.group/u/pnk)\
**Post date:** [November 22, 2025, 4:56am UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/1 "2025-11-22T04:56:39Z")

</div>

I understand that, in Cuda, “shared memory draws from a pool shared with the L1 cache” and that “the programmer can configure how much is allocated to shared memory, up to 96KB.” Can the Numba Cuda programmer configure this? If so, how?

Thanks!

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [November 24, 2025, 10:37pm UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/2 "2025-11-24T22:37:56Z")

</div>

The quick answer is that yes, the Numba CUDA programmer can configure this, but there’s no good API for it. I’m sure I’ve written about this somewhere before, but I haven’t been able to locate the guidance on this that I gave to someone else who asked the same question.

I’ll continue searching and if I don’t find it, I’ll put together another example illustrating how to configure the shared memory / L1 carve out with Numba-CUDA.

---

<div class="post-metadata">

**Author:** ![pnk](https://avatars.discourse-cdn.com/v4/letter/p/848f3c/32.png) [@pnk](https://numba.discourse.group/u/pnk)\
**Post date:** [November 25, 2025, 11:54pm UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/3 "2025-11-25T23:54:49Z")

</div>

I would greatly appreciate it!

---

<div class="post-metadata">

**Author:** ![pnk](https://avatars.discourse-cdn.com/v4/letter/p/848f3c/32.png) [@pnk](https://numba.discourse.group/u/pnk)\
**Post date:** [November 28, 2025, 3:38am UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/4 "2025-11-28T03:38:29Z")

</div>

Were you able to find the guidance? I’m kinda stuck; I’ve augmented the algorithm to handle larger data but, for the modified code, the number of registers is too large so I get CUDA\_ERROR\_COOPERATIVE\_LAUNCH\_TOO\_LARGE.  
I think configuring the shared memory to be slightly larger would enable the code to handle the dataset i need it to handle.

I’d be very grateful to get guidance on this as soon as convenient for you.

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [November 28, 2025, 12:15pm UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/5 "2025-11-28T12:15:42Z")

</div>

I didn’t find an example. There are some bits of functionality for this in Numba-CUDA, but it turns out they’re untested and perhaps bitrotted. I was trying:

```python

from numba import cuda, types

sig = (types.float32[::1], types.float32[::1])

@cuda.jit(sig)
def add_one(r, x):
    i = cuda.grid(1)
    if i < len(r):
        r[i] = x[i] + 1

kernel = add_one.overloads[sig]
cufunc = kernel._codelibrary.get_cufunc()
cufunc.cache_config(prefer_shared=True)

```

but this fails with:

```auto
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/issues/discourse-3080/repro.py", line 13, in <module>
    cufunc.cache_config(prefer_shared=True)
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/cudadrv/driver.py", line 2418, in cache_config
    flag = attr.CU_FUNC_CACHE_PREFER_SHARED
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
AttributeError: type object 'CUfunction_attribute' has no attribute 'CU_FUNC_CACHE_PREFER_SHARED'

```

I suspect it worked with the ctypes bindings but no longer works now we use the CUDA Python bindings.

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [November 28, 2025, 1:01pm UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/6 "2025-11-28T13:01:46Z")

</div>

I’ve created a feature request on the Numba-CUDA issue tracker for this: [[FEA] Support for L1 cache / shared memory configuration · Issue #617 · NVIDIA/numba-cuda · GitHub](https://github.com/NVIDIA/numba-cuda/issues/617)

---

<div class="post-metadata">

**Author:** ![pnk](https://avatars.discourse-cdn.com/v4/letter/p/848f3c/32.png) [@pnk](https://numba.discourse.group/u/pnk)\
**Post date:** [November 29, 2025, 4:41pm UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/7 "2025-11-29T16:41:24Z")

</div>

Thanks. Is there even a hacky approach to doing it that will tide me over until a feature is added to Numba? Maybe involving calling some c code?

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [November 30, 2025, 11:20pm UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/8 "2025-11-30T23:20:43Z")

</div>

You’re right, there will be a hacky way, but it won’t need C code. Just using ctypes should be enough. Let me try a couple of things and get back to you.

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [December 1, 2025, 9:53am UTC](https://numba.discourse.group/t/configuring-how-much-shared-memory-is-available/3080/9 "2025-12-01T09:53:35Z")

</div>

I found this workaround: [[FEA] Provide a way to set the maximum dynamic shared memory size · Issue #94 · NVIDIA/numba-cuda · GitHub](https://github.com/NVIDIA/numba-cuda/issues/94#issuecomment-2553864989) - I need to try it to make sure it still works, but it provides a starting point for me to fix up if not.