Skip to content
Merged
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
changes as per suggestions by @gmarkall
  • Loading branch information
k1m190r committed Mar 22, 2023
commit 217ae5cb2abcd4ce552a8f2685259c21fd4d4607
33 changes: 23 additions & 10 deletions docs/source/cuda/memory.rst
Original file line number Diff line number Diff line change
Expand Up @@ -159,7 +159,8 @@ traditional dynamic memory management.
Dynamic Shared Memory
---------------------

In order to use dynamic shared memory in kernel code declare shared array of size 0:
In order to use dynamic shared memory in kernel code declare a shared array of
size 0:

.. code-block:: python

Expand All @@ -168,19 +169,21 @@ In order to use dynamic shared memory in kernel code declare shared array of siz
dyn_arr = cuda.shared.array(0, dtype=np.float32)
...

and specify the size of dynamic shared array in bytes during kernel invokation (dispatcher object instantiation):
and specify the size of dynamic shared array in bytes during kernel invocation:
Comment thread
k1m190r marked this conversation as resolved.
Resolved

.. code-block:: python

kernel_func[32, 32, 0, 128](x)

In above code dispatcher object is instantiated with 4 parameters:
In the above code the kernel launch is configured with 4 parameters:

.. code-block:: python

kernel_func[grid_dim, block_dim, stream_id, dyn_shared_mem_size]
kernel_func[grid_dim, block_dim, stream, dyn_shared_mem_size]

**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. For example, consider:
**Note:** that all dynamic shared memory arrays *alias*, so if you want to have
Comment thread
k1m190r marked this conversation as resolved.
Resolved
multiple dynamic shared arrays, you need to take *disjoint* views of the arrays.
For example, consider:

.. code-block:: python

Expand All @@ -198,19 +201,27 @@ In above code dispatcher object is instantiated with 4 parameters:
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:
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:

.. code-block:: pycon
.. code-block:: python

3.140000
1078523331

because `1078523331` is the `int32` represented by the bits of the `float32` IEEE754 `3.14`.
because 1078523331 is the ``int32`` represented by the bits of the ``float32``
value 3.14.

If we take disjoint views of the dynamic shared memory:

.. code-block:: python

from numba import cuda
import numpy as np

@cuda.jit
Comment thread
k1m190r marked this conversation as resolved.
Resolved
def f_with_view():
f32_arr = cuda.shared.array(0, dtype=np.float32)
Expand All @@ -223,9 +234,11 @@ If we take disjoint views of the dynamic shared memory:
f_with_view[1, 1, 0, 8]()
cuda.synchronize()

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

.. code-block:: pycon
.. code-block:: python

3.140000
1
Expand Down