Skip to content

Commit

Permalink
Docs: Fix unified memory code highlighting
Browse files Browse the repository at this point in the history
  • Loading branch information
@MKKnorr
MKKnorr committed Oct 25, 2024
1 parent 454511c commit 74ebf79
Showing 1 changed file with 58 additions and 64 deletions.
122 changes: 58 additions & 64 deletions docs/how-to/hip_runtime_api/memory_management/unified_memory.rst
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,8 +37,6 @@ model is shown in the following figure.

.. figure:: ../../../data/how-to/hip_runtime_api/memory_management/unified_memory/um.svg

.. _Pinned host memory can be made accessible to the device, meaning the device can read from or write to the pinned memory on the host via the PCIe bus without actually copying the memory to the device. This can be useful for sparse accesses, when it is not known, what regions of the memroy will be accessed by a kernel. This however has the disadvantage, that the memory has to be read from the host over the PCIe bus every time it is accessed (if it's not in the GPUs cache any more), as it is not stored in device memory. All of this however has nothing to do with unified memory. It doesn't even need unified virtual addressing, the host-pointer has to be specifically mapped to the device.
Unified memory enables the access to memory located on other devices via
several methods, depending on whether hardware support is available or has to be
managed by the driver.
Expand All @@ -59,9 +57,6 @@ allocated and used for a GPU, than the GPU itself has physically available.
This level of unified memory support can be very beneficial for sparse accesses
to an array, that is not often used on the device.


.. _What about mapping memory? Does this require hardware support? This is also a form of unified memory
Driver managed page migration
--------------------------------------------------------------------------------

Expand Down Expand Up @@ -267,7 +262,7 @@ explicit memory management example is presented in the last tab.
.. tab-item:: hipMallocManaged()

.. code-block:: cpp
:emphasize-lines: 12-15
:emphasize-lines: 22-25
#include <hip/hip_runtime.h>
#include <iostream>
Expand Down Expand Up @@ -319,19 +314,19 @@ explicit memory management example is presented in the last tab.
.. tab-item:: __managed__

.. code-block:: cpp
:emphasize-lines: 9-10
:emphasize-lines: 19-20
#include <hip/hip_runtime.h>
#include <iostream>
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
// Addition of two values.
Expand Down Expand Up @@ -362,22 +357,21 @@ explicit memory management example is presented in the last tab.
.. tab-item:: new

.. code-block:: cpp
:emphasize-lines: 12-15
:emphasize-lines: 20-23
#include <hip/hip_runtime.h>
#include <iostream>
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
// Addition of two values.
__global__ void add(int* a, int* b, int* c) {
*c = *a + *b;
Expand Down Expand Up @@ -413,19 +407,19 @@ explicit memory management example is presented in the last tab.
.. tab-item:: Explicit Memory Management

.. code-block:: cpp
:emphasize-lines: 17-24, 29-30
:emphasize-lines: 27-34, 39-40
#include <hip/hip_runtime.h>
#include <iostream>
#define HIP_CHECK(func) \
{ \
const hipError_t err = func; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
// Addition of two values.
Expand Down Expand Up @@ -506,20 +500,20 @@ application by moving data to the desired device before it's actually
needed. ``hipCpuDeviceId`` is a special constant to specify the CPU as target.

.. code-block:: cpp
:emphasize-lines: 20-23,31-32
:emphasize-lines: 33-36,41-42
#include <hip/hip_runtime.h>
#include <iostream>
#define HIP_CHECK(func) \
{ \
const hipError_t err = func; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
// Addition of two values.
__global__ void add(int *a, int *b, int *c) {
Expand Down Expand Up @@ -589,20 +583,20 @@ The following is the updated version of the example above with memory advice
instead of prefetching.

.. code-block:: cpp
:emphasize-lines: 17-26
:emphasize-lines: 29-41
#include <hip/hip_runtime.h>
#include <iostream>
#define HIP_CHECK(func) \
{ \
const hipError_t err = func; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
// Addition of two values.
__global__ void add(int *a, int *b, int *c) {
Expand Down Expand Up @@ -661,20 +655,20 @@ Memory range attributes
memory range. The attributes are given in :cpp:enum:`hipMemRangeAttribute`.

.. code-block:: cpp
:emphasize-lines: 29-34
:emphasize-lines: 44-49
#include <hip/hip_runtime.h>
#include <iostream>
#define HIP_CHECK(func) \
{ \
const hipError_t err = func; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
#define HIP_CHECK(expression) \
{ \
const hipError_t err = expression; \
if(err != hipSuccess){ \
std::cerr << "HIP error: " \
<< hipGetErrorString(err) \
<< " at " << __LINE__ << "\n"; \
} \
}
// Addition of two values.
__global__ void add(int *a, int *b, int *c) {
Expand Down

0 comments on commit 74ebf79

Please sign in to comment.