# How do I use \`sharedmem\` parameter in kernel?

**URL:** <https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841>\
**Category:** Support: How do I do ...?\
**Created:** [March 16, 2023, 10:17pm UTC](https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841 "2023-03-16T22:17:09Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![biosckon](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/biosckon/32/775_2.png) [@biosckon](https://numba.discourse.group/u/biosckon)\
**Post date:** [March 16, 2023, 10:17pm UTC](https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841/1 "2023-03-16T22:17:09Z")

</div>

Docs describe that dispatcher object can be configured with **sharedmem** :

```python
func[griddim, blockdim, stream, sharedmem]

```

How do I refer to this value of **sharedmem** from within kernel function?

Thank 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:** [March 17, 2023, 12:28pm UTC](https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841/2 "2023-03-17T12:28:40Z")

</div>

You declare a shared array of 0 size, like:

```auto
cuda.shared.array(0, dtype=np.float32)

```

Note that all dynamic shared memory arrays alias, so if you want to have multiple dynamic shared arrays, you need to take disjoint views of the arrays. As an example, consider:

```python
from numba import cuda
import numpy as np

@cuda.jit
def f():
    f32_arr = cuda.shared.array(0, dtype=np.float32)
    i32_arr = cuda.shared.array(0, dtype=np.int32)
    f32_arr[0] = 3.14
    print(f32_arr[0])
    print(i32_arr[0])

f[1, 1, 0, 4]()
cuda.synchronize()

```

This allocates 4 bytes of shared memory (large enough for one `int32` or one `float32`) and declares dynamic shared memory arrays of type `int32 ` and of type `float32`. When `f32_arr[0]` is set, this also sets the value of `i32_arr[0]`, because they’re pointing at the same memory. So we see as output:

```auto
$ python repro.py 
3.140000
1078523331

```

because `1078523331` is the `int32` represented by the bits of the `float32` `3.14`. If we take disjoint views of the dynamic shared memory:

```python
@cuda.jit
def f_with_view():
    f32_arr = cuda.shared.array(0, dtype=np.float32)
    i32_arr = cuda.shared.array(0, dtype=np.int32)[1:]
    f32_arr[0] = 3.14
    i32_arr[0] = 1
    print(f32_arr[0])
    print(i32_arr[0])

f_with_view[1, 1, 0, 8]()
cuda.synchronize()

```

This time we declare 8 dynamic shared memory bytes, using the first four for a `float32` value and the next four for an `int32` value - we see we can set both the `int32` and `float32` value without them aliasing:

```auto
$ python repro.py
3.140000
1

```

---

<div class="post-metadata">

**Author:** ![biosckon](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/biosckon/32/775_2.png) [@biosckon](https://numba.discourse.group/u/biosckon)\
**Post date:** [March 17, 2023, 5:28pm UTC](https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841/3 "2023-03-17T17:28:39Z")

</div>

Thank you Graham. This informative answer belongs to the docs! 🙂

Unless it’s already there, somewhere…

---

<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:** [March 18, 2023, 7:57am UTC](https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841/4 "2023-03-18T07:57:09Z")

</div>

You’re right, it does, and it’s not really written anywhere - I think when the docs were first written it was assumed that readers would learn or know of the semantics from the CUDA C programming guide, but these days it’s reasonable to expect to only need to use the Numba docs for things like this.

---

<div class="post-metadata">

**Author:** ![biosckon](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/biosckon/32/775_2.png) [@biosckon](https://numba.discourse.group/u/biosckon)\
**Post date:** [March 18, 2023, 11:11am UTC](https://numba.discourse.group/t/how-do-i-use-sharedmem-parameter-in-kernel/1841/5 "2023-03-18T11:11:12Z")

</div>

Indeed and I have just enough of CUDA C to be horrible at it but I wasn’t certain. My first instinct was to set the size to 0 but didn’t dare to be certain…Anyway if it’s up in docs it’ll help. Thanks!
