SWDEV-514685 - Update documentation 2025-02-11 (#25)
Co-authored-by: Istvan Kiss <neon60@gmail.com>
[ROCm/hip commit: 69becfc7a8]
This commit is contained in:
Kiss, Istvan
GitHub
@@ -1,5 +1,6 @@
|
||||
.*
|
||||
!.gitignore
|
||||
!.spellcheck.local.yaml
|
||||
*.o
|
||||
*.exe
|
||||
*.swp
|
||||
|
||||
@@ -0,0 +1,10 @@
|
||||
matrix:
|
||||
- name: Markdown
|
||||
sources:
|
||||
- ['!docs/doxygen/mainpage.md']
|
||||
- name: reST
|
||||
sources:
|
||||
- []
|
||||
- name: Cpp
|
||||
sources:
|
||||
- ['include/hip/*']
|
||||
@@ -7,16 +7,19 @@ APUs
|
||||
AQL
|
||||
AXPY
|
||||
asm
|
||||
Asynchrony
|
||||
asynchrony
|
||||
backtrace
|
||||
Bitcode
|
||||
bitcode
|
||||
bitcodes
|
||||
blockDim
|
||||
blockIdx
|
||||
builtins
|
||||
Builtins
|
||||
CAS
|
||||
clr
|
||||
compilable
|
||||
constexpr
|
||||
coroutines
|
||||
Ctx
|
||||
cuBLASLt
|
||||
@@ -51,6 +54,7 @@ FNUZ
|
||||
fp
|
||||
gedit
|
||||
GPGPU
|
||||
gridDim
|
||||
GROMACS
|
||||
GWS
|
||||
hardcoded
|
||||
@@ -87,6 +91,7 @@ iteratively
|
||||
Lapack
|
||||
latencies
|
||||
libc
|
||||
libhipcxx
|
||||
libstdc
|
||||
lifecycle
|
||||
linearizing
|
||||
@@ -97,6 +102,7 @@ makefile
|
||||
Malloc
|
||||
malloc
|
||||
MALU
|
||||
maxregcount
|
||||
MiB
|
||||
memset
|
||||
multicore
|
||||
@@ -118,6 +124,7 @@ overindexing
|
||||
oversubscription
|
||||
overutilized
|
||||
parallelizable
|
||||
parallelized
|
||||
pixelated
|
||||
pragmas
|
||||
preallocated
|
||||
@@ -125,6 +132,7 @@ preconditioners
|
||||
predefining
|
||||
prefetched
|
||||
preprocessor
|
||||
printf
|
||||
profilers
|
||||
PTX
|
||||
PyHIP
|
||||
@@ -137,6 +145,7 @@ rocgdb
|
||||
ROCm's
|
||||
rocTX
|
||||
roundtrip
|
||||
rst
|
||||
RTC
|
||||
RTTI
|
||||
rvalue
|
||||
@@ -149,10 +158,12 @@ sinewave
|
||||
SOMA
|
||||
SPMV
|
||||
structs
|
||||
struct's
|
||||
SYCL
|
||||
syntaxes
|
||||
texel
|
||||
texels
|
||||
threadIdx
|
||||
tradeoffs
|
||||
templated
|
||||
toolkits
|
||||
@@ -167,5 +178,6 @@ unregister
|
||||
upscaled
|
||||
variadic
|
||||
vulkan
|
||||
warpSize
|
||||
WinGDB
|
||||
zc
|
||||
|
||||
+16
-23
@@ -12,6 +12,9 @@ Key features include:
|
||||
|
||||
New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms. Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application. HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port.
|
||||
|
||||
> [!NOTE]
|
||||
> The published documentation is available at [HIP documentation](https://rocm.docs.amd.com/projects/HIP/en/latest/index.html) in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the `HIP/docs` folder of this GitHub repository. As with all ROCm projects, the documentation is open source. For more information on contributing to the documentation, see [Contribute to ROCm documentation](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).
|
||||
|
||||
## DISCLAIMER
|
||||
|
||||
The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions, and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard versionchanges, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. Any computer system has risks of security vulnerabilities that cannot be completely prevented or mitigated.AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to time to the content hereof without obligation of AMD to notify any person of such revisions or changes.THIS INFORMATION IS PROVIDED ‘AS IS.” AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS, OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY RELIANCE, DIRECT, INDIRECT, SPECIAL, OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES. AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies.
|
||||
@@ -39,8 +42,8 @@ HIP releases are typically naming convention for each ROCM release to help diffe
|
||||
* [HIP FAQ](docs/faq.rst)
|
||||
* [HIP C++ Language Extensions](docs/reference/cpp_language_extensions.rst)
|
||||
* [HIP Porting Guide](docs/how-to/hip_porting_guide.md)
|
||||
* [HIP Porting Driver Guide](docs/how-to/hip_porting_driver_api.md)
|
||||
* [HIP Programming Guide](docs/how-to/programming_manual.md)
|
||||
* [HIP Porting Driver Guide](docs/how-to/hip_porting_driver_api.rst)
|
||||
* [HIP Programming Guide](docs/programming_guide.rst)
|
||||
* [HIP Logging](docs/how-to/logging.rst)
|
||||
* [Building HIP From Source](docs/install/build.rst)
|
||||
* [HIP Debugging](docs/how-to/debugging.rst)
|
||||
@@ -48,15 +51,15 @@ HIP releases are typically naming convention for each ROCM release to help diffe
|
||||
* [HIP Terminology](docs/reference/terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/OpenCL)
|
||||
* [HIPIFY](https://github.com/ROCm/HIPIFY/blob/amd-staging/README.md)
|
||||
* Supported CUDA APIs:
|
||||
* [Runtime API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUDA_Runtime_API_functions_supported_by_HIP.md)
|
||||
* [Driver API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUDA_Driver_API_functions_supported_by_HIP.md)
|
||||
* [cuComplex API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/cuComplex_API_supported_by_HIP.md)
|
||||
* [Device API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUDA_Device_API_supported_by_HIP.md)
|
||||
* [cuBLAS](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUBLAS_API_supported_by_ROC.md)
|
||||
* [cuRAND](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CURAND_API_supported_by_HIP.md)
|
||||
* [cuDNN](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUDNN_API_supported_by_HIP.md)
|
||||
* [cuFFT](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUFFT_API_supported_by_HIP.md)
|
||||
* [cuSPARSE](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/tables/CUSPARSE_API_supported_by_HIP.md)
|
||||
* [Runtime API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUDA_Runtime_API_functions_supported_by_HIP.md)
|
||||
* [Driver API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUDA_Driver_API_functions_supported_by_HIP.md)
|
||||
* [cuComplex API](https://github.com/ROCm/HIPIFY/blob/amd-staging/reference/docs/tables/cuComplex_API_supported_by_HIP.md)
|
||||
* [Device API](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUDA_Device_API_supported_by_HIP.md)
|
||||
* [cuBLAS](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUBLAS_API_supported_by_ROC.md)
|
||||
* [cuRAND](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CURAND_API_supported_by_HIP.md)
|
||||
* [cuDNN](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUDNN_API_supported_by_HIP.md)
|
||||
* [cuFFT](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUFFT_API_supported_by_HIP.md)
|
||||
* [cuSPARSE](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/reference/tables/CUSPARSE_API_supported_by_HIP.md)
|
||||
* [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
|
||||
* [Release Notes](RELEASE.md)
|
||||
|
||||
@@ -124,19 +127,9 @@ provides source portability to either platform. HIP provides the _hipcc_ compi
|
||||
|
||||
## Examples and Getting Started
|
||||
|
||||
* A sample and [blog](https://github.com/ROCm/hip-tests/tree/develop/samples/0_Intro/square) that uses any of [HIPIFY](https://github.com/ROCm/HIPIFY/blob/amd-staging/README.md) tools to convert a simple app from CUDA to HIP:
|
||||
* The [ROCm-examples](https://github.com/ROCm/rocm-examples) repository includes many examples with explanations that help users getting started with HIP, as well as providing advanced examples for HIP and its libraries.
|
||||
|
||||
```shell
|
||||
cd samples/01_Intro/square
|
||||
# follow README / blog steps to hipify the application.
|
||||
```
|
||||
|
||||
* Guide to [Porting a New Cuda Project](https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html#porting-a-new-cuda-project)
|
||||
|
||||
## More Examples
|
||||
|
||||
The GitHub repository [HIP-Examples](https://github.com/ROCm/HIP-Examples) contains a hipified version of benchmark suite.
|
||||
Besides, there are more samples in Github [HIP samples](https://github.com/ROCm/hip-tests/tree/develop/samples), showing how to program with different features, build and run.
|
||||
* HIP's documentation includes a guide for [Porting a New Cuda Project](https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html#porting-a-new-cuda-project).
|
||||
|
||||
## Tour of the HIP Directories
|
||||
|
||||
|
||||
@@ -6,3 +6,4 @@
|
||||
/doxygen/html
|
||||
/doxygen/xml
|
||||
/sphinx/_toc.yml
|
||||
__pycache__
|
||||
|
||||
@@ -5,6 +5,8 @@
|
||||
# https://www.sphinx-doc.org/en/master/usage/configuration.html
|
||||
|
||||
import re
|
||||
import sys
|
||||
from pathlib import Path
|
||||
from typing import Any, Dict, List
|
||||
|
||||
from rocm_docs import ROCmDocs
|
||||
@@ -38,7 +40,10 @@ external_projects_current_project = "hip"
|
||||
for sphinx_var in ROCmDocs.SPHINX_VARS:
|
||||
globals()[sphinx_var] = getattr(docs_core, sphinx_var)
|
||||
|
||||
extensions += ["sphinxcontrib.doxylink"]
|
||||
# Add the _extensions directory to Python's search path
|
||||
sys.path.append(str(Path(__file__).parent / 'extension'))
|
||||
|
||||
extensions += ["sphinxcontrib.doxylink", "custom_directive"]
|
||||
|
||||
cpp_id_attributes = ["__global__", "__device__", "__host__", "__forceinline__", "static"]
|
||||
cpp_paren_attributes = ["__declspec"]
|
||||
@@ -50,5 +55,6 @@ numfig = False
|
||||
exclude_patterns = [
|
||||
"doxygen/mainpage.md",
|
||||
"understand/glossary.md",
|
||||
'how-to/debugging_env.rst'
|
||||
'how-to/debugging_env.rst',
|
||||
"data/env_variables_hip.rst"
|
||||
]
|
||||
@@ -0,0 +1,263 @@
|
||||
.. meta::
|
||||
:description: HIP environment variables
|
||||
:keywords: AMD, HIP, environment variables, environment
|
||||
|
||||
The GPU isolation environment variables in HIP are collected in the following table.
|
||||
|
||||
.. _hip-env-isolation:
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 70,30
|
||||
|
||||
* - **Environment variable**
|
||||
- **Value**
|
||||
|
||||
* - | ``ROCR_VISIBLE_DEVICES``
|
||||
| A list of device indices or UUIDs that will be exposed to applications.
|
||||
- Example: ``0,GPU-DEADBEEFDEADBEEF``
|
||||
|
||||
* - | ``GPU_DEVICE_ORDINAL``
|
||||
| Devices indices exposed to OpenCL and HIP applications.
|
||||
- Example: ``0,2``
|
||||
|
||||
* - | ``HIP_VISIBLE_DEVICES`` or ``CUDA_VISIBLE_DEVICES``
|
||||
| Device indices exposed to HIP applications.
|
||||
- Example: ``0,2``
|
||||
|
||||
The profiling environment variables in HIP are collected in the following table.
|
||||
|
||||
.. _hip-env-prof:
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 70,30
|
||||
|
||||
* - **Environment variable**
|
||||
- **Value**
|
||||
|
||||
* - | ``HSA_CU_MASK``
|
||||
| Sets the mask on a lower level of queue creation in the driver,
|
||||
| this mask will also be set for queues being profiled.
|
||||
- Example: ``1:0-8``
|
||||
|
||||
* - | ``ROC_GLOBAL_CU_MASK``
|
||||
| Sets the mask on queues created by the HIP or the OpenCL runtimes,
|
||||
| this mask will also be set for queues being profiled.
|
||||
- Example: ``0xf``, enables only 4 CUs
|
||||
|
||||
* - | ``HIP_FORCE_QUEUE_PROFILING``
|
||||
| Used to run the app as if it were run in rocprof. Forces command queue
|
||||
| profiling on by default.
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
The debugging environment variables in HIP are collected in the following table.
|
||||
|
||||
.. _hip-env-debug:
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 35,14,51
|
||||
|
||||
* - **Environment variable**
|
||||
- **Default value**
|
||||
- **Value**
|
||||
|
||||
* - | ``AMD_LOG_LEVEL``
|
||||
| Enables HIP log on various level.
|
||||
- ``0``
|
||||
- | 0: Disable log.
|
||||
| 1: Enables error logs.
|
||||
| 2: Enables warning logs next to lower-level logs.
|
||||
| 3: Enables information logs next to lower-level logs.
|
||||
| 4: Enables debug logs next to lower-level logs.
|
||||
| 5: Enables debug extra logs next to lower-level logs.
|
||||
|
||||
* - | ``AMD_LOG_LEVEL_FILE``
|
||||
| Sets output file for ``AMD_LOG_LEVEL``.
|
||||
- stderr output
|
||||
-
|
||||
|
||||
* - | ``AMD_LOG_MASK``
|
||||
| Specifies HIP log filters. Here is the ` complete list of log masks <https://github.com/ROCm/clr/blob/develop/rocclr/utils/debug.hpp#L40>`_.
|
||||
- ``0x7FFFFFFF``
|
||||
- | 0x1: Log API calls.
|
||||
| 0x2: Kernel and copy commands and barriers.
|
||||
| 0x4: Synchronization and waiting for commands to finish.
|
||||
| 0x8: Decode and display AQL packets.
|
||||
| 0x10: Queue commands and queue contents.
|
||||
| 0x20: Signal creation, allocation, pool.
|
||||
| 0x40: Locks and thread-safety code.
|
||||
| 0x80: Kernel creations and arguments, etc.
|
||||
| 0x100: Copy debug.
|
||||
| 0x200: Detailed copy debug.
|
||||
| 0x400: Resource allocation, performance-impacting events.
|
||||
| 0x800: Initialization and shutdown.
|
||||
| 0x1000: Misc debug, not yet classified.
|
||||
| 0x2000: Show raw bytes of AQL packet.
|
||||
| 0x4000: Show code creation debug.
|
||||
| 0x8000: More detailed command info, including barrier commands.
|
||||
| 0x10000: Log message location.
|
||||
| 0x20000: Memory allocation.
|
||||
| 0x40000: Memory pool allocation, including memory in graphs.
|
||||
| 0x80000: Timestamp details.
|
||||
| 0xFFFFFFFF: Log always even mask flag is zero.
|
||||
|
||||
* - | ``HIP_LAUNCH_BLOCKING``
|
||||
| Used for serialization on kernel execution.
|
||||
- ``0``
|
||||
- | 0: Disable. Kernel executes normally.
|
||||
| 1: Enable. Serializes kernel enqueue, behaves the same as ``AMD_SERIALIZE_KERNEL``.
|
||||
|
||||
* - | ``HIP_VISIBLE_DEVICES`` (or ``CUDA_VISIBLE_DEVICES``)
|
||||
| Only devices whose index is present in the sequence are visible to HIP
|
||||
- Unset by default.
|
||||
- 0,1,2: Depending on the number of devices on the system.
|
||||
|
||||
* - | ``GPU_DUMP_CODE_OBJECT``
|
||||
| Dump code object.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``AMD_SERIALIZE_KERNEL``
|
||||
| Serialize kernel enqueue.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Wait for completion before enqueue.
|
||||
| 2: Wait for completion after enqueue.
|
||||
| 3: Both
|
||||
|
||||
* - | ``AMD_SERIALIZE_COPY``
|
||||
| Serialize copies
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Wait for completion before enqueue.
|
||||
| 2: Wait for completion after enqueue.
|
||||
| 3: Both
|
||||
|
||||
* - | ``AMD_DIRECT_DISPATCH``
|
||||
| Enable direct kernel dispatch (Currently for Linux; under development for Windows).
|
||||
- ``1``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``GPU_MAX_HW_QUEUES``
|
||||
| The maximum number of hardware queues allocated per device.
|
||||
- ``4``
|
||||
- The variable controls how many independent hardware queues HIP runtime can create per process,
|
||||
per device. If an application allocates more HIP streams than this number, then HIP runtime reuses
|
||||
the same hardware queues for the new streams in a round-robin manner. Note that this maximum
|
||||
number does not apply to hardware queues that are created for CU-masked HIP streams, or
|
||||
cooperative queues for HIP Cooperative Groups (single queue per device).
|
||||
|
||||
The memory management related environment variables in HIP are collected in the
|
||||
following table.
|
||||
|
||||
.. _hip-env-memory:
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 35,14,51
|
||||
|
||||
* - **Environment variable**
|
||||
- **Default value**
|
||||
- **Value**
|
||||
|
||||
* - | ``HIP_HIDDEN_FREE_MEM``
|
||||
| Amount of memory to hide from the free memory reported by hipMemGetInfo.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| Unit: megabyte (MB)
|
||||
|
||||
* - | ``HIP_HOST_COHERENT``
|
||||
| Specifies if the memory is coherent between the host and GPU in ``hipHostMalloc``.
|
||||
- ``0``
|
||||
- | 0: Memory is not coherent.
|
||||
| 1: Memory is coherent.
|
||||
| Environment variable has effect, if the following conditions are statisfied:
|
||||
| - One of the ``hipHostMallocDefault``, ``hipHostMallocPortable``, ``hipHostMallocWriteCombined`` or ``hipHostMallocNumaUser`` flag set to 1.
|
||||
| - ``hipHostMallocCoherent``, ``hipHostMallocNonCoherent`` and ``hipHostMallocMapped`` flags set to 0.
|
||||
|
||||
* - | ``HIP_INITIAL_DM_SIZE``
|
||||
| Set initial heap size for device malloc.
|
||||
- ``8388608``
|
||||
- | Unit: Byte
|
||||
| The default value corresponds to 8 MB.
|
||||
|
||||
* - | ``HIP_MEM_POOL_SUPPORT``
|
||||
| Enables memory pool support in HIP.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``HIP_MEM_POOL_USE_VM``
|
||||
| Enables memory pool support in HIP.
|
||||
- | ``0``: other OS
|
||||
| ``1``: Windows
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``HIP_VMEM_MANAGE_SUPPORT``
|
||||
| Virtual Memory Management Support.
|
||||
- ``1``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``GPU_MAX_HEAP_SIZE``
|
||||
| Set maximum size of the GPU heap to % of board memory.
|
||||
- ``100``
|
||||
- | Unit: Percentage
|
||||
|
||||
* - | ``GPU_MAX_REMOTE_MEM_SIZE``
|
||||
| Maximum size that allows device memory substitution with system.
|
||||
- ``2``
|
||||
- | Unit: kilobyte (KB)
|
||||
|
||||
* - | ``GPU_NUM_MEM_DEPENDENCY``
|
||||
| Number of memory objects for dependency tracking.
|
||||
- ``256``
|
||||
-
|
||||
|
||||
* - | ``GPU_STREAMOPS_CP_WAIT``
|
||||
| Force the stream memory operation to wait on CP.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``HSA_LOCAL_MEMORY_ENABLE``
|
||||
| Enable HSA device local memory usage.
|
||||
- ``1``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``PAL_ALWAYS_RESIDENT``
|
||||
| Force memory resources to become resident at allocation time.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``PAL_PREPINNED_MEMORY_SIZE``
|
||||
| Size of prepinned memory.
|
||||
- ``64``
|
||||
- | Unit: kilobyte (KB)
|
||||
|
||||
* - | ``REMOTE_ALLOC``
|
||||
| Use remote memory for the global heap allocation.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
The following table lists environment variables that are useful but relate to
|
||||
different features in HIP.
|
||||
|
||||
.. _hip-env-other:
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 35,14,51
|
||||
|
||||
* - **Environment variable**
|
||||
- **Default value**
|
||||
- **Value**
|
||||
|
||||
* - | ``HIPRTC_COMPILE_OPTIONS_APPEND``
|
||||
| Sets compile options needed for ``hiprtc`` compilation.
|
||||
- None
|
||||
- ``--gpu-architecture=gfx906:sramecc+:xnack``, ``-fgpu-rdc``
|
||||
+274
@@ -0,0 +1,274 @@
|
||||
<mxfile host="65bd71144e">
|
||||
<diagram id="zBbb_w2fufU70cdOGtND" name="1 oldal">
|
||||
<mxGraphModel dx="1310" dy="1222" grid="1" gridSize="10" guides="1" tooltips="1" connect="1" arrows="1" fold="1" page="0" pageScale="1" pageWidth="660" pageHeight="610" background="none" math="0" shadow="0">
|
||||
<root>
|
||||
<mxCell id="0"/>
|
||||
<mxCell id="1" parent="0"/>
|
||||
<mxCell id="5965" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#5E5B61;fontColor=#FFFFFF;strokeColor=none;spacing=0;" parent="1" vertex="1">
|
||||
<mxGeometry y="-190" width="680" height="610" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5966" value="<font face="Helvetica">time</font>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="10" y="380" width="55" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5967" value="" style="endArrow=classic;startArrow=none;html=1;rounded=0;strokeWidth=2;startFill=0;strokeColor=#FFFFFF;exitX=1;exitY=0.5;exitDx=0;exitDy=0;" parent="1" source="5966" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="510" y="140" as="sourcePoint"/>
|
||||
<mxPoint x="671" y="396" as="targetPoint"/>
|
||||
<Array as="points">
|
||||
<mxPoint x="190" y="396"/>
|
||||
</Array>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="5968" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#333333;fontColor=#FFFFFF;strokeColor=none;" parent="1" vertex="1">
|
||||
<mxGeometry x="10" y="-130" width="660" height="70" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5969" value="<font face="Helvetica">default stream</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="20" y="-110" width="121" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5970" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="200" y="-120" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5971" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="5970" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5972" value="<div>H2D<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5970" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5973" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="250" y="-120" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5974" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="5973" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5975" value="<div>H2D<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5973" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5976" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#808080;entryX=0;entryY=1;entryDx=0;entryDy=0;exitX=0;exitY=0;exitDx=0;exitDy=0;" parent="5973" source="5974" target="5974" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="-0.6900000000002251" y="260" as="sourcePoint"/>
|
||||
<mxPoint x="-0.6899999999999977" y="105" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="5977" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="300" y="-120" width="120" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5978" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#C23555;fontColor=#FFFFFF;strokeColor=#A20025;" parent="5977" vertex="1">
|
||||
<mxGeometry width="120.00000000000001" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5979" value="<div>kernel<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5977" vertex="1">
|
||||
<mxGeometry x="4.485465996156356" y="7.499999999999999" width="109.1463392398047" height="29.999999999999996" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5980" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="420" y="-120" width="120" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5981" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#C23555;fontColor=#FFFFFF;strokeColor=#A20025;" parent="5980" vertex="1">
|
||||
<mxGeometry width="120.00000000000001" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5982" value="<div>kernel<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5980" vertex="1">
|
||||
<mxGeometry x="4.485465996156356" y="7.499999999999999" width="109.1463392398047" height="29.999999999999996" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5983" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="540" y="-120" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5984" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="5983" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5985" value="<div>D2H<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5983" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5986" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="590" y="-120" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5987" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="5986" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5988" value="<div>D2H<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5986" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5989" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#808080;entryX=0;entryY=1;entryDx=0;entryDy=0;exitX=0;exitY=0;exitDx=0;exitDy=0;" parent="5986" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="0.8999999999999773" as="sourcePoint"/>
|
||||
<mxPoint x="0.8999999999999773" y="45" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="5990" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#333333;fontColor=#FFFFFF;strokeColor=none;" parent="1" vertex="1">
|
||||
<mxGeometry x="10" y="75" width="660" height="70" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5991" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="250" y="87.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5992" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="5991" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5993" value="<div>H2D<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5991" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5994" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="370" y="87.5" width="120" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5995" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#C23555;fontColor=#FFFFFF;strokeColor=#A20025;" parent="5994" vertex="1">
|
||||
<mxGeometry width="120.00000000000001" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5996" value="<div>kernel<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5994" vertex="1">
|
||||
<mxGeometry x="4.485465996156356" y="7.499999999999999" width="109.1463392398047" height="29.999999999999996" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5997" value="<font face="Helvetica">stream2</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="20" y="95" width="120" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5998" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="490" y="87.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="5999" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="5998" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6000" value="<div>D2H<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="5998" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6001" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#333333;fontColor=#FFFFFF;strokeColor=none;" parent="1" vertex="1">
|
||||
<mxGeometry x="10" y="-10" width="660" height="70" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6002" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="200" y="2.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6003" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="6002" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6004" value="<div>H2D<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6002" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6005" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="250" y="2.5" width="120" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6006" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#C23555;fontColor=#FFFFFF;strokeColor=#A20025;" parent="6005" vertex="1">
|
||||
<mxGeometry width="120.00000000000001" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6007" value="<div>kernel<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6005" vertex="1">
|
||||
<mxGeometry x="4.485465996156356" y="7.499999999999999" width="109.1463392398047" height="29.999999999999996" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6008" value="<font face="Helvetica">stream1</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="20" y="10" width="120" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6009" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="370" y="2.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6010" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="6009" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6011" value="<div>D2H<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6009" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6014" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#808080;entryX=0;entryY=0;entryDx=0;entryDy=0;exitX=0;exitY=1;exitDx=0;exitDy=0;" parent="1" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="249.99999999999955" y="144.99999999999977" as="sourcePoint"/>
|
||||
<mxPoint x="250" y="-10.000000000000227" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="6015" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#808080;entryX=0;entryY=0;entryDx=0;entryDy=0;exitX=0;exitY=1;exitDx=0;exitDy=0;" parent="1" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="370" y="144.99999999999977" as="sourcePoint"/>
|
||||
<mxPoint x="370" y="-10.000000000000227" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="6016" value="<font face="Helvetica">Seqeuntial calls:</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="10" y="-170" width="170" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6017" value="<font face="Helvetica">Asynchronous calls:</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="10" y="-50" width="170" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6018" value="<font face="Helvetica">Asynchronous calls with hipEvent:<br></font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="11" y="160" width="300" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6019" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#333333;fontColor=#FFFFFF;strokeColor=none;" parent="1" vertex="1">
|
||||
<mxGeometry x="11" y="290" width="660" height="70" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6020" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="251" y="302.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6021" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="6020" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6022" value="<div>H2D<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6020" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6023" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="371" y="302.5" width="120" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6024" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#C23555;fontColor=#FFFFFF;strokeColor=#A20025;" parent="6023" vertex="1">
|
||||
<mxGeometry width="120.00000000000001" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6025" value="<div>kernel<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6023" vertex="1">
|
||||
<mxGeometry x="4.485465996156356" y="7.499999999999999" width="109.1463392398047" height="29.999999999999996" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6026" value="<font face="Helvetica">stream2</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="21" y="310" width="120" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6027" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#333333;fontColor=#FFFFFF;strokeColor=none;" parent="1" vertex="1">
|
||||
<mxGeometry x="11" y="205" width="660" height="70" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6028" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="201" y="217.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6029" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="6028" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6030" value="<div>H2D<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6028" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6031" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="251" y="217.5" width="120" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6032" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#C23555;fontColor=#FFFFFF;strokeColor=#A20025;" parent="6031" vertex="1">
|
||||
<mxGeometry width="120.00000000000001" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6033" value="<div>kernel<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6031" vertex="1">
|
||||
<mxGeometry x="4.485465996156356" y="7.499999999999999" width="109.1463392398047" height="29.999999999999996" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6034" value="<font face="Helvetica">stream1</font>" style="text;html=1;strokeColor=none;fillColor=none;align=left;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="21" y="225" width="120" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6035" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="491" y="302.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6036" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="6035" vertex="1">
|
||||
<mxGeometry width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6037" value="<div>D2H<br><font style="font-size: 12px;">data2</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6035" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6040" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#808080;entryX=0;entryY=0;entryDx=0;entryDy=0;exitX=0;exitY=1;exitDx=0;exitDy=0;" parent="1" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="250.30999999999972" y="359.9999999999998" as="sourcePoint"/>
|
||||
<mxPoint x="250.30999999999983" y="204.99999999999977" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="6041" value="<font style="font-size: 12px;" face="Helvetica">event</font>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="341" y="175" width="55" height="30" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6043" value="" style="group" parent="1" vertex="1" connectable="0">
|
||||
<mxGeometry x="372" y="217.5" width="50" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6044" value="" style="rounded=0;whiteSpace=wrap;html=1;fillColor=#4f1623;fontColor=#FFFFFF;strokeColor=none;" parent="6043" vertex="1">
|
||||
<mxGeometry x="-1" width="50.000000000000014" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6045" value="<div>D2H<br><font style="font-size: 12px;">data1</font></div>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="6043" vertex="1">
|
||||
<mxGeometry x="0.8964285714285714" y="8.75" width="48.21071428571429" height="27.5" as="geometry"/>
|
||||
</mxCell>
|
||||
<mxCell id="6047" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#FF0000;entryX=1;entryY=0;entryDx=0;entryDy=0;exitX=1;exitY=1;exitDx=0;exitDy=0;" parent="1" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="370.9999999999998" y="359.9999999999998" as="sourcePoint"/>
|
||||
<mxPoint x="371" y="204.99999999999977" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="6051" value="" style="endArrow=none;html=1;rounded=0;fillColor=#eeeeee;strokeColor=#FF0000;entryX=1;entryY=0;entryDx=0;entryDy=0;exitX=1;exitY=1;exitDx=0;exitDy=0;" parent="1" edge="1">
|
||||
<mxGeometry width="50" height="50" relative="1" as="geometry">
|
||||
<mxPoint x="540.9999999999998" y="359.9999999999998" as="sourcePoint"/>
|
||||
<mxPoint x="541" y="204.99999999999977" as="targetPoint"/>
|
||||
</mxGeometry>
|
||||
</mxCell>
|
||||
<mxCell id="6052" value="<font style="font-size: 12px;" face="Helvetica">eventA<br>eventB<br></font>" style="text;html=1;strokeColor=none;fillColor=none;align=center;verticalAlign=middle;whiteSpace=wrap;rounded=0;fontSize=17;fontColor=#FFFFFF;" parent="1" vertex="1">
|
||||
<mxGeometry x="511" y="160" width="55" height="45" as="geometry"/>
|
||||
</mxCell>
|
||||
</root>
|
||||
</mxGraphModel>
|
||||
</diagram>
|
||||
</mxfile>
|
||||
File diff suppressed because one or more lines are too long
|
After Leveys: | Korkeus: | Koko: 33 KiB |
@@ -0,0 +1,59 @@
|
||||
import os
|
||||
import re
|
||||
from docutils.parsers.rst import Directive
|
||||
from docutils.statemachine import StringList
|
||||
|
||||
class TableInclude(Directive):
|
||||
required_arguments = 1
|
||||
optional_arguments = 0
|
||||
final_argument_whitespace = True
|
||||
option_spec = {
|
||||
'table': str
|
||||
}
|
||||
|
||||
def run(self):
|
||||
# Get the file path from the first argument
|
||||
file_path = self.arguments[0]
|
||||
|
||||
# Get the environment to resolve the full path
|
||||
env = self.state.document.settings.env
|
||||
src_dir = os.path.abspath(env.srcdir)
|
||||
full_file_path = os.path.join(src_dir, file_path)
|
||||
|
||||
# Check if the file exists
|
||||
if not os.path.exists(full_file_path):
|
||||
raise self.error(f"RST file {full_file_path} does not exist.")
|
||||
|
||||
# Read the entire file content
|
||||
with open(full_file_path, 'r', encoding='utf-8') as f:
|
||||
content = f.read()
|
||||
|
||||
# Find all tables with named targets
|
||||
table_pattern = r'(?:^\.\.\ _(.+?):\n)(.. list-table::.*?(?:\n\s*\*\s*-.*?)+)(?=\n\n|\Z)'
|
||||
table_matches = list(re.finditer(table_pattern, content, re.MULTILINE | re.DOTALL))
|
||||
|
||||
# Get the specific table name from options
|
||||
table_name = self.options.get('table')
|
||||
|
||||
# If no table specified, merge compatible tables
|
||||
if not table_name:
|
||||
raise self.error("The ':table:' option is required to specify which table to include.")
|
||||
|
||||
# Find the specific table
|
||||
matching_tables = [
|
||||
match for match in table_matches
|
||||
if match.group(1).strip() == table_name
|
||||
]
|
||||
|
||||
if not matching_tables:
|
||||
raise self.error(f"Table '{table_name}' not found in {full_file_path}")
|
||||
|
||||
# Extract the matched table content
|
||||
table_content = matching_tables[0].group(2)
|
||||
|
||||
# Insert the table content into the current document
|
||||
self.state_machine.insert_input(table_content.splitlines(), full_file_path)
|
||||
return []
|
||||
|
||||
def setup(app):
|
||||
app.add_directive('include-table', TableInclude)
|
||||
@@ -47,7 +47,7 @@ The :doc:`HIP API documentation <doxygen/html/index>` describes each API and
|
||||
its limitations, if any, compared with the equivalent CUDA API.
|
||||
|
||||
The kernel language features are documented in the
|
||||
:doc:`/reference/cpp_language_extensions` page.
|
||||
:doc:`/how-to/hip_cpp_language_extensions` page.
|
||||
|
||||
Relation to other GPGPU frameworks
|
||||
==================================
|
||||
@@ -65,39 +65,8 @@ platforms.
|
||||
Additional porting might be required to deal with architecture feature
|
||||
queries or CUDA capabilities that HIP doesn't support.
|
||||
|
||||
How does HIP compare with OpenCL?
|
||||
---------------------------------
|
||||
|
||||
HIP offers several benefits over OpenCL:
|
||||
|
||||
* Device code can be written in modern C++, including templates, lambdas,
|
||||
classes and so on.
|
||||
* Host and device code can be mixed in the source files.
|
||||
* The HIP API is less verbose than OpenCL and is familiar to CUDA developers.
|
||||
* Porting from CUDA to HIP is significantly easier than from CUDA to OpenCL.
|
||||
* HIP uses development tools specialized for each platform: :doc:`amdclang++ <llvm-project:index>`
|
||||
for AMD GPUs or `nvcc <https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html>`_
|
||||
for NVIDIA GPUs, and profilers like :doc:`ROCm Compute Profiler <rocprofiler-compute:index>` or
|
||||
`Nsight Systems <https://developer.nvidia.com/nsight-systems>`_.
|
||||
* HIP provides
|
||||
* pointers and host-side pointer arithmetic.
|
||||
* device-level control over memory allocation and placement.
|
||||
* an offline compilation model.
|
||||
|
||||
How does porting CUDA to HIP compare to porting CUDA to OpenCL?
|
||||
---------------------------------------------------------------
|
||||
|
||||
OpenCL differs from HIP and CUDA when considering the host runtime,
|
||||
but even more so when considering the kernel code.
|
||||
The HIP device code is a C++ dialect, while OpenCL is C99-based.
|
||||
OpenCL does not support single-source compilation.
|
||||
|
||||
As a result, the OpenCL syntax differs significantly from HIP, and porting tools
|
||||
must perform complex transformations, especially regarding templates or other
|
||||
C++ features in kernels.
|
||||
|
||||
To better understand the syntax differences, see :doc:`here<reference/terms>` or
|
||||
the :doc:`HIP porting guide <how-to/hip_porting_guide>`.
|
||||
To better understand the syntax differences, see :doc:`CUDA to HIP API Function Comparison <reference/api_syntax>`
|
||||
or the :doc:`HIP porting guide <how-to/hip_porting_guide>`.
|
||||
|
||||
Can I install CUDA and ROCm on the same machine?
|
||||
------------------------------------------------
|
||||
|
||||
@@ -273,7 +273,8 @@ HIP environment variable summary
|
||||
|
||||
Here are some of the more commonly used environment variables:
|
||||
|
||||
.. include:: ../how-to/debugging_env.rst
|
||||
.. include-table:: data/env_variables_hip.rst
|
||||
:table: hip-env-debug
|
||||
|
||||
General debugging tips
|
||||
======================================================
|
||||
|
||||
@@ -1,95 +0,0 @@
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 35,14,51
|
||||
|
||||
* - **Environment variable**
|
||||
- **Default value**
|
||||
- **Value**
|
||||
|
||||
* - | ``AMD_LOG_LEVEL``
|
||||
| Enables HIP log on various level.
|
||||
- ``0``
|
||||
- | 0: Disable log.
|
||||
| 1: Enables error logs.
|
||||
| 2: Enables warning logs next to lower-level logs.
|
||||
| 3: Enables information logs next to lower-level logs.
|
||||
| 4: Enables debug logs next to lower-level logs.
|
||||
| 5: Enables debug extra logs next to lower-level logs.
|
||||
|
||||
* - | ``AMD_LOG_LEVEL_FILE``
|
||||
| Sets output file for ``AMD_LOG_LEVEL``.
|
||||
- stderr output
|
||||
-
|
||||
|
||||
* - | ``AMD_LOG_MASK``
|
||||
| Specifies HIP log filters. Here is the ` complete list of log masks <https://github.com/ROCm/clr/blob/develop/rocclr/utils/debug.hpp#L40>`_.
|
||||
- ``0x7FFFFFFF``
|
||||
- | 0x1: Log API calls.
|
||||
| 0x2: Kernel and copy commands and barriers.
|
||||
| 0x4: Synchronization and waiting for commands to finish.
|
||||
| 0x8: Decode and display AQL packets.
|
||||
| 0x10: Queue commands and queue contents.
|
||||
| 0x20: Signal creation, allocation, pool.
|
||||
| 0x40: Locks and thread-safety code.
|
||||
| 0x80: Kernel creations and arguments, etc.
|
||||
| 0x100: Copy debug.
|
||||
| 0x200: Detailed copy debug.
|
||||
| 0x400: Resource allocation, performance-impacting events.
|
||||
| 0x800: Initialization and shutdown.
|
||||
| 0x1000: Misc debug, not yet classified.
|
||||
| 0x2000: Show raw bytes of AQL packet.
|
||||
| 0x4000: Show code creation debug.
|
||||
| 0x8000: More detailed command info, including barrier commands.
|
||||
| 0x10000: Log message location.
|
||||
| 0x20000: Memory allocation.
|
||||
| 0x40000: Memory pool allocation, including memory in graphs.
|
||||
| 0x80000: Timestamp details.
|
||||
| 0xFFFFFFFF: Log always even mask flag is zero.
|
||||
|
||||
* - | ``HIP_LAUNCH_BLOCKING``
|
||||
| Used for serialization on kernel execution.
|
||||
- ``0``
|
||||
- | 0: Disable. Kernel executes normally.
|
||||
| 1: Enable. Serializes kernel enqueue, behaves the same as ``AMD_SERIALIZE_KERNEL``.
|
||||
|
||||
* - | ``HIP_VISIBLE_DEVICES`` (or ``CUDA_VISIBLE_DEVICES``)
|
||||
| Only devices whose index is present in the sequence are visible to HIP
|
||||
- Unset by default.
|
||||
- 0,1,2: Depending on the number of devices on the system.
|
||||
|
||||
* - | ``GPU_DUMP_CODE_OBJECT``
|
||||
| Dump code object.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``AMD_SERIALIZE_KERNEL``
|
||||
| Serialize kernel enqueue.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Wait for completion before enqueue.
|
||||
| 2: Wait for completion after enqueue.
|
||||
| 3: Both
|
||||
|
||||
* - | ``AMD_SERIALIZE_COPY``
|
||||
| Serialize copies
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Wait for completion before enqueue.
|
||||
| 2: Wait for completion after enqueue.
|
||||
| 3: Both
|
||||
|
||||
* - | ``AMD_DIRECT_DISPATCH``
|
||||
| Enable direct kernel dispatch (Currently for Linux; under development for Windows).
|
||||
- ``1``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``GPU_MAX_HW_QUEUES``
|
||||
| The maximum number of hardware queues allocated per device.
|
||||
- ``4``
|
||||
- The variable controls how many independent hardware queues HIP runtime can create per process,
|
||||
per device. If an application allocates more HIP streams than this number, then HIP runtime reuses
|
||||
the same hardware queues for the new streams in a round-robin manner. Note that this maximum
|
||||
number does not apply to hardware queues that are created for CU-masked HIP streams, or
|
||||
cooperative queues for HIP Cooperative Groups (single queue per device).
|
||||
@@ -0,0 +1,922 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the built-in variables and functions that
|
||||
are accessible from HIP kernels and HIP's C++ support. It's
|
||||
intended for users who are familiar with CUDA kernel syntax and
|
||||
want to learn how HIP differs from CUDA.
|
||||
:keywords: AMD, ROCm, HIP, CUDA, c++ language extensions, HIP functions
|
||||
|
||||
################################################################################
|
||||
HIP C++ language extensions
|
||||
################################################################################
|
||||
|
||||
HIP extends the C++ language with additional features designed for programming
|
||||
heterogeneous applications. These extensions mostly relate to the kernel
|
||||
language, but some can also be applied to host functionality.
|
||||
|
||||
********************************************************************************
|
||||
HIP qualifiers
|
||||
********************************************************************************
|
||||
|
||||
Function-type qualifiers
|
||||
================================================================================
|
||||
|
||||
HIP introduces three different function qualifiers to mark functions for
|
||||
execution on the device or the host, and also adds new qualifiers to control
|
||||
inlining of functions.
|
||||
|
||||
.. _host_attr:
|
||||
|
||||
__host__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The ``__host__`` qualifier is used to specify functions for execution
|
||||
on the host. This qualifier is implicitly defined for any function where no
|
||||
``__host__``, ``__device__`` or ``__global__`` qualifier is added, in order to
|
||||
not break compatibility with existing C++ functions.
|
||||
|
||||
You can't combine ``__host__`` with ``__global__``.
|
||||
|
||||
__device__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The ``__device__`` qualifier is used to specify functions for execution on the
|
||||
device. They can only be called from other ``__device__`` functions or from
|
||||
``__global__`` functions.
|
||||
|
||||
You can combine it with the ``__host__`` qualifier and mark functions
|
||||
``__host__ __device__``. In this case, the function is compiled for the host and
|
||||
the device. Note that these functions can't use the HIP built-ins (e.g.,
|
||||
:ref:`threadIdx.x <thread_and_block_idx>` or :ref:`warpSize <warp_size>`), as
|
||||
they are not available on the host. If you need to use HIP grid coordinate
|
||||
functions, you can pass the necessary coordinate information as an argument.
|
||||
|
||||
__global__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Functions marked ``__global__`` are executed on the device and are referred to
|
||||
as kernels. Their return type must be ``void``. Kernels have a special launch
|
||||
mechanism, and have to be launched from the host.
|
||||
|
||||
There are some restrictions on the parameters of kernels. Kernels can't:
|
||||
|
||||
* have a parameter of type ``std::initializer_list`` or ``va_list``
|
||||
* have a variable number of arguments
|
||||
* use references as parameters
|
||||
* use parameters having different sizes in host and device code, e.g. long double arguments, or structs containing long double members.
|
||||
* use struct-type arguments which have different layouts in host and device code.
|
||||
|
||||
Kernels can have variadic template parameters, but only one parameter pack,
|
||||
which must be the last item in the template parameter list.
|
||||
|
||||
.. note::
|
||||
Unlike CUDA, HIP does not support dynamic parallelism, meaning that kernels
|
||||
can not be called from the device.
|
||||
|
||||
Calling __global__ functions
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
|
||||
The launch mechanism for kernels differs from standard function calls, as they
|
||||
need an additional configuration, that specifies the grid and block dimensions
|
||||
(i.e. the amount of threads to be launched), as well as specifying the amount of
|
||||
shared memory per block and which stream to execute the kernel on.
|
||||
|
||||
Kernels are called using the triple chevron ``<<<>>>`` syntax known from CUDA,
|
||||
but HIP also supports the ``hipLaunchKernelGGL`` macro.
|
||||
|
||||
When using ``hipLaunchKernelGGL``, the first five configuration parameters must
|
||||
be:
|
||||
|
||||
* ``symbol kernelName``: The name of the kernel you want to launch. To support
|
||||
template kernels that contain several template parameters separated by use the
|
||||
``HIP_KERNEL_NAME`` macro to wrap the template instantiation
|
||||
(:doc:`HIPIFY <hipify:index>` inserts this automatically).
|
||||
* ``dim3 gridDim``: 3D-grid dimensions that specifies the number of blocks to
|
||||
launch.
|
||||
* ``dim3 blockDim``: 3D-block dimensions that specifies the number of threads in
|
||||
each block.
|
||||
* ``size_t dynamicShared``: The amount of additional shared dynamic memory to
|
||||
allocate per block.
|
||||
* ``hipStream_t``: The stream on which to run the kernel. A value of ``0``
|
||||
corresponds to the default stream.
|
||||
|
||||
The kernel arguments are listed after the configuration parameters.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#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"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Performs a simple initialization of an array with the thread's index variables.
|
||||
// This function is only available in device code.
|
||||
__device__ void init_array(float * const a, const unsigned int arraySize){
|
||||
// globalIdx uniquely identifies a thread in a 1D launch configuration.
|
||||
const int globalIdx = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
// Each thread initializes a single element of the array.
|
||||
if(globalIdx < arraySize){
|
||||
a[globalIdx] = globalIdx;
|
||||
}
|
||||
}
|
||||
|
||||
// Rounds a value up to the next multiple.
|
||||
// This function is available in host and device code.
|
||||
__host__ __device__ constexpr int round_up_to_nearest_multiple(int number, int multiple){
|
||||
return (number + multiple - 1)/multiple;
|
||||
}
|
||||
|
||||
__global__ void example_kernel(float * const a, const unsigned int N)
|
||||
{
|
||||
// Initialize array.
|
||||
init_array(a, N);
|
||||
// Perform additional work:
|
||||
// - work with the array
|
||||
// - use the array in a different kernel
|
||||
// - ...
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int N = 100000000; // problem size
|
||||
constexpr int blockSize = 256; //configurable block size
|
||||
|
||||
//needed number of blocks for the given problem size
|
||||
constexpr int gridSize = round_up_to_nearest_multiple(N, blockSize);
|
||||
|
||||
float *a;
|
||||
// allocate memory on the GPU
|
||||
HIP_CHECK(hipMalloc(&a, sizeof(*a) * N));
|
||||
|
||||
std::cout << "Launching kernel." << std::endl;
|
||||
example_kernel<<<dim3(gridSize), dim3(blockSize), 0/*example doesn't use shared memory*/, 0/*default stream*/>>>(a, N);
|
||||
// make sure kernel execution is finished by synchronizing. The CPU can also
|
||||
// execute other instructions during that time
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
std::cout << "Kernel execution finished." << std::endl;
|
||||
|
||||
HIP_CHECK(hipFree(a));
|
||||
}
|
||||
|
||||
Inline qualifiers
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
HIP adds the ``__noinline__`` and ``__forceinline__`` function qualifiers.
|
||||
|
||||
``__noinline__`` is a hint to the compiler to not inline the function, whereas
|
||||
``__forceinline__`` forces the compiler to inline the function. These qualifiers
|
||||
can be applied to both ``__host__`` and ``__device__`` functions.
|
||||
|
||||
``__noinline__`` and ``__forceinline__`` can not be used in combination.
|
||||
|
||||
__launch_bounds__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
GPU multiprocessors have a fixed pool of resources (primarily registers and
|
||||
shared memory) which are shared by the actively running warps. Using more
|
||||
resources per thread can increase executed instructions per cycle but reduces
|
||||
the resources available for other warps and may therefore limit the occupancy,
|
||||
i.e. the number of warps that can be executed simultaneously. Thus GPUs have to
|
||||
balance resource usage between instruction- and thread-level parallelism.
|
||||
|
||||
``__launch_bounds__`` allows the application to provide hints that influence the
|
||||
resource (primarily registers) usage of the generated code. It is a function
|
||||
attribute that must be attached to a __global__ function:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void __launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_WARPS_PER_EXECUTION_UNIT)
|
||||
kernel_name(/*args*/);
|
||||
|
||||
The ``__launch_bounds__`` parameters are explained in the following sections:
|
||||
|
||||
MAX_THREADS_PER_BLOCK
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
|
||||
This parameter is a guarantee from the programmer, that kernel will not be
|
||||
launched with more threads than ``MAX_THREADS_PER_BLOCK``.
|
||||
|
||||
If no ``__launch_bounds__`` are specified, ``MAX_THREADS_PER_BLOCK`` is
|
||||
the maximum block size supported by the device (see
|
||||
:doc:`../reference/hardware_features`). Reducing ``MAX_THREADS_PER_BLOCK``
|
||||
allows the compiler to use more resources per thread than an unconstrained
|
||||
compilation. This might however reduce the amount of blocks that can run
|
||||
concurrently on a CU, thereby reducing occupancy and trading thread-level
|
||||
parallelism for instruction-level parallelism.
|
||||
|
||||
``MAX_THREADS_PER_BLOCK`` is particularly useful in cases, where the compiler is
|
||||
constrained by register usage in order to meet requirements of large block sizes
|
||||
that are never used at launch time.
|
||||
|
||||
The compiler can only use the hints to manage register usage, and does not
|
||||
automatically reduce shared memory usage. The compilation fails, if the compiler
|
||||
can not generate code that satisfies the launch bounds.
|
||||
|
||||
On NVCC this parameter maps to the ``.maxntid`` PTX directive.
|
||||
|
||||
When launching kernels HIP will validate the launch configuration to make sure
|
||||
the requested block size is not larger than ``MAX_THREADS_PER_BLOCK`` and
|
||||
return an error if it is exceeded.
|
||||
|
||||
If :doc:`AMD_LOG_LEVEL <./logging>` is set, detailed information will be shown
|
||||
in the error log message, including the launch configuration of the kernel and
|
||||
the specified ``__launch_bounds__``.
|
||||
|
||||
MIN_WARPS_PER_EXECUTION_UNIT
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
|
||||
This parameter specifies the minimum number of warps that must be able to run
|
||||
concurrently on an execution unit.
|
||||
``MIN_WARPS_PER_EXECUTION_UNIT`` is optional and defaults to 1 if not specified.
|
||||
Since active warps compete for the same fixed pool of resources, the compiler
|
||||
must constrain the resource usage of the warps. This option gives a lower
|
||||
bound to the occupancy of the kernel.
|
||||
|
||||
From this parameter, the compiler derives a maximum number of registers that can
|
||||
be used in the kernel. The amount of registers that can be used at most is
|
||||
:math:`\frac{\text{available registers}}{\text{MIN_WARPS_PER_EXECUTION_UNIT}}`,
|
||||
but it might also have other, architecture specific, restrictions.
|
||||
|
||||
The available registers per Compute Unit are listed in
|
||||
:doc:`rocm:reference/gpu-arch-specs`. Beware that these values are per Compute
|
||||
Unit, not per Execution Unit. On AMD GPUs a Compute Unit consists of 4 Execution
|
||||
Units, also known as SIMDs, each with their own register file. For more
|
||||
information see :doc:`../understand/hardware_implementation`.
|
||||
:cpp:struct:`hipDeviceProp_t` also has a field ``executionUnitsPerMultiprocessor``.
|
||||
|
||||
Porting from CUDA __launch_bounds__
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
|
||||
CUDA also defines a ``__launch_bounds__`` qualifier which works similar to HIP's
|
||||
implementation, however it uses different parameters:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MULTIPROCESSOR)
|
||||
|
||||
The first parameter is the same as HIP's implementation, but
|
||||
``MIN_BLOCKS_PER_MULTIPROCESSOR`` must be converted to
|
||||
``MIN_WARPS_PER_EXECUTION``, which uses warps and execution units rather than
|
||||
blocks and multiprocessors. This conversion is performed automatically by
|
||||
:doc:`HIPIFY <hipify:index>`, or can be done manually with the following
|
||||
equation.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
MIN_WARPS_PER_EXECUTION_UNIT = (MIN_BLOCKS_PER_MULTIPROCESSOR * MAX_THREADS_PER_BLOCK) / warpSize
|
||||
|
||||
Directly controlling the warps per execution unit makes it easier to reason
|
||||
about the occupancy, unlike with blocks, where the occupancy depends on the
|
||||
block size.
|
||||
|
||||
The use of execution units rather than multiprocessors also provides support for
|
||||
architectures with multiple execution units per multiprocessor. For example, the
|
||||
AMD GCN architecture has 4 execution units per multiprocessor.
|
||||
|
||||
maxregcount
|
||||
""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""""
|
||||
|
||||
Unlike ``nvcc``, ``amdclang++`` does not support the ``--maxregcount`` option.
|
||||
Instead, users are encouraged to use the ``__launch_bounds__`` directive since
|
||||
the parameters are more intuitive and portable than micro-architecture details
|
||||
like registers. The directive allows per-kernel control.
|
||||
|
||||
Memory space qualifiers
|
||||
================================================================================
|
||||
|
||||
HIP adds qualifiers to specify the memory space in which the variables are
|
||||
located.
|
||||
|
||||
Generally, variables allocated in host memory are not directly accessible within
|
||||
device code, while variables allocated in device memory are not directly
|
||||
accessible from the host code. More details on this can be found in
|
||||
:ref:`unified_memory`.
|
||||
|
||||
__device__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Variables marked with ``__device__`` reside in device memory. It can be
|
||||
combined together with one of the following qualifiers, however these qualifiers
|
||||
also imply the ``__device__`` qualifier.
|
||||
|
||||
By default it can only be accessed from the threads on the device. In order to
|
||||
access it from the host, its address and size need to be queried using
|
||||
:cpp:func:`hipGetSymbolAddress` and :cpp:func:`hipGetSymbolSize` and copied with
|
||||
:cpp:func:`hipMemcpyToSymbol` or :cpp:func:`hipMemcpyFromSymbol`.
|
||||
|
||||
__constant__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Variables marked with ``__constant__`` reside in device memory. Variables in
|
||||
that address space are routed through the constant cache, but that address space
|
||||
has a limited logical size.
|
||||
This memory space is read-only from within kernels and can only be set by the
|
||||
host before kernel execution.
|
||||
|
||||
To get the best performance benefit, these variables need a special access
|
||||
pattern to benefit from the constant cache - the access has to be uniform within
|
||||
a warp, otherwise the accesses are serialized.
|
||||
|
||||
The constant cache reduces the pressure on the other caches and may enable
|
||||
higher throughput and lower latency accesses.
|
||||
|
||||
To set the ``__constant__`` variables the host must copy the data to the device
|
||||
using :cpp:func:`hipMemcpyToSymbol`, for example:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__constant__ int const_array[8];
|
||||
|
||||
void set_constant_memory(){
|
||||
int host_data[8] {1,2,3,4,5,6,7,8};
|
||||
|
||||
hipMemcpyToSymbol(const_array, host_data, sizeof(int) * 8);
|
||||
|
||||
// call kernel that accesses const_array
|
||||
}
|
||||
|
||||
__shared__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Variables marked with ``__shared__`` are only accessible by threads within the
|
||||
same block and have the lifetime of that block. It is usually backed by on-chip
|
||||
shared memory, providing fast access to all threads within a block, which makes
|
||||
it perfectly suited for sharing variables.
|
||||
|
||||
Shared memory can be allocated statically within the kernel, but the size
|
||||
of it has to be known at compile time.
|
||||
|
||||
In order to dynamically allocate shared memory during runtime, but before the
|
||||
kernel is launched, the variable has to be declared ``extern``, and the kernel
|
||||
launch has to specify the needed amount of ``extern`` shared memory in the launch
|
||||
configuration. The statically allocated shared memory is allocated without this
|
||||
parameter.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
extern __shared__ int shared_array[];
|
||||
|
||||
__global__ void kernel(){
|
||||
// initialize shared memory
|
||||
shared_array[threadIdx.x] = threadIdx.x;
|
||||
// use shared memory - synchronize to make sure, that all threads of the
|
||||
// block see all changes to shared memory
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
int main(){
|
||||
//shared memory in this case depends on the configurable block size
|
||||
constexpr int blockSize = 256;
|
||||
constexpr int sharedMemSize = blockSize * sizeof(int);
|
||||
constexpr int gridSize = 2;
|
||||
|
||||
kernel<<<dim3(gridSize), dim3(blockSize), sharedMemSize, 0>>>();
|
||||
}
|
||||
|
||||
__managed__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Managed memory is a special qualifier, that makes the marked memory available on
|
||||
the device and on the host. For more details see :ref:`unified_memory`.
|
||||
|
||||
__restrict__
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The ``__restrict__`` keyword tells the compiler that the associated memory
|
||||
pointer does not alias with any other pointer in the function. This can help the
|
||||
compiler perform better optimizations. For best results, every pointer passed to
|
||||
a function should use this keyword.
|
||||
|
||||
********************************************************************************
|
||||
Built-in constants
|
||||
********************************************************************************
|
||||
|
||||
HIP defines some special built-in constants for use in device code.
|
||||
|
||||
These built-ins are not implicitly defined by the compiler, the
|
||||
``hip_runtime.h`` header has to be included instead.
|
||||
|
||||
Index built-ins
|
||||
================================================================================
|
||||
|
||||
Kernel code can use these identifiers to distinguish between the different
|
||||
threads and blocks within a kernel.
|
||||
|
||||
These built-ins are of type dim3, and are constant for each thread, but differ
|
||||
between the threads or blocks, and are initialized at kernel launch.
|
||||
|
||||
blockDim and gridDim
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
``blockDim`` and ``gridDim`` contain the sizes specified at kernel launch.
|
||||
``blockDim`` contains the amount of threads in the x-, y- and z-dimensions of
|
||||
the block of threads. Similarly ``gridDim`` contains the amount of blocks in the
|
||||
grid.
|
||||
|
||||
.. _thread_and_block_idx:
|
||||
|
||||
threadIdx and blockIdx
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
``threadIdx`` and ``blockIdx`` can be used to identify the threads and blocks
|
||||
within the kernel.
|
||||
|
||||
``threadIdx`` identifies the thread within a block, meaning its values are
|
||||
within ``0`` and ``blockDim.{x,y,z} - 1``. Likewise ``blockIdx`` identifies the
|
||||
block within the grid, and the values are within ``0`` and ``gridDim.{} - 1``.
|
||||
|
||||
A global unique identifier of a three-dimensional grid can be calculated using
|
||||
the following code:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
(threadIdx.x + blockIdx.x * blockDim.x) +
|
||||
(threadIdx.y + blockIdx.y * blockDim.y) * blockDim.x +
|
||||
(threadIdx.z + blockIdx.z * blockDim.z) * blockDim.x * blockDim.y
|
||||
|
||||
.. _warp_size:
|
||||
|
||||
warpSize
|
||||
================================================================================
|
||||
|
||||
The ``warpSize`` constant contains the number of threads per warp for the given
|
||||
target device. It can differ between different architectures, and on RDNA
|
||||
architectures it can even differ between kernel launches, depending on whether
|
||||
they run in CU or WGP mode. See the
|
||||
:doc:`hardware features <../reference/hardware_features>` for more
|
||||
information.
|
||||
|
||||
Since ``warpSize`` can differ between devices, it can not be assumed to be a
|
||||
compile-time constant on the host. It has to be queried using
|
||||
:cpp:func:`hipDeviceGetAttribute` or :cpp:func:`hipDeviceGetProperties`, e.g.:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
int val;
|
||||
hipDeviceGetAttribute(&val, hipDeviceAttributeWarpSize, deviceId);
|
||||
|
||||
.. note::
|
||||
|
||||
``warpSize`` should not be assumed to be a specific value in portable HIP
|
||||
applications. NVIDIA devices return 32 for this variable; AMD devices return
|
||||
64 for gfx9 and 32 for gfx10 and above. While code that assumes a ``warpSize``
|
||||
of 32 can run on devices with a ``warpSize`` of 64, it only utilizes half of
|
||||
the the compute resources.
|
||||
|
||||
********************************************************************************
|
||||
Vector types
|
||||
********************************************************************************
|
||||
|
||||
These types are not automatically provided by the compiler. The
|
||||
``hip_vector_types.h`` header, which is also included by ``hip_runtime.h`` has
|
||||
to be included to use these types.
|
||||
|
||||
Fundamental vector types
|
||||
================================================================================
|
||||
|
||||
Fundamental vector types derive from the `fundamental C++ integral and
|
||||
floating-point types <https://en.cppreference.com/w/cpp/language/types>`_. These
|
||||
types are defined in ``hip_vector_types.h``, which is included by
|
||||
``hip_runtime.h``.
|
||||
|
||||
All vector types can be created with ``1``, ``2``, ``3`` or ``4`` elements, the
|
||||
corresponding type is ``<fundamental_type>i``, where ``i`` is the number of
|
||||
elements.
|
||||
|
||||
All vector types support a constructor function of the form
|
||||
``make_<type_name>()``. For example,
|
||||
``float3 make_float3(float x, float y, float z)`` creates a vector of type
|
||||
``float3`` with value ``(x,y,z)``.
|
||||
The elements of the vectors can be accessed using their members ``x``, ``y``,
|
||||
``z``, and ``w``.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
double2 d2_vec = make_double2(2.0, 4.0);
|
||||
double first_elem = d2_vec.x;
|
||||
|
||||
HIP supports vectors created from the following fundamental types:
|
||||
|
||||
.. list-table::
|
||||
:widths: 50 50
|
||||
|
||||
*
|
||||
- **Integral Types**
|
||||
-
|
||||
*
|
||||
- ``char``
|
||||
- ``uchar``
|
||||
*
|
||||
- ``short``
|
||||
- ``ushort``
|
||||
*
|
||||
- ``int``
|
||||
- ``uint``
|
||||
*
|
||||
- ``long``
|
||||
- ``ulong``
|
||||
*
|
||||
- ``longlong``
|
||||
- ``ulonglong``
|
||||
*
|
||||
- **Floating-Point Types**
|
||||
-
|
||||
*
|
||||
- ``float``
|
||||
-
|
||||
*
|
||||
- ``double``
|
||||
-
|
||||
|
||||
.. _dim3:
|
||||
|
||||
dim3
|
||||
================================================================================
|
||||
|
||||
``dim3`` is a special three-dimensional unsigned integer vector type that is
|
||||
commonly used to specify grid and group dimensions for kernel launch
|
||||
configurations.
|
||||
|
||||
Its constructor accepts up to three arguments. The unspecified dimensions are
|
||||
initialized to 1.
|
||||
|
||||
********************************************************************************
|
||||
Built-in device functions
|
||||
********************************************************************************
|
||||
|
||||
.. _memory_fence_instructions:
|
||||
|
||||
Memory fence instructions
|
||||
================================================================================
|
||||
|
||||
HIP does not enforce strict ordering on memory operations, meaning, that the
|
||||
order in which memory accesses are executed, is not necessarily the order in
|
||||
which other threads observe these changes. So it can not be assumed, that data
|
||||
written by one thread is visible by another thread without synchronization.
|
||||
|
||||
Memory fences are a way to enforce a sequentially consistent order on the memory
|
||||
operations. This means, that all writes to memory made before a memory fence are
|
||||
observed by all threads after the fence. The scope of these fences depends on
|
||||
what specific memory fence is called.
|
||||
|
||||
HIP supports ``__threadfence()``, ``__threadfence_block()`` and
|
||||
``__threadfence_system()``:
|
||||
|
||||
* ``__threadfence_block()`` orders memory accesses for all threads within a thread block.
|
||||
* ``__threadfence()`` orders memory accesses for all threads on a device.
|
||||
* ``__threadfence_system()`` orders memory accesses for all threads in the system, making writes to memory visible to other devices and the host
|
||||
|
||||
.. _synchronization_functions:
|
||||
|
||||
Synchronization functions
|
||||
================================================================================
|
||||
|
||||
Synchronization functions cause all threads in a group to wait at this
|
||||
synchronization point until all threads reached it. These functions implicitly
|
||||
include a :ref:`threadfence <memory_fence_instructions>`, thereby ensuring
|
||||
visibility of memory accesses for the threads in the group.
|
||||
|
||||
The ``__syncthreads()`` function comes in different versions.
|
||||
|
||||
``void __syncthreads()`` simply synchronizes the threads of a block. The other
|
||||
versions additionally evaluate a predicate:
|
||||
|
||||
``int __syncthreads_count(int predicate)`` returns the number of threads for
|
||||
which the predicate evaluates to non-zero.
|
||||
|
||||
``int __syncthreads_and(int predicate)`` returns non-zero if the predicate
|
||||
evaluates to non-zero for all threads.
|
||||
|
||||
``int __syncthreads_or(int predicate)`` returns non-zero if any of the
|
||||
predicates evaluates to non-zero.
|
||||
|
||||
The Cooperative Groups API offers options to synchronize threads on a developer
|
||||
defined set of thread groups. For further information, check the
|
||||
:ref:`Cooperative Groups API reference <cooperative_groups_reference>` or the
|
||||
:ref:`Cooperative Groups section in the programming guide
|
||||
<cooperative_groups_how-to>`.
|
||||
|
||||
Math functions
|
||||
================================================================================
|
||||
|
||||
HIP-Clang supports a set of math operations that are callable from the device.
|
||||
HIP supports most of the device functions supported by CUDA. These are described
|
||||
on :ref:`Math API page <math_api_reference>`.
|
||||
|
||||
Texture functions
|
||||
================================================================================
|
||||
|
||||
The supported texture functions are listed in ``texture_fetch_functions.h`` and
|
||||
``texture_indirect_functions.h`` header files in the
|
||||
`HIP-AMD backend repository <https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail>`_.
|
||||
|
||||
Texture functions are not supported on some devices. To determine if texture functions are supported
|
||||
on your device, use ``Macro __HIP_NO_IMAGE_SUPPORT == 1``. You can query the attribute
|
||||
``hipDeviceAttributeImageSupport`` to check if texture functions are supported in the host runtime
|
||||
code.
|
||||
|
||||
Surface functions
|
||||
================================================================================
|
||||
|
||||
The supported surface functions are located on :ref:`Surface object reference
|
||||
page <surface_object_reference>`.
|
||||
|
||||
Timer functions
|
||||
================================================================================
|
||||
|
||||
HIP provides device functions to read a high-resolution timer from within the
|
||||
kernel.
|
||||
|
||||
The following functions count the cycles on the device, where the rate varies
|
||||
with the actual frequency.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
clock_t clock()
|
||||
long long int clock64()
|
||||
|
||||
.. note::
|
||||
|
||||
``clock()`` and ``clock64()`` do not work properly on AMD RDNA3 (GFX11) graphic processors.
|
||||
|
||||
The difference between the returned values represents the cycles used.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global void kernel(){
|
||||
long long int start = clock64();
|
||||
// kernel code
|
||||
long long int stop = clock64();
|
||||
long long int cycles = stop - start;
|
||||
}
|
||||
|
||||
``long long int wall_clock64()`` returns the wall clock time on the device, with a constant, fixed frequency.
|
||||
The frequency is device dependent and can be queried using:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
int wallClkRate = 0; //in kilohertz
|
||||
hipDeviceGetAttribute(&wallClkRate, hipDeviceAttributeWallClockRate, deviceId);
|
||||
|
||||
.. _atomic functions:
|
||||
|
||||
Atomic functions
|
||||
================================================================================
|
||||
|
||||
Atomic functions are read-modify-write (RMW) operations, whose result is visible
|
||||
to all other threads on the scope of the atomic operation, once the operation
|
||||
completes.
|
||||
|
||||
If multiple instructions from different devices or threads target the same
|
||||
memory location, the instructions are serialized in an undefined order.
|
||||
|
||||
Atomic operations in kernels can operate on block scope (i.e. shared memory),
|
||||
device scope (global memory), or system scope (system memory), depending on
|
||||
:doc:`hardware support <../reference/hardware_features>`.
|
||||
|
||||
The listed functions are also available with the ``_system`` (e.g.
|
||||
``atomicAdd_system``) suffix, operating on system scope, which includes host
|
||||
memory and other GPUs' memory. The functions without suffix operate on shared
|
||||
or global memory on the executing device, depending on the memory space of the
|
||||
variable.
|
||||
|
||||
HIP supports the following atomic operations, where ``TYPE`` is one of ``int``,
|
||||
``unsigned int``, ``unsigned long``, ``unsigned long long``, ``float`` or
|
||||
``double``, while ``INTEGER`` is ``int``, ``unsigned int``, ``unsigned long``,
|
||||
``unsigned long long``:
|
||||
|
||||
.. list-table:: Atomic operations
|
||||
|
||||
* - ``TYPE atomicAdd(TYPE* address, TYPE val)``
|
||||
|
||||
* - ``TYPE atomicSub(TYPE* address, TYPE val)``
|
||||
|
||||
* - ``TYPE atomicMin(TYPE* address, TYPE val)``
|
||||
* - ``long long atomicMin(long long* address, long long val)``
|
||||
|
||||
* - ``TYPE atomicMax(TYPE* address, TYPE val)``
|
||||
* - ``long long atomicMax(long long* address, long long val)``
|
||||
|
||||
* - ``TYPE atomicExch(TYPE* address, TYPE val)``
|
||||
|
||||
* - ``TYPE atomicCAS(TYPE* address, TYPE compare, TYPE val)``
|
||||
|
||||
* - ``INTEGER atomicAnd(INTEGER* address, INTEGER val)``
|
||||
|
||||
* - ``INTEGER atomicOr(INTEGER* address, INTEGER val)``
|
||||
|
||||
* - ``INTEGER atomicXor(INTEGER* address, INTEGER val)``
|
||||
|
||||
* - ``unsigned int atomicInc(unsigned int* address)``
|
||||
|
||||
* - ``unsigned int atomicDec(unsigned int* address)``
|
||||
|
||||
Unsafe floating-point atomic operations
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Some HIP devices support fast atomic operations on floating-point values. For
|
||||
example, ``atomicAdd`` on single- or double-precision floating-point values may
|
||||
generate a hardware instruction that is faster than emulating the atomic
|
||||
operation using an atomic compare-and-swap (CAS) loop.
|
||||
|
||||
On some devices, fast atomic instructions can produce results that differ from
|
||||
the version implemented with atomic CAS loops. For example, some devices
|
||||
will use different rounding or denormal modes, and some devices produce
|
||||
incorrect answers if fast floating-point atomic instructions target fine-grained
|
||||
memory allocations.
|
||||
|
||||
The HIP-Clang compiler offers compile-time options to control the generation of
|
||||
unsafe atomic instructions. By default the compiler does not generate unsafe
|
||||
instructions. This is the same behaviour as with the ``-mno-unsafe-fp-atomics``
|
||||
compilation flag. The ``-munsafe-fp-atomics`` flag indicates to the compiler
|
||||
that all floating-point atomic function calls are allowed to use an unsafe
|
||||
version, if one exists. For example, on some devices, this flag indicates to the
|
||||
compiler that no floating-point ``atomicAdd`` function can target fine-grained
|
||||
memory. These options are applied globally for the entire compilation.
|
||||
|
||||
HIP provides special functions that override the global compiler option for safe
|
||||
or unsafe atomic functions.
|
||||
|
||||
The ``safe`` prefix always generates safe atomic operations, even when
|
||||
``-munsafe-fp-atomics`` is used, whereas ``unsafe`` always generates fast atomic
|
||||
instructions, even when ``-mno-unsafe-fp-atomics``. The following table lists
|
||||
the safe and unsafe atomic functions, where ``FLOAT_TYPE`` is either ``float``
|
||||
or ``double``.
|
||||
|
||||
.. list-table:: AMD specific atomic operations
|
||||
|
||||
* - ``FLOAT_TYPE unsafeAtomicAdd(FLOAT_TYPE* address, FLOAT_TYPE val)``
|
||||
|
||||
* - ``FLOAT_TYPE safeAtomicAdd(FLOAT_TYPE* address, FLOAT_TYPE val)``
|
||||
|
||||
.. _warp-cross-lane:
|
||||
|
||||
Warp cross-lane functions
|
||||
================================================================================
|
||||
|
||||
Threads in a warp are referred to as ``lanes`` and are numbered from ``0`` to
|
||||
``warpSize - 1``. Warp cross-lane functions cooperate across all lanes in a
|
||||
warp. AMD GPUs guarantee, that all warp lanes are executed in lockstep, whereas
|
||||
NVIDIA GPUs that support Independent Thread Scheduling might require additional
|
||||
synchronization, or the use of the ``__sync`` variants.
|
||||
|
||||
Note that different devices can have different warp sizes. You should query the
|
||||
:ref:`warpSize <warp_size>` in portable code and not assume a fixed warp size.
|
||||
|
||||
All mask values returned or accepted by these built-ins are 64-bit unsigned
|
||||
integer values, even when compiled for a device with 32 threads per warp. On
|
||||
such devices the higher bits are unused. CUDA code ported to HIP requires
|
||||
changes to ensure that the correct type is used.
|
||||
|
||||
Note that the ``__sync`` variants are made available in ROCm 6.2, but disabled by
|
||||
default to help with the transition to 64-bit masks. They can be enabled by
|
||||
setting the preprocessor macro ``HIP_ENABLE_WARP_SYNC_BUILTINS``. These built-ins
|
||||
will be enabled unconditionally in the next ROCm release. Wherever possible, the
|
||||
implementation includes a static assert to check that the program source uses
|
||||
the correct type for the mask.
|
||||
|
||||
The ``_sync`` variants require a 64-bit unsigned integer mask argument that
|
||||
specifies the lanes of the warp that will participate. Each participating thread
|
||||
must have its own bit set in its mask argument, and all active threads specified
|
||||
in any mask argument must execute the same call with the same mask, otherwise
|
||||
the result is undefined.
|
||||
|
||||
.. _warp_vote_functions:
|
||||
|
||||
Warp vote and ballot functions
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
int __all(int predicate)
|
||||
int __any(int predicate)
|
||||
unsigned long long __ballot(int predicate)
|
||||
unsigned long long __activemask()
|
||||
|
||||
int __all_sync(unsigned long long mask, int predicate)
|
||||
int __any_sync(unsigned long long mask, int predicate)
|
||||
unsigned long long __ballot_sync(unsigned long long mask, int predicate)
|
||||
|
||||
You can use ``__any`` and ``__all`` to get a summary view of the predicates evaluated by the
|
||||
participating lanes.
|
||||
|
||||
* ``__any()``: Returns 1 if the predicate is non-zero for any participating lane, otherwise it returns 0.
|
||||
|
||||
* ``__all()``: Returns 1 if the predicate is non-zero for all participating lanes, otherwise it returns 0.
|
||||
|
||||
To determine if the target platform supports the any/all instruction, you can
|
||||
query the ``hasWarpVote`` device property on the host or use the
|
||||
``HIP_ARCH_HAS_WARP_VOTE`` compiler definition in device code.
|
||||
|
||||
``__ballot`` returns a bit mask containing the 1-bit predicate value from each
|
||||
lane. The nth bit of the result contains the bit contributed by the nth lane.
|
||||
|
||||
``__activemask()`` returns a bit mask of currently active warp lanes. The nth
|
||||
bit of the result is 1 if the nth lane is active.
|
||||
|
||||
Note that the ``__ballot`` and ``__activemask`` built-ins in HIP have a 64-bit return
|
||||
value (unlike the 32-bit value returned by the CUDA built-ins). Code ported from
|
||||
CUDA should be adapted to support the larger warp sizes that the HIP version
|
||||
requires.
|
||||
|
||||
Applications can test whether the target platform supports the ``__ballot`` or
|
||||
``__activemask`` instructions using the ``hasWarpBallot`` device property in host
|
||||
code or the ``HIP_ARCH_HAS_WARP_BALLOT`` macro defined by the compiler for device
|
||||
code.
|
||||
|
||||
Warp match functions
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
unsigned long long __match_any(T value)
|
||||
unsigned long long __match_all(T value, int *pred)
|
||||
|
||||
unsigned long long __match_any_sync(unsigned long long mask, T value)
|
||||
unsigned long long __match_all_sync(unsigned long long mask, T value, int *pred)
|
||||
|
||||
``T`` can be a 32-bit integer type, 64-bit integer type or a single precision or
|
||||
double precision floating point type.
|
||||
|
||||
``__match_any`` returns a bit mask where the n-th bit is set to 1 if the n-th
|
||||
lane has the same ``value`` as the current lane, and 0 otherwise.
|
||||
|
||||
``__match_all`` returns a bit mask with the bits of the participating lanes are
|
||||
set to 1 if all lanes have the same ``value``, and 0 otherwise.
|
||||
The predicate ``pred`` is set to true if all participating threads have the same
|
||||
``value``, and false otherwise.
|
||||
|
||||
Warp shuffle functions
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
T __shfl (T var, int srcLane, int width=warpSize);
|
||||
T __shfl_up (T var, unsigned int delta, int width=warpSize);
|
||||
T __shfl_down (T var, unsigned int delta, int width=warpSize);
|
||||
T __shfl_xor (T var, int laneMask, int width=warpSize);
|
||||
|
||||
T __shfl_sync (unsigned long long mask, T var, int srcLane, int width=warpSize);
|
||||
T __shfl_up_sync (unsigned long long mask, T var, unsigned int delta, int width=warpSize);
|
||||
T __shfl_down_sync (unsigned long long mask, T var, unsigned int delta, int width=warpSize);
|
||||
T __shfl_xor_sync (unsigned long long mask, T var, int laneMask, int width=warpSize);
|
||||
|
||||
``T`` can be a 32-bit integer type, 64-bit integer type or a single precision or
|
||||
double precision floating point type.
|
||||
|
||||
The warp shuffle functions exchange values between threads within a warp.
|
||||
|
||||
The optional ``width`` argument specifies subgroups, in which the warp can be
|
||||
divided to share the variables.
|
||||
It has to be a power of two smaller than or equal to ``warpSize``. If it is
|
||||
smaller than ``warpSize``, the warp is grouped into separate groups, that are each
|
||||
indexed from 0 to width as if it was its own entity, and only the lanes within
|
||||
that subgroup participate in the shuffle. The lane indices in the subgroup are
|
||||
given by ``laneIdx % width``.
|
||||
|
||||
The different shuffle functions behave as following:
|
||||
|
||||
``__shfl``
|
||||
The thread reads the value from the lane specified in ``srcLane``.
|
||||
|
||||
``__shfl_up``
|
||||
The thread reads ``var`` from lane ``laneIdx - delta``, thereby "shuffling"
|
||||
the values of the lanes of the warp "up". If the resulting source lane is out
|
||||
of range, the thread returns its own ``var``.
|
||||
|
||||
``__shfl_down``
|
||||
The thread reads ``var`` from lane ``laneIdx - delta``, thereby "shuffling"
|
||||
the values of the lanes of the warp "down". If the resulting source lane is
|
||||
out of range, the thread returns its own ``var``.
|
||||
|
||||
``__shfl_xor``
|
||||
The thread reads ``var`` from lane ``laneIdx xor lane_mask``. If ``width`` is
|
||||
smaller than ``warpSize``, the threads can read values from subgroups before
|
||||
the current subgroup. If it tries to read values from later subgroups, the
|
||||
function returns the ``var`` of the calling thread.
|
||||
|
||||
Warp matrix functions
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Warp matrix functions allow a warp to cooperatively operate on small matrices
|
||||
that have elements spread over lanes in an unspecified manner.
|
||||
|
||||
HIP does not support warp matrix types or functions.
|
||||
|
||||
Cooperative groups functions
|
||||
================================================================================
|
||||
|
||||
You can use cooperative groups to synchronize groups of threads across thread
|
||||
blocks. It also provide a way of communicating between these groups.
|
||||
|
||||
For further information, check the :ref:`Cooperative Groups API reference
|
||||
<cooperative_groups_reference>` or the :ref:`Cooperative Groups programming
|
||||
guide <cooperative_groups_how-to>`.
|
||||
@@ -1,3 +1,9 @@
|
||||
<head>
|
||||
<meta charset="UTF-8">
|
||||
<meta name="description" content="HIP porting guide describing how to port CUDA code to HIP.">
|
||||
<meta name="keywords" content="HIP, Heterogeneous-computing Interface for Portability, HIP porting guide">
|
||||
</head>
|
||||
|
||||
# HIP porting guide
|
||||
|
||||
In addition to providing a portable C++ programming environment for GPUs, HIP is designed to ease
|
||||
@@ -373,7 +379,9 @@ run hipcc when appropriate.
|
||||
|
||||
### ``warpSize``
|
||||
|
||||
Code should not assume a warp size of 32 or 64. See [Warp Cross-Lane Functions](https://rocm.docs.amd.com/projects/HIP/en/latest/reference/cpp_language_extensions.html#warp-cross-lane-functions) for information on how to write portable wave-aware code.
|
||||
Code should not assume a warp size of 32 or 64. See the
|
||||
:ref:`HIP language extension for warpSize <warp_size>` for information on how
|
||||
to write portable wave-aware code.
|
||||
|
||||
### Kernel launch with group size > 256
|
||||
|
||||
|
||||
@@ -1,13 +1,21 @@
|
||||
<head>
|
||||
<meta charset="UTF-8">
|
||||
<meta name="description" content="HIP Runtime Compiler API.">
|
||||
<meta name="keywords" content="HIP, Heterogeneous-computing Interface for Portability, HIP runtime compiler">
|
||||
</head>
|
||||
|
||||
# Programming for HIP runtime compiler (RTC)
|
||||
|
||||
HIP lets you compile kernels at runtime with the `hiprtc*` APIs.
|
||||
Kernels can be stored as a text string and can be passed to HIPRTC APIs alongside options to guide the compilation.
|
||||
|
||||
NOTE:
|
||||
:::{note}
|
||||
|
||||
* This library can be used on systems without HIP installed nor AMD GPU driver installed at all (offline compilation). Therefore, it does not depend on any HIP runtime library.
|
||||
* But it does depend on Code Object Manager (comgr). You may try to statically link comgr into HIPRTC to avoid any ambiguity.
|
||||
* Developers can decide to bundle this library with their application.
|
||||
* This library can be used on systems without HIP installed nor AMD GPU driver installed at all (offline compilation). Therefore, it doesn't depend on any HIP runtime library.
|
||||
* This library depends on Code Object Manager (comgr). You can try to statically link comgr into HIPRTC to avoid ambiguity.
|
||||
* Developers can bundle this library with their application.
|
||||
|
||||
:::
|
||||
|
||||
## Compilation APIs
|
||||
|
||||
@@ -224,6 +232,22 @@ int main() {
|
||||
}
|
||||
```
|
||||
|
||||
## Kernel Compilation Cache
|
||||
|
||||
HIPRTC incorporates a cache to avoid recompiling kernels between program executions. The contents of the cache include the kernel source code (including the contents of any `#include` headers), the compilation flags, and the compiler version. After a ROCm version update, the kernels are progressively recompiled, and the new results are cached. When the cache is disabled, each kernel is recompiled every time it is requested.
|
||||
|
||||
Use the following environment variables to manage the cache status as enabled or disabled, the location for storing the cache contents, and the cache eviction policy:
|
||||
|
||||
* `AMD_COMGR_CACHE` By default this variable has a value of `0` and the compilation cache feature is disabled. To enable the feature set the environment variable to a value of `1` (or any value other than `0`). This behavior may change in a future release.
|
||||
|
||||
* `AMD_COMGR_CACHE_DIR`: By default the value of this environment variable is defined as `$XDG_CACHE_HOME/comgr_cache`, which defaults to `$USER/.cache/comgr_cache` on Linux, and `%LOCALAPPDATA%\cache\comgr_cache` on Windows. You can specify a different directory for the environment variable to change the path for cache storage. If the runtime fails to access the specified cache directory, or the environment variable is set to an empty string (""), the cache is disabled.
|
||||
|
||||
* `AMD_COMGR_CACHE_POLICY`: If assigned a value, the string is interpreted and applied to the cache pruning policy. The string format is consistent with [Clang's ThinLTO cache pruning policy](https://rocm.docs.amd.com/projects/llvm-project/en/latest/LLVM/clang/html/ThinLTO.html#cache-pruning). The default policy is defined as: `prune_interval=1h:prune_expiration=0h:cache_size=75%:cache_size_bytes=30g:cache_size_files=0`. If the runtime fails to parse the defined string, or the environment variable is set to an empty string (""), the cache is disabled.
|
||||
|
||||
:::{note}
|
||||
This cache is also shared with the OpenCL runtime shipped with ROCm.
|
||||
:::
|
||||
|
||||
## HIPRTC specific options
|
||||
|
||||
HIPRTC provides a few HIPRTC specific flags
|
||||
|
||||
@@ -40,6 +40,7 @@ Here are the various HIP Runtime API high level functions:
|
||||
* :doc:`./hip_runtime_api/initialization`
|
||||
* :doc:`./hip_runtime_api/memory_management`
|
||||
* :doc:`./hip_runtime_api/error_handling`
|
||||
* :doc:`./hip_runtime_api/asynchronous`
|
||||
* :doc:`./hip_runtime_api/cooperative_groups`
|
||||
* :doc:`./hip_runtime_api/hipgraph`
|
||||
* :doc:`./hip_runtime_api/call_stack`
|
||||
|
||||
@@ -0,0 +1,534 @@
|
||||
.. meta::
|
||||
:description: This topic describes asynchronous concurrent execution in HIP
|
||||
:keywords: AMD, ROCm, HIP, asynchronous concurrent execution, asynchronous, async, concurrent, concurrency
|
||||
|
||||
.. _asynchronous_how-to:
|
||||
|
||||
*******************************************************************************
|
||||
Asynchronous concurrent execution
|
||||
*******************************************************************************
|
||||
|
||||
Asynchronous concurrent execution is important for efficient parallelism and
|
||||
resource utilization, with techniques such as overlapping computation and data
|
||||
transfer, managing concurrent kernel execution with streams on single or
|
||||
multiple devices, or using HIP graphs.
|
||||
|
||||
Streams and concurrent execution
|
||||
===============================================================================
|
||||
|
||||
All asynchronous APIs, such as kernel execution, data movement and potentially
|
||||
data allocation/freeing all happen in the context of device streams.
|
||||
|
||||
Streams are FIFO buffers of commands to execute in order on a given device.
|
||||
Commands which enqueue tasks on a stream all return promptly and the task is
|
||||
executed asynchronously. Multiple streams can point to the same device and
|
||||
those streams might be fed from multiple concurrent host-side threads. Multiple
|
||||
streams tied to the same device are not guaranteed to execute their commands in
|
||||
order.
|
||||
|
||||
Managing streams
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Streams enable the overlap of computation and data transfer, ensuring
|
||||
continuous GPU activity. By enabling tasks to run concurrently within the same
|
||||
GPU or across different GPUs, streams improve performance and throughput in
|
||||
high-performance computing (HPC).
|
||||
|
||||
To create a stream, the following functions are used, each defining a handle
|
||||
to the newly created stream:
|
||||
|
||||
- :cpp:func:`hipStreamCreate`: Creates a stream with default settings.
|
||||
- :cpp:func:`hipStreamCreateWithFlags`: Creates a stream, with specific
|
||||
flags, listed below, enabling more control over stream behavior:
|
||||
|
||||
- ``hipStreamDefault``: creates a default stream suitable for most
|
||||
operations. The default stream is a blocking operation.
|
||||
- ``hipStreamNonBlocking``: creates a non-blocking stream, allowing
|
||||
concurrent execution of operations. It ensures that tasks can run
|
||||
simultaneously without waiting for each other to complete, thus improving
|
||||
overall performance.
|
||||
|
||||
- :cpp:func:`hipStreamCreateWithPriority`: Allows creating a stream with a
|
||||
specified priority, enabling prioritization of certain tasks.
|
||||
|
||||
The :cpp:func:`hipStreamSynchronize` function is used to block the calling host
|
||||
thread until all previously submitted tasks in a specified HIP stream have
|
||||
completed. It ensures that all operations in the given stream, such as kernel
|
||||
executions or memory transfers, are finished before the host thread proceeds.
|
||||
|
||||
.. note::
|
||||
|
||||
If the :cpp:func:`hipStreamSynchronize` function input stream is 0 (or the
|
||||
default stream), it waits for all operations in the default stream to
|
||||
complete.
|
||||
|
||||
Concurrent execution between host and device
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Concurrent execution between the host (CPU) and device (GPU) allows the CPU to
|
||||
perform other tasks while the GPU is executing kernels. Kernels are launched
|
||||
asynchronously using ``hipLaunchKernelGGL`` or using the triple chevron with a stream,
|
||||
enabling the CPU to continue executing other code while the GPU processes the
|
||||
kernel. Similarly, memory operations like :cpp:func:`hipMemcpyAsync` are
|
||||
performed asynchronously, allowing data transfers between the host and device
|
||||
without blocking the CPU.
|
||||
|
||||
Concurrent kernel execution
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Concurrent execution of multiple kernels on the GPU allows different kernels to
|
||||
run simultaneously to maximize GPU resource usage. Managing dependencies
|
||||
between kernels is crucial for ensuring correct execution order. This can be
|
||||
achieved using :cpp:func:`hipStreamWaitEvent`, which allows a kernel to wait
|
||||
for a specific event before starting execution.
|
||||
|
||||
Independent kernels can only run concurrently if there are enough registers
|
||||
and shared memory for the kernels. To enable concurrent kernel executions, the
|
||||
developer may have to reduce the block size of the kernels. The kernel runtimes
|
||||
can be misleading for concurrent kernel runs, that is why during optimization
|
||||
it is a good practice to check the trace files, to see if one kernel is blocking
|
||||
another kernel, while they are running in parallel. For more information about
|
||||
the application tracing, check::doc:`rocprofiler:/how-to/using-rocprof`.
|
||||
|
||||
When running kernels in parallel, the execution time can increase due to
|
||||
contention for shared resources. This is because multiple kernels may attempt
|
||||
to access the same GPU resources simultaneously, leading to delays.
|
||||
|
||||
Multiple kernels executing concurrently is only beneficial under specific conditions. It
|
||||
is most effective when the kernels do not fully utilize the GPU's resources. In
|
||||
such cases, overlapping kernel execution can improve overall throughput and
|
||||
efficiency by keeping the GPU busy without exceeding its capacity.
|
||||
|
||||
Overlap of data transfer and kernel execution
|
||||
===============================================================================
|
||||
|
||||
One of the primary benefits of asynchronous operations and multiple streams is
|
||||
the ability to overlap data transfer with kernel execution, leading to better
|
||||
resource utilization and improved performance.
|
||||
|
||||
Asynchronous execution is particularly advantageous in iterative processes. For
|
||||
instance, if a kernel is initiated, it can be efficient to prepare the input
|
||||
data simultaneously, provided that this preparation does not depend on the
|
||||
kernel's execution. Such iterative data transfer and kernel execution overlap
|
||||
can be find in the :ref:`async_example`.
|
||||
|
||||
Querying device capabilities
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Some AMD HIP-enabled devices can perform asynchronous memory copy operations to
|
||||
or from the GPU concurrently with kernel execution. Applications can query this
|
||||
capability by checking the ``asyncEngineCount`` device property. Devices with
|
||||
an ``asyncEngineCount`` greater than zero support concurrent data transfers.
|
||||
Additionally, if host memory is involved in the copy, it should be page-locked
|
||||
to ensure optimal performance. Page-locking (or pinning) host memory increases
|
||||
the bandwidth between the host and the device, reducing the overhead associated
|
||||
with data transfers. For more details, visit :ref:`host_memory` page.
|
||||
|
||||
Asynchronous memory operations
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Asynchronous memory operations do not block the host while copying data and,
|
||||
when used with multiple streams, allow data to be transferred between the host
|
||||
and device while kernels are executed on the same GPU. Using operations like
|
||||
:cpp:func:`hipMemcpyAsync` or :cpp:func:`hipMemcpyPeerAsync`, developers can
|
||||
initiate data transfers without waiting for the previous operation to complete.
|
||||
This overlap of computation and data transfer ensures that the GPU is not idle
|
||||
while waiting for data. :cpp:func:`hipMemcpyPeerAsync` enables data transfers
|
||||
between different GPUs, facilitating multi-GPU communication.
|
||||
|
||||
:ref:`async_example`` include launching kernels in one stream while performing
|
||||
data transfers in another. This technique is especially useful in applications
|
||||
with large data sets that need to be processed quickly.
|
||||
|
||||
Concurrent data transfers with intra-device copies
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Devices that support the ``concurrentKernels`` property can perform
|
||||
intra-device copies concurrently with kernel execution. Additionally, devices
|
||||
that support the ``asyncEngineCount`` property can perform data transfers to
|
||||
or from the GPU simultaneously with kernel execution. Intra-device copies can
|
||||
be initiated using standard memory copy functions with destination and source
|
||||
addresses residing on the same device.
|
||||
|
||||
Synchronization, event management and synchronous calls
|
||||
===============================================================================
|
||||
|
||||
Synchronization and event management are important for coordinating tasks and
|
||||
ensuring correct execution order, and synchronous calls are necessary for
|
||||
maintaining data consistency.
|
||||
|
||||
Synchronous calls
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Synchronous calls ensure task completion before moving to the next operation.
|
||||
For example, :cpp:func:`hipMemcpy` for data transfers waits for completion
|
||||
before returning control to the host. Similarly, synchronous kernel launches
|
||||
are used when immediate completion is required. When a synchronous function is
|
||||
called, control is not returned to the host thread before the device has
|
||||
completed the requested task. The behavior of the host thread—whether to yield,
|
||||
block, or spin—can be specified using :cpp:func:`hipSetDeviceFlags` with
|
||||
appropriate flags. Understanding when to use synchronous calls is important for
|
||||
managing execution flow and avoiding data races.
|
||||
|
||||
Events for synchronization
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
By creating an event with :cpp:func:`hipEventCreate` and recording it with
|
||||
:cpp:func:`hipEventRecord`, developers can synchronize operations across
|
||||
streams, ensuring correct task execution order. :cpp:func:`hipEventSynchronize`
|
||||
lets the application wait for an event to complete before proceeding with the next
|
||||
operation.
|
||||
|
||||
Programmatic dependent launch and synchronization
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
While CUDA supports programmatic dependent launches allowing a secondary kernel
|
||||
to start before the primary kernel finishes, HIP achieves similar functionality
|
||||
using streams and events. By employing :cpp:func:`hipStreamWaitEvent`, it is
|
||||
possible to manage the execution order without explicit hardware support. This
|
||||
mechanism allows a secondary kernel to launch as soon as the necessary
|
||||
conditions are met, even if the primary kernel is still running.
|
||||
|
||||
.. _async_example:
|
||||
|
||||
Example
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
The examples shows the difference between sequential, asynchronous calls and
|
||||
asynchronous calls with ``hipEvents``.
|
||||
|
||||
.. figure:: ../../data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg
|
||||
:alt: Compare the different calls
|
||||
:align: center
|
||||
|
||||
The example codes
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Sequential
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
// GPU Kernels
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += 1.0;}
|
||||
};
|
||||
__global__ void kernelB(double* arrayA, double* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] += arrayA[x] + 3.0;}
|
||||
};
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int numOfBlocks = 1 << 20;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr int numberOfIterations = 50;
|
||||
// The array size smaller to avoid the relatively short kernel launch compared to memory copies
|
||||
constexpr size_t arraySize = 1U << 25;
|
||||
double *d_dataA;
|
||||
double *d_dataB;
|
||||
|
||||
double initValueA = 0.0;
|
||||
double initValueB = 2.0;
|
||||
|
||||
std::vector<double> vectorA(arraySize, initValueA);
|
||||
std::vector<double> vectorB(arraySize, initValueB);
|
||||
// Allocate device memory
|
||||
HIP_CHECK(hipMalloc(&d_dataA, arraySize * sizeof(*d_dataA)));
|
||||
HIP_CHECK(hipMalloc(&d_dataB, arraySize * sizeof(*d_dataB)));
|
||||
for(int iteration = 0; iteration < numberOfIterations; iteration++)
|
||||
{
|
||||
// Host to Device copies
|
||||
HIP_CHECK(hipMemcpy(d_dataA, vectorA.data(), arraySize * sizeof(*d_dataA), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(d_dataB, vectorB.data(), arraySize * sizeof(*d_dataB), hipMemcpyHostToDevice));
|
||||
// Launch the GPU kernels
|
||||
hipLaunchKernelGGL(kernelA, dim3(numOfBlocks), dim3(threadsPerBlock), 0, 0, d_dataA, arraySize);
|
||||
hipLaunchKernelGGL(kernelB, dim3(numOfBlocks), dim3(threadsPerBlock), 0, 0, d_dataA, d_dataB, arraySize);
|
||||
// Device to Host copies
|
||||
HIP_CHECK(hipMemcpy(vectorA.data(), d_dataA, arraySize * sizeof(*vectorA.data()), hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipMemcpy(vectorB.data(), d_dataB, arraySize * sizeof(*vectorB.data()), hipMemcpyDeviceToHost));
|
||||
}
|
||||
// Wait for all operations to complete
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Verify results
|
||||
const double expectedA = (double)numberOfIterations;
|
||||
const double expectedB =
|
||||
initValueB + (3.0 * numberOfIterations) +
|
||||
(expectedA * (expectedA + 1.0)) / 2.0;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(vectorA[i] != expectedA){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedA << " got " << vectorA[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
if(vectorB[i] != expectedB){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedB << " got " << vectorB[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if(passed){
|
||||
std::cout << "Sequential execution completed successfully." << std::endl;
|
||||
}else{
|
||||
std::cerr << "Sequential execution failed." << std::endl;
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipFree(d_dataA));
|
||||
HIP_CHECK(hipFree(d_dataB));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: Asynchronous
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
// GPU Kernels
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += 1.0;}
|
||||
};
|
||||
__global__ void kernelB(double* arrayA, double* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] += arrayA[x] + 3.0;}
|
||||
};
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int numOfBlocks = 1 << 20;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr int numberOfIterations = 50;
|
||||
// The array size smaller to avoid the relatively short kernel launch compared to memory copies
|
||||
constexpr size_t arraySize = 1U << 25;
|
||||
double *d_dataA;
|
||||
double *d_dataB;
|
||||
|
||||
double initValueA = 0.0;
|
||||
double initValueB = 2.0;
|
||||
|
||||
std::vector<double> vectorA(arraySize, initValueA);
|
||||
std::vector<double> vectorB(arraySize, initValueB);
|
||||
// Allocate device memory
|
||||
HIP_CHECK(hipMalloc(&d_dataA, arraySize * sizeof(*d_dataA)));
|
||||
HIP_CHECK(hipMalloc(&d_dataB, arraySize * sizeof(*d_dataB)));
|
||||
// Create streams
|
||||
hipStream_t streamA, streamB;
|
||||
HIP_CHECK(hipStreamCreate(&streamA));
|
||||
HIP_CHECK(hipStreamCreate(&streamB));
|
||||
for(unsigned int iteration = 0; iteration < numberOfIterations; iteration++)
|
||||
{
|
||||
// Stream 1: Host to Device 1
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataA, vectorA.data(), arraySize * sizeof(*d_dataA), hipMemcpyHostToDevice, streamA));
|
||||
// Stream 2: Host to Device 2
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataB, vectorB.data(), arraySize * sizeof(*d_dataB), hipMemcpyHostToDevice, streamB));
|
||||
// Stream 1: Kernel 1
|
||||
hipLaunchKernelGGL(kernelA, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamA, d_dataA, arraySize);
|
||||
// Wait for streamA finish
|
||||
HIP_CHECK(hipStreamSynchronize(streamA));
|
||||
// Stream 2: Kernel 2
|
||||
hipLaunchKernelGGL(kernelB, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamB, d_dataA, d_dataB, arraySize);
|
||||
// Stream 1: Device to Host 2 (after Kernel 1)
|
||||
HIP_CHECK(hipMemcpyAsync(vectorA.data(), d_dataA, arraySize * sizeof(*vectorA.data()), hipMemcpyDeviceToHost, streamA));
|
||||
// Stream 2: Device to Host 2 (after Kernel 2)
|
||||
HIP_CHECK(hipMemcpyAsync(vectorB.data(), d_dataB, arraySize * sizeof(*vectorB.data()), hipMemcpyDeviceToHost, streamB));
|
||||
}
|
||||
// Wait for all operations in both streams to complete
|
||||
HIP_CHECK(hipStreamSynchronize(streamA));
|
||||
HIP_CHECK(hipStreamSynchronize(streamB));
|
||||
// Verify results
|
||||
double expectedA = (double)numberOfIterations;
|
||||
double expectedB =
|
||||
initValueB + (3.0 * numberOfIterations) +
|
||||
(expectedA * (expectedA + 1.0)) / 2.0;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(vectorA[i] != expectedA){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedA << " got " << vectorA[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
if(vectorB[i] != expectedB){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedB << " got " << vectorB[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if(passed){
|
||||
std::cout << "Asynchronous execution completed successfully." << std::endl;
|
||||
}else{
|
||||
std::cerr << "Asynchronous execution failed." << std::endl;
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipStreamDestroy(streamA));
|
||||
HIP_CHECK(hipStreamDestroy(streamB));
|
||||
HIP_CHECK(hipFree(d_dataA));
|
||||
HIP_CHECK(hipFree(d_dataB));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: hipStreamWaitEvent
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
// GPU Kernels
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += 1.0;}
|
||||
};
|
||||
__global__ void kernelB(double* arrayA, double* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] += arrayA[x] + 3.0;}
|
||||
};
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int numOfBlocks = 1 << 20;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr int numberOfIterations = 50;
|
||||
// The array size smaller to avoid the relatively short kernel launch compared to memory copies
|
||||
constexpr size_t arraySize = 1U << 25;
|
||||
double *d_dataA;
|
||||
double *d_dataB;
|
||||
double initValueA = 0.0;
|
||||
double initValueB = 2.0;
|
||||
|
||||
std::vector<double> vectorA(arraySize, initValueA);
|
||||
std::vector<double> vectorB(arraySize, initValueB);
|
||||
// Allocate device memory
|
||||
HIP_CHECK(hipMalloc(&d_dataA, arraySize * sizeof(*d_dataA)));
|
||||
HIP_CHECK(hipMalloc(&d_dataB, arraySize * sizeof(*d_dataB)));
|
||||
// Create streams
|
||||
hipStream_t streamA, streamB;
|
||||
HIP_CHECK(hipStreamCreate(&streamA));
|
||||
HIP_CHECK(hipStreamCreate(&streamB));
|
||||
// Create events
|
||||
hipEvent_t event, eventA, eventB;
|
||||
HIP_CHECK(hipEventCreate(&event));
|
||||
HIP_CHECK(hipEventCreate(&eventA));
|
||||
HIP_CHECK(hipEventCreate(&eventB));
|
||||
for(unsigned int iteration = 0; iteration < numberOfIterations; iteration++)
|
||||
{
|
||||
// Stream 1: Host to Device 1
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataA, vectorA.data(), arraySize * sizeof(*d_dataA), hipMemcpyHostToDevice, streamA));
|
||||
// Stream 2: Host to Device 2
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataB, vectorB.data(), arraySize * sizeof(*d_dataB), hipMemcpyHostToDevice, streamB));
|
||||
// Stream 1: Kernel 1
|
||||
hipLaunchKernelGGL(kernelA, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamA, d_dataA, arraySize);
|
||||
// Record event after the GPU kernel in Stream 1
|
||||
HIP_CHECK(hipEventRecord(event, streamA));
|
||||
// Stream 2: Wait for event before starting Kernel 2
|
||||
HIP_CHECK(hipStreamWaitEvent(streamB, event, 0));
|
||||
// Stream 2: Kernel 2
|
||||
hipLaunchKernelGGL(kernelB, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamB, d_dataA, d_dataB, arraySize);
|
||||
// Stream 1: Device to Host 2 (after Kernel 1)
|
||||
HIP_CHECK(hipMemcpyAsync(vectorA.data(), d_dataA, arraySize * sizeof(*vectorA.data()), hipMemcpyDeviceToHost, streamA));
|
||||
// Stream 2: Device to Host 2 (after Kernel 2)
|
||||
HIP_CHECK(hipMemcpyAsync(vectorB.data(), d_dataB, arraySize * sizeof(*vectorB.data()), hipMemcpyDeviceToHost, streamB));
|
||||
// Wait for all operations in both streams to complete
|
||||
HIP_CHECK(hipEventRecord(eventA, streamA));
|
||||
HIP_CHECK(hipEventRecord(eventB, streamB));
|
||||
HIP_CHECK(hipStreamWaitEvent(streamA, eventA, 0));
|
||||
HIP_CHECK(hipStreamWaitEvent(streamB, eventB, 0));
|
||||
}
|
||||
// Verify results
|
||||
double expectedA = (double)numberOfIterations;
|
||||
double expectedB =
|
||||
initValueB + (3.0 * numberOfIterations) +
|
||||
(expectedA * (expectedA + 1.0)) / 2.0;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(vectorA[i] != expectedA){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedA << " got " << vectorA[i] << std::endl;
|
||||
break;
|
||||
}
|
||||
if(vectorB[i] != expectedB){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedB << " got " << vectorB[i] << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if(passed){
|
||||
std::cout << "Asynchronous execution with events completed successfully." << std::endl;
|
||||
}else{
|
||||
std::cerr << "Asynchronous execution with events failed." << std::endl;
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipEventDestroy(event));
|
||||
HIP_CHECK(hipEventDestroy(eventA));
|
||||
HIP_CHECK(hipEventDestroy(eventB));
|
||||
HIP_CHECK(hipStreamDestroy(streamA));
|
||||
HIP_CHECK(hipStreamDestroy(streamB));
|
||||
HIP_CHECK(hipFree(d_dataA));
|
||||
HIP_CHECK(hipFree(d_dataB));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
HIP Graphs
|
||||
===============================================================================
|
||||
|
||||
HIP graphs offer an efficient alternative to the standard method of launching
|
||||
GPU tasks via streams. Comprising nodes for operations and edges for
|
||||
dependencies, HIP graphs reduce kernel launch overhead and provide a high-level
|
||||
abstraction for managing dependencies and synchronization. By representing
|
||||
sequences of kernels and memory operations as a single graph, they simplify
|
||||
complex workflows and enhance performance, particularly for applications with
|
||||
intricate dependencies and multiple execution stages.
|
||||
For more details, see the :ref:`how_to_HIP_graph` documentation.
|
||||
+259
-26
@@ -1,52 +1,285 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the device memory of the HIP ecosystem
|
||||
ROCm software.
|
||||
:keywords: AMD, ROCm, HIP, device memory
|
||||
:keywords: AMD, ROCm, HIP, GPU, device memory, global, constant, texture, surface, shared
|
||||
|
||||
.. _device_memory:
|
||||
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
Device memory
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
|
||||
Device memory exists on the device, e.g. on GPUs in the video random access
|
||||
memory (VRAM), and is accessible by the kernels operating on the device. Recent
|
||||
architectures use graphics double data rate (GDDR) synchronous dynamic
|
||||
random-access memory (SDRAM) such as GDDR6, or high-bandwidth memory (HBM) such
|
||||
as HBM2e. Device memory can be allocated as global memory, constant, texture or
|
||||
surface memory.
|
||||
Device memory is random access memory that is physically located on a GPU. In
|
||||
general it is memory with a bandwidth that is an order of magnitude higher
|
||||
compared to RAM available to the host. That high bandwidth is only available to
|
||||
on-device accesses, accesses from the host or other devices have to go over a
|
||||
special interface which is considerably slower, usually the PCIe bus or the AMD
|
||||
Infinity Fabric.
|
||||
|
||||
On certain architectures like APUs, the GPU and CPU share the same physical
|
||||
memory.
|
||||
|
||||
There is also a special local data share on-chip directly accessible to the
|
||||
:ref:`compute units <hardware_implementation>`, that can be used for shared
|
||||
memory.
|
||||
|
||||
The physical device memory can be used to back up several different memory
|
||||
spaces in HIP, as described in the following.
|
||||
|
||||
Global memory
|
||||
================================================================================
|
||||
|
||||
Read-write storage visible to all threads on a given device. There are
|
||||
specialized versions of global memory with different usage semantics which are
|
||||
typically backed by the same hardware, but can use different caching paths.
|
||||
Global memory is the general read-write accessible memory visible to all threads
|
||||
on a given device. Since variables located in global memory have to be marked
|
||||
with the ``__device__`` qualifier, this memory space is also referred to as
|
||||
device memory.
|
||||
|
||||
Without explicitly copying it, it can only be accessed by the threads within a
|
||||
kernel operating on the device, however :ref:`unified_memory` can be used to
|
||||
let the runtime manage this, if desired.
|
||||
|
||||
Allocating global memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
This memory needs to be explicitly allocated.
|
||||
|
||||
It can be allocated from the host via the :ref:`HIP runtime memory management
|
||||
functions <memory_management_reference>` like :cpp:func:`hipMalloc`, or can be
|
||||
defined using the ``__device__`` qualifier on variables.
|
||||
|
||||
It can also be allocated within a kernel using ``malloc`` or ``new``.
|
||||
The specified amount of memory is allocated by each thread that executes the
|
||||
instructions. The recommended way to allocate the memory depends on the use
|
||||
case. If the memory is intended to be shared between the threads of a block, it
|
||||
is generally beneficial to allocate one large block of memory, due to the way
|
||||
the memory is accessed.
|
||||
|
||||
.. note::
|
||||
Memory allocated within a kernel can only be freed in kernels, not by the HIP
|
||||
runtime on the host, like :cpp:func:`hipFree`. It is also not possible to
|
||||
free device memory allocated on the host, with :cpp:func:`hipMalloc` for
|
||||
example, in a kernel.
|
||||
|
||||
|
||||
An example for how to share memory allocated within a kernel by only one thread
|
||||
is given in the following example. In case the device memory is only needed for
|
||||
communication between the threads in a single block, :ref:`shared_memory` is the
|
||||
better option, but is also limited in size.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void kernel_memory_allocation(TYPE* pointer){
|
||||
// The pointer is stored in shared memory, so that all
|
||||
// threads of the block can access the pointer
|
||||
__shared__ int *memory;
|
||||
|
||||
size_t blockSize = blockDim.x;
|
||||
constexpr size_t elementsPerThread = 1024;
|
||||
if(threadIdx.x == 0){
|
||||
// allocate memory in one contiguous block
|
||||
memory = new int[blockDim.x * elementsPerThread];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// load pointer into thread-local variable to avoid
|
||||
// unnecessary accesses to shared memory
|
||||
int *localPtr = memory;
|
||||
|
||||
// work with allocated memory, e.g. initialization
|
||||
for(int i = 0; i < elementsPerThread; ++i){
|
||||
// access in a contiguous way
|
||||
localPtr[i * blockSize + threadIdx.x] = i;
|
||||
}
|
||||
|
||||
// synchronize to make sure no thread is accessing the memory before freeing
|
||||
__syncthreads();
|
||||
if(threadIdx.x == 0){
|
||||
delete[] memory;
|
||||
}
|
||||
}
|
||||
|
||||
Copying between device and host
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
When not using :ref:`unified_memory`, memory has to be explicitly copied between
|
||||
the device and the host, using the HIP runtime API.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
size_t elements = 1 << 20;
|
||||
size_t size_bytes = elements * sizeof(int);
|
||||
|
||||
// allocate host and device memory
|
||||
int *host_pointer = new int[elements];
|
||||
int *device_input, *device_result;
|
||||
HIP_CHECK(hipMalloc(&device_input, size_bytes));
|
||||
HIP_CHECK(hipMalloc(&device_result, size_bytes));
|
||||
|
||||
// copy from host to the device
|
||||
HIP_CHECK(hipMemcpy(device_input, host_pointer, size_bytes, hipMemcpyHostToDevice));
|
||||
|
||||
// Use memory on the device, i.e. execute kernels
|
||||
|
||||
// copy from device to host, to e.g. get results from the kernel
|
||||
HIP_CHECK(hipMemcpy(host_pointer, device_result, size_bytes, hipMemcpyDeviceToHost));
|
||||
|
||||
// free memory when not needed any more
|
||||
HIP_CHECK(hipFree(device_result));
|
||||
HIP_CHECK(hipFree(device_input));
|
||||
delete[] host_pointer;
|
||||
|
||||
Constant memory
|
||||
================================================================================
|
||||
|
||||
Read-only storage visible to all threads on a given device. It is a limited
|
||||
segment backed by device memory with queryable size. It needs to be set by the
|
||||
host before kernel execution. Constant memory provides the best performance
|
||||
benefit when all threads within a warp access the same address.
|
||||
Constant memory is read-only storage visible to all threads on a given device.
|
||||
It is a limited segment backed by device memory, that takes a different caching
|
||||
route than normal device memory accesses. It needs to be set by the host before
|
||||
kernel execution.
|
||||
|
||||
In order to get the highest bandwidth from the constant memory, all threads of
|
||||
a warp have to access the same memory address. If they access different
|
||||
addresses, the accesses get serialized and the bandwidth is therefore reduced.
|
||||
|
||||
Using constant memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Constant memory can not be dynamically allocated, and the size has to be
|
||||
specified during compile time. If the values can not be specified during compile
|
||||
time, they have to be set by the host before the kernel, that accesses the
|
||||
constant memory, is called.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
constexpr size_t const_array_size = 32;
|
||||
__constant__ double const_array[const_array_size];
|
||||
|
||||
void set_constant_memory(double* values){
|
||||
hipMemcpyToSymbol(const_array, values, const_array_size * sizeof(double));
|
||||
}
|
||||
|
||||
__global__ void kernel_using_const_memory(double* array){
|
||||
|
||||
int warpIdx = threadIdx.x / warpSize;
|
||||
// uniform access of warps to const_array for best performance
|
||||
array[blockDim.x] *= const_array[warpIdx];
|
||||
}
|
||||
|
||||
Texture memory
|
||||
================================================================================
|
||||
|
||||
Read-only storage visible to all threads on a given device and accessible
|
||||
through additional APIs. Its origins come from graphics APIs, and provides
|
||||
performance benefits when accessing memory in a pattern where the
|
||||
addresses are close to each other in a 2D representation of the memory.
|
||||
Texture memory is special read-only memory visible to all threads on a given
|
||||
device and accessible through additional APIs. Its origins come from graphics
|
||||
APIs, and provides performance benefits when accessing memory in a pattern where
|
||||
the addresses are close to each other in a 2D or 3D representation of the
|
||||
memory. It also provides additional features like filtering and addressing for
|
||||
out-of-bounds accesses, which are further explained in :ref:`texture_fetching`.
|
||||
|
||||
The :ref:`texture management module <texture_management_reference>` of the HIP
|
||||
runtime API reference contains the functions of texture memory.
|
||||
The original use of the texture cache was also to take pressure off the global
|
||||
memory and other caches, however on modern GPUs, that support textures, the L1
|
||||
cache and texture cache are combined, so the main purpose is to make use of the
|
||||
texture specific features.
|
||||
|
||||
To find out whether textures are supported on a device, query
|
||||
:cpp:enumerator:`hipDeviceAttributeImageSupport`.
|
||||
|
||||
Using texture memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Textures are more complex than just a region of memory, so their layout has to
|
||||
be specified. They are represented by ``hipTextureObject_t`` and created using
|
||||
:cpp:func:`hipCreateTextureObject`.
|
||||
|
||||
The underlying memory is a 1D, 2D or 3D ``hipArray_t``, that needs to be
|
||||
allocated using :cpp:func:`hipMallocArray`.
|
||||
|
||||
On the device side, texture objects are accessed using the ``tex1D/2D/3D``
|
||||
functions.
|
||||
|
||||
The texture management functions can be found in the :ref:`Texture management
|
||||
API reference <texture_management_reference>`
|
||||
|
||||
A full example for how to use textures can be found in the `ROCm texture
|
||||
management example <https://github.com/ROCm/rocm-examples/blob/develop/HIP-Basic/texture_management/main.hip>`_
|
||||
|
||||
Surface memory
|
||||
================================================================================
|
||||
|
||||
A read-write version of texture memory, which can be useful for applications
|
||||
that require direct manipulation of 1D, 2D, or 3D hipArray_t.
|
||||
A read-write version of texture memory. It is created in the same way as a
|
||||
texture, but with :cpp:func:`hipCreateSurfaceObject`.
|
||||
|
||||
Since surfaces are also cached in the read-only texture cache, the changes
|
||||
written back to the surface can't be observed in the same kernel. A new kernel
|
||||
has to be launched in order to see the updated surface.
|
||||
|
||||
The corresponding functions are listed in the :ref:`Surface object API reference
|
||||
<surface_object_reference>`.
|
||||
|
||||
.. _shared_memory:
|
||||
|
||||
Shared memory
|
||||
================================================================================
|
||||
|
||||
Shared memory is read-write memory, that is only visible to the threads within a
|
||||
block. It is allocated per thread block, and needs to be either statically
|
||||
allocated at compile time, or can be dynamically allocated when launching the
|
||||
kernel, but not during kernel execution. Its general use-case is to share
|
||||
variables between the threads within a block, but can also be used as scratch
|
||||
pad memory.
|
||||
|
||||
Shared memory is not backed by the same physical memory as the other address
|
||||
spaces. It is on-chip memory local to the :ref:`compute units
|
||||
<hardware_implementation>`, providing low-latency, high-bandwidth access,
|
||||
comparable to the L1 cache. It is however limited in size, and as it is
|
||||
allocated per block, can restrict how many blocks can be scheduled to a compute
|
||||
unit concurrently, thereby potentially reducing occupancy.
|
||||
|
||||
An overview of the size of the local data share (LDS), that backs up shared
|
||||
memory, is given in the
|
||||
:doc:`GPU hardware specifications <rocm:reference/gpu-arch-specs>`.
|
||||
|
||||
Allocate shared memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Memory can be dynamically allocated by declaring an ``extern __shared__`` array,
|
||||
whose size can be set during kernel launch, which can then be accessed in the
|
||||
kernel.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
extern __shared__ int dynamic_shared[];
|
||||
__global__ void kernel(int array1SizeX, int array1SizeY, int array2Size){
|
||||
// at least (array1SizeX * array1SizeY + array2Size) * sizeof(int) bytes
|
||||
// dynamic shared memory need to be allocated when the kernel is launched
|
||||
int* array1 = dynamic_shared;
|
||||
// array1 is interpreted as 2D of size:
|
||||
int array1Size = array1SizeX * array1SizeY;
|
||||
|
||||
int* array2 = &(array1[array1Size]);
|
||||
|
||||
if(threadIdx.x < array1SizeX && threadIdx.y < array1SizeY){
|
||||
// access array1 with threadIdx.x + threadIdx.y * array1SizeX
|
||||
}
|
||||
if(threadIdx.x < array2Size){
|
||||
// access array2 threadIdx.x
|
||||
}
|
||||
}
|
||||
|
||||
A more in-depth example on dynamically allocated shared memory can be found in
|
||||
the `ROCm dynamic shared example
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/dynamic_shared>`_.
|
||||
|
||||
To statically allocate shared memory, just declare it in the kernel. The memory
|
||||
is allocated per block, not per thread. If the kernel requires more shared
|
||||
memory than is available to the architecture, the compilation fails.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void kernel(){
|
||||
__shared__ int array[128];
|
||||
__shared__ double result;
|
||||
}
|
||||
|
||||
A more in-depth example on statically allocated shared memory can be found in
|
||||
the `ROCm shared memory example
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/shared_memory>`_.
|
||||
|
||||
The :ref:`surface objects module <surface_object_reference>` of HIP runtime API
|
||||
contains the functions for creating, destroying and reading surface memory.
|
||||
+66
-39
@@ -5,56 +5,67 @@
|
||||
|
||||
.. _texture_fetching:
|
||||
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
Texture fetching
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
|
||||
`Textures <../../../../doxygen/html/group___texture.html>`_ are more than just a buffer
|
||||
interpreted as a 1D, 2D, or 3D array.
|
||||
Textures give access to specialized hardware on GPUs that is usually used in
|
||||
graphics processing. In particular, textures use a different way of accessing
|
||||
their underlying device memory. Memory accesses to textures are routed through
|
||||
a special read-only texture cache, that is optimized for logical spatial
|
||||
locality, e.g. locality in 2D grids. This can also benefit certain algorithms
|
||||
used in GPGPU computing, when the access pattern is the same as used when
|
||||
accessing normal textures.
|
||||
|
||||
As textures are associated with graphics, they are indexed using floating-point
|
||||
values. The index can be in the range of [0 to size-1] or [0 to 1].
|
||||
Additionally, textures can be indexed using floating-point values. This is used
|
||||
in graphics applications to interpolate between neighboring values of a texture.
|
||||
Depending on the interpolation mode the index can be in the range of ``0`` to
|
||||
``size - 1`` or ``0`` to ``1``. Textures also have a way of handling
|
||||
out-of-bounds accesses.
|
||||
|
||||
Depending on the index, texture sampling or texture addressing is performed,
|
||||
which decides the return value.
|
||||
Depending on the value of the index, :ref:`texture filtering <texture_filtering>`
|
||||
or :ref:`texture addressing <texture_addressing>` is performed.
|
||||
|
||||
**Texture sampling**: When a texture is indexed with a fraction, the queried
|
||||
value is often between two or more texels (texture elements). The sampling
|
||||
method defines what value to return in such cases.
|
||||
|
||||
**Texture addressing**: Sometimes, the index is outside the bounds of the
|
||||
texture. This condition might look like a problem but helps to put a texture on
|
||||
a surface multiple times or to create a visible sign of out-of-bounds indexing,
|
||||
in computer graphics. The addressing mode defines what value to return when
|
||||
indexing a texture out of bounds.
|
||||
|
||||
The different sampling and addressing modes are described in the following
|
||||
sections.
|
||||
|
||||
Here is the sample texture used in this document for demonstration purposes. It
|
||||
Here is the example texture used in this document for demonstration purposes. It
|
||||
is 2x2 texels and indexed in the [0 to 1] range.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/original.png
|
||||
:width: 150
|
||||
:alt: Sample texture
|
||||
:alt: Example texture
|
||||
:align: center
|
||||
|
||||
Texture used as example
|
||||
|
||||
Texture sampling
|
||||
===============================================================================
|
||||
In HIP textures objects are of type :cpp:struct:`hipTextureObject_t` and created
|
||||
using :cpp:func:`hipCreateTextureObject`.
|
||||
|
||||
Texture sampling handles the usage of fractional indices. It is the method that
|
||||
describes, which nearby values will be used, and how they are combined into the
|
||||
resulting value.
|
||||
For a full list of available texture functions see the :ref:`HIP texture API
|
||||
reference <texture_management_reference>`.
|
||||
|
||||
The various texture sampling methods are discussed in the following sections.
|
||||
A code example for how to use textures can be found in the `ROCm texture
|
||||
management example <https://github.com/ROCm/rocm-examples/blob/develop/HIP-Basic/texture_management/main.hip>`_
|
||||
|
||||
.. _texture_filtering:
|
||||
|
||||
Texture filtering
|
||||
================================================================================
|
||||
|
||||
Texture filtering handles the usage of fractional indices. When the index is a
|
||||
fraction, the queried value lies between two or more texels (texture elements),
|
||||
depending on the dimensionality of the texture. The filtering method defines how
|
||||
to interpolate between these values.
|
||||
|
||||
The filter modes are specified in :cpp:enumerator:`hipTextureFilterMode`.
|
||||
|
||||
The various texture filtering methods are discussed in the following sections.
|
||||
|
||||
.. _texture_fetching_nearest:
|
||||
|
||||
Nearest point sampling
|
||||
Nearest point filtering
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This filter mode corresponds to ``hipFilterModePoint``.
|
||||
|
||||
In this method, the modulo of index is calculated as:
|
||||
|
||||
``tex(x) = T[floor(x)]``
|
||||
@@ -70,22 +81,24 @@ of the nearest texel.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/nearest.png
|
||||
:width: 300
|
||||
:alt: Texture upscaled with nearest point sampling
|
||||
:alt: Texture upscaled with nearest point filtering
|
||||
:align: center
|
||||
|
||||
Texture upscaled with nearest point sampling
|
||||
Texture upscaled with nearest point filtering
|
||||
|
||||
.. _texture_fetching_linear:
|
||||
|
||||
Linear filtering
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This filter mode corresponds to ``hipFilterModeLinear``.
|
||||
|
||||
The linear filtering method does a linear interpolation between values. Linear
|
||||
interpolation is used to create a linear transition between two values. The
|
||||
formula used is ``(1-t)P1 + tP2`` where ``P1`` and ``P2`` are the values and
|
||||
``t`` is within the [0 to 1] range.
|
||||
|
||||
In the case of texture sampling the following formulas are used:
|
||||
In the case of linear texture filtering the following formulas are used:
|
||||
|
||||
* For one dimensional textures: ``tex(x) = (1-α)T[i] + αT[i+1]``
|
||||
* For two dimensional textures: ``tex(x,y) = (1-α)(1-β)T[i,j] + α(1-β)T[i+1,j] + (1-α)βT[i,j+1] + αβT[i+1,j+1]``
|
||||
@@ -95,7 +108,7 @@ Where x, y, and, z are the floating-point indices. i, j, and, k are the integer
|
||||
indices and, α, β, and, γ values represent how far along the sampled point is on
|
||||
the three axes. These values are calculated by these formulas: ``i = floor(x')``, ``α = frac(x')``, ``x' = x - 0.5``, ``j = floor(y')``, ``β = frac(y')``, ``y' = y - 0.5``, ``k = floor(z')``, ``γ = frac(z')`` and ``z' = z - 0.5``
|
||||
|
||||
This following image shows a texture stretched out to a 4x4 pixel quad, but
|
||||
The following image shows a texture stretched out to a 4x4 pixel quad, but
|
||||
still indexed in the [0 to 1] range. The in-between values are interpolated
|
||||
between the neighboring texels.
|
||||
|
||||
@@ -106,12 +119,18 @@ between the neighboring texels.
|
||||
|
||||
Texture upscaled with linear filtering
|
||||
|
||||
.. _texture_addressing:
|
||||
|
||||
Texture addressing
|
||||
===============================================================================
|
||||
|
||||
Texture addressing mode handles the index that is out of bounds of the texture.
|
||||
This mode describes which values of the texture or a preset value to use when
|
||||
the index is out of bounds.
|
||||
The texture addressing modes are specified in
|
||||
:cpp:enumerator:`hipTextureAddressMode`.
|
||||
|
||||
The texture addressing mode handles out-of-bounds accesses to the texture. This
|
||||
can be used in graphics applications to e.g. repeat a texture on a surface
|
||||
multiple times in various ways or create visible signs of out-of-bounds
|
||||
indexing.
|
||||
|
||||
The following sections describe the various texture addressing methods.
|
||||
|
||||
@@ -120,8 +139,10 @@ The following sections describe the various texture addressing methods.
|
||||
Address mode border
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
In this method, the texture fetching returns a border value when indexing out of
|
||||
bounds. The border value must be set before texture fetching.
|
||||
This addressing mode is set using ``hipAddressModeBorder``.
|
||||
|
||||
This addressing mode returns a border value when indexing out of bounds. The
|
||||
border value must be set before texture fetching.
|
||||
|
||||
The following image shows the texture on a 4x4 pixel quad, indexed in the
|
||||
[0 to 3] range. The out-of-bounds values are the border color, which is yellow.
|
||||
@@ -141,6 +162,8 @@ the addressing begins.
|
||||
Address mode clamp
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This addressing mode is set using ``hipAddressModeClamp``.
|
||||
|
||||
This mode clamps the index between [0 to size-1]. Due to this, when indexing
|
||||
out-of-bounds, the values on the edge of the texture repeat. The clamp mode is
|
||||
the default addressing mode.
|
||||
@@ -164,6 +187,8 @@ the addressing begins.
|
||||
Address mode wrap
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This addressing mode is set using ``hipAddressModeWrap``.
|
||||
|
||||
Wrap mode addressing is only available for normalized texture coordinates. In
|
||||
this addressing mode, the fractional part of the index is used:
|
||||
|
||||
@@ -189,6 +214,8 @@ the addressing begins.
|
||||
Address mode mirror
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This addressing mode is set using ``hipAddressModeMirror``.
|
||||
|
||||
Similar to the wrap mode the mirror mode is only available for normalized
|
||||
texture coordinates and also creates a repeating image, but mirroring the
|
||||
neighboring instances.
|
||||
|
||||
@@ -111,8 +111,7 @@ allocator can be used.
|
||||
❌: **Unsupported**
|
||||
|
||||
:sup:`1` Works only with ``HSA_XNACK=1`` and kernels with HMM support. First GPU
|
||||
access causes recoverable page-fault. For more details, visit `GPU memory
|
||||
<https://rocm.docs.amd.com/en/latest/conceptual/gpu-memory.html#xnack>`_.
|
||||
access causes recoverable page-fault.
|
||||
|
||||
.. _unified memory allocators:
|
||||
|
||||
@@ -144,8 +143,7 @@ GPUs, it is essential to set the environment variable ``HSA_XNACK=1`` and use
|
||||
a GPU kernel mode driver that supports HMM
|
||||
<https://www.kernel.org/doc/html/latest/mm/hmm.html>`_. Without this
|
||||
configuration, the behavior will be similar to that of systems without HMM
|
||||
support. For more details, visit
|
||||
`GPU memory <https://rocm.docs.amd.com/en/latest/conceptual/gpu-memory.html#xnack>`_.
|
||||
support.
|
||||
|
||||
The table below illustrates the expected behavior of managed and unified memory
|
||||
functions on ROCm and CUDA, both with and without HMM support.
|
||||
|
||||
+271
-42
@@ -28,7 +28,7 @@ reduce memory usage and unnecessary ``memcpy`` calls.
|
||||
.. _memory_allocation_virtual_memory:
|
||||
|
||||
Memory allocation
|
||||
================================================================================
|
||||
=================
|
||||
|
||||
Standard memory allocation uses the :cpp:func:`hipMalloc` function to allocate a
|
||||
block of memory on the device. However, when using virtual memory, this process
|
||||
@@ -37,10 +37,34 @@ is separated into multiple steps using the :cpp:func:`hipMemCreate`,
|
||||
:cpp:func:`hipMemSetAccess` functions. This guide explains what these functions
|
||||
do and how you can use them for virtual memory management.
|
||||
|
||||
Allocate physical memory
|
||||
--------------------------------------------------------------------------------
|
||||
.. _vmm_support:
|
||||
|
||||
The first step is to allocate the physical memory itself with the
|
||||
Virtual memory management support
|
||||
---------------------------------
|
||||
|
||||
The first step is to check if the targeted device or GPU supports virtual memory management.
|
||||
Use the :cpp:func:`hipDeviceGetAttribute` function to get the
|
||||
``hipDeviceAttributeVirtualMemoryManagementSupported`` attribute for a specific GPU, as shown in the following example.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
int vmm = 0, currentDev = 0;
|
||||
hipDeviceGetAttribute(
|
||||
&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, currentDev
|
||||
);
|
||||
|
||||
if (vmm == 0) {
|
||||
std::cout << "GPU " << currentDev << " doesn't support virtual memory management." << std::endl;
|
||||
} else {
|
||||
std::cout << "GPU " << currentDev << " support virtual memory management." << std::endl;
|
||||
}
|
||||
|
||||
.. _allocate_physical_memory:
|
||||
|
||||
Allocate physical memory
|
||||
------------------------
|
||||
|
||||
The next step is to allocate the physical memory using the
|
||||
:cpp:func:`hipMemCreate` function. This function accepts the size of the buffer,
|
||||
an ``unsigned long long`` variable for the flags, and a
|
||||
:cpp:struct:`hipMemAllocationProp` variable. :cpp:struct:`hipMemAllocationProp`
|
||||
@@ -48,42 +72,54 @@ contains the properties of the memory to be allocated, such as where the memory
|
||||
is physically located and what kind of shareable handles are available. If the
|
||||
allocation is successful, the function returns a value of
|
||||
:cpp:enumerator:`hipSuccess`, with :cpp:type:`hipMemGenericAllocationHandle_t`
|
||||
representing a valid physical memory allocation. The allocated memory size must
|
||||
be aligned with the granularity appropriate for the properties of the
|
||||
allocation. You can use the :cpp:func:`hipMemGetAllocationGranularity` function
|
||||
to determine the correct granularity.
|
||||
representing a valid physical memory allocation.
|
||||
|
||||
The allocated memory must be aligned with the appropriate granularity. The
|
||||
granularity value can be queried with :cpp:func:`hipMemGetAllocationGranularity`,
|
||||
and its value depends on the target device hardware and the type of memory
|
||||
allocation. If the allocation size is not aligned, meaning it is not cleanly
|
||||
divisible by the minimum granularity value, :cpp:func:`hipMemCreate` will return
|
||||
an out-of-memory error.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
size_t granularity = 0;
|
||||
hipMemGenericAllocationHandle_t allocHandle;
|
||||
hipMemAllocationProp prop = {};
|
||||
prop.type = HIP_MEM_ALLOCATION_TYPE_PINNED;
|
||||
prop.location.type = HIP_MEM_LOCATION_TYPE_DEVICE;
|
||||
// The pinned allocation type cannot be migrated from its current location
|
||||
// while the application is actively using it.
|
||||
prop.type = hipMemAllocationTypePinned;
|
||||
// Set the location type to device, currently there are no other valid option.
|
||||
prop.location.type = hipMemLocationTypeDevice;
|
||||
// Set the device id, where the memory will be allocated.
|
||||
prop.location.id = currentDev;
|
||||
hipMemGetAllocationGranularity(&granularity, &prop, HIP_MEM_ALLOC_GRANULARITY_MINIMUM);
|
||||
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum);
|
||||
padded_size = ROUND_UP(size, granularity);
|
||||
hipMemCreate(&allocHandle, padded_size, &prop, 0);
|
||||
|
||||
Reserve virtual address range
|
||||
--------------------------------------------------------------------------------
|
||||
.. _reserve_virtual_address:
|
||||
|
||||
After you have acquired an allocation of physical memory, you must map it before
|
||||
you can use it. To do so, you need a virtual address to map it to. Mapping
|
||||
means the physical memory allocation is available from the virtual address range
|
||||
it is mapped to. To reserve a virtual memory range, use the
|
||||
:cpp:func:`hipMemAddressReserve` function. The size of the virtual memory must
|
||||
match the amount of physical memory previously allocated. You can then map the
|
||||
physical memory allocation to the newly-acquired virtual memory address range
|
||||
using the :cpp:func:`hipMemMap` function.
|
||||
Reserve virtual address range
|
||||
-----------------------------
|
||||
|
||||
After you have acquired an allocation of physical memory, you must map it to a
|
||||
virtual address before you can use it. Mapping means the physical memory
|
||||
allocation is available from the virtual address range it is mapped to. To
|
||||
reserve a virtual memory range, use the :cpp:func:`hipMemAddressReserve`
|
||||
function. The size of the virtual memory must match the amount of physical
|
||||
memory previously allocated. You can then map the physical memory allocation to
|
||||
the newly-acquired virtual memory address range using the :cpp:func:`hipMemMap`
|
||||
function.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&ptr, padded_size, 0, 0, 0);
|
||||
hipMemMap(ptr, padded_size, 0, allocHandle, 0);
|
||||
|
||||
.. _set_memory_access:
|
||||
|
||||
Set memory access
|
||||
--------------------------------------------------------------------------------
|
||||
-----------------
|
||||
|
||||
Finally, use the :cpp:func:`hipMemSetAccess` function to enable memory access.
|
||||
It accepts the pointer to the virtual memory, the size, and a
|
||||
@@ -103,16 +139,39 @@ devices.
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAccessDesc accessDesc = {};
|
||||
accessDesc.location.type = HIP_MEM_LOCATION_TYPE_DEVICE;
|
||||
accessDesc.location.type = hipMemLocationTypeDevice;
|
||||
accessDesc.location.id = currentDev;
|
||||
accessDesc.flags = HIP_MEM_ACCESS_FLAGS_PROT_READWRITE;
|
||||
accessDesc.flags = hipMemAccessFlagsProtReadwrite;
|
||||
hipMemSetAccess(ptr, padded_size, &accessDesc, 1);
|
||||
|
||||
At this point the memory is allocated, mapped, and ready for use. You can read
|
||||
and write to it, just like you would a C style memory allocation.
|
||||
|
||||
.. _usage_virtual_memory:
|
||||
|
||||
Dynamically increase allocation size
|
||||
------------------------------------
|
||||
|
||||
To increase the amount of pre-allocated memory, use
|
||||
:cpp:func:`hipMemAddressReserve`, which accepts the starting address, and the
|
||||
size of the reservation in bytes. This allows you to have a continuous virtual
|
||||
address space without worrying about the underlying physical allocation.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&new_ptr, (new_size - padded_size), 0, ptr + padded_size, 0);
|
||||
hipMemMap(new_ptr, (new_size - padded_size), 0, newAllocHandle, 0);
|
||||
hipMemSetAccess(new_ptr, (new_size - padded_size), &accessDesc, 1);
|
||||
|
||||
The code sample above assumes that :cpp:func:`hipMemAddressReserve` was able to
|
||||
reserve the memory address at the specified location. However, this isn't
|
||||
guaranteed to be true, so you should validate that ``new_ptr`` points to a
|
||||
specific virtual address before using it.
|
||||
|
||||
.. _free_virtual_memory:
|
||||
|
||||
Free virtual memory
|
||||
--------------------------------------------------------------------------------
|
||||
-------------------
|
||||
|
||||
To free the memory allocated in this manner, use the corresponding free
|
||||
functions. To unmap the memory, use :cpp:func:`hipMemUnmap`. To release the
|
||||
@@ -128,27 +187,197 @@ synchronizes the device. This causes worse resource usage and performance.
|
||||
hipMemRelease(allocHandle);
|
||||
hipMemAddressFree(ptr, size);
|
||||
|
||||
.. _usage_virtual_memory:
|
||||
Example code
|
||||
============
|
||||
|
||||
Memory usage
|
||||
================================================================================
|
||||
The virtual memory management example follows these steps:
|
||||
|
||||
Dynamically increase allocation size
|
||||
--------------------------------------------------------------------------------
|
||||
1. Check virtual memory management :ref:`support <vmm_support>`:
|
||||
The :cpp:func:`hipDeviceGetAttribute` function is used to check the virtual
|
||||
memory management support of the GPU with ID 0.
|
||||
|
||||
The :cpp:func:`hipMemAddressReserve` function allows you to increase the amount
|
||||
of pre-allocated memory. This function accepts a parameter representing the
|
||||
requested starting address of the virtual memory. This allows you to have a
|
||||
continuous virtual address space without worrying about the underlying physical
|
||||
allocation.
|
||||
2. Physical memory :ref:`allocation <allocate_physical_memory>`: Physical memory
|
||||
is allocated using :cpp:func:`hipMemCreate` with pinned memory on the
|
||||
device.
|
||||
|
||||
3. Virtual memory :ref:`reservation <reserve_virtual_address>`: Virtual address
|
||||
range is reserved using :cpp:func:`hipMemAddressReserve`.
|
||||
|
||||
4. Mapping virtual address to physical memory: The physical memory is mapped
|
||||
to a virtual address (``virtualPointer``) using :cpp:func:`hipMemMap`.
|
||||
|
||||
5. Memory :ref:`access permissions<set_memory_access>`: Permission is set for
|
||||
pointer to allow read and write access using :cpp:func:`hipMemSetAccess`.
|
||||
|
||||
6. Memory operation: Data is written to the memory via ``virtualPointer``.
|
||||
|
||||
7. Launch kernels: The ``zeroAddr`` and ``fillAddr`` kernels are
|
||||
launched using the virtual memory pointer.
|
||||
|
||||
8. :ref:`Cleanup <free_virtual_memory>`: The mappings, physical memory, and
|
||||
virtual address are released at the end to avoid memory leaks.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&new_ptr, (new_size - padded_size), 0, ptr + padded_size, 0);
|
||||
hipMemMap(new_ptr, (new_size - padded_size), 0, newAllocHandle, 0);
|
||||
hipMemSetAccess(new_ptr, (new_size - padded_size), &accessDesc, 1);
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define ROUND_UP(SIZE,GRANULARITY) ((1 + SIZE / GRANULARITY) * GRANULARITY)
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void zeroAddr(int* pointer) {
|
||||
*pointer = 0;
|
||||
}
|
||||
|
||||
__global__ void fillAddr(int* pointer) {
|
||||
*pointer = 42;
|
||||
}
|
||||
|
||||
|
||||
int main() {
|
||||
|
||||
int currentDev = 0;
|
||||
|
||||
// Step 1: Check virtual memory management support on device 0
|
||||
int vmm = 0;
|
||||
HIP_CHECK(
|
||||
hipDeviceGetAttribute(
|
||||
&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, currentDev
|
||||
)
|
||||
);
|
||||
|
||||
std::cout << "Virtual memory management support value: " << vmm << std::endl;
|
||||
|
||||
if (vmm == 0) {
|
||||
std::cout << "GPU 0 doesn't support virtual memory management.";
|
||||
return 0;
|
||||
}
|
||||
|
||||
// Size of memory to allocate
|
||||
size_t size = 4 * 1024;
|
||||
|
||||
// Step 2: Allocate physical memory
|
||||
hipMemGenericAllocationHandle_t allocHandle;
|
||||
hipMemAllocationProp prop = {};
|
||||
prop.type = hipMemAllocationTypePinned;
|
||||
prop.location.type = hipMemLocationTypeDevice;
|
||||
prop.location.id = currentDev;
|
||||
size_t granularity = 0;
|
||||
HIP_CHECK(
|
||||
hipMemGetAllocationGranularity(
|
||||
&granularity,
|
||||
&prop,
|
||||
hipMemAllocationGranularityMinimum));
|
||||
size_t padded_size = ROUND_UP(size, granularity);
|
||||
HIP_CHECK(hipMemCreate(&allocHandle, padded_size * 2, &prop, 0));
|
||||
|
||||
// Step 3: Reserve a virtual memory address range
|
||||
void* virtualPointer = nullptr;
|
||||
HIP_CHECK(hipMemAddressReserve(&virtualPointer, padded_size, granularity, nullptr, 0));
|
||||
|
||||
// Step 4: Map the physical memory to the virtual address range
|
||||
HIP_CHECK(hipMemMap(virtualPointer, padded_size, 0, allocHandle, 0));
|
||||
|
||||
// Step 5: Set memory access permission for pointer
|
||||
hipMemAccessDesc accessDesc = {};
|
||||
accessDesc.location.type = hipMemLocationTypeDevice;
|
||||
accessDesc.location.id = currentDev;
|
||||
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
|
||||
|
||||
HIP_CHECK(hipMemSetAccess(virtualPointer, padded_size, &accessDesc, 1));
|
||||
|
||||
// Step 6: Perform memory operation
|
||||
int value = 42;
|
||||
HIP_CHECK(hipMemcpy(virtualPointer, &value, sizeof(int), hipMemcpyHostToDevice));
|
||||
|
||||
int result = 1;
|
||||
HIP_CHECK(hipMemcpy(&result, virtualPointer, sizeof(int), hipMemcpyDeviceToHost));
|
||||
if( result == 42) {
|
||||
std::cout << "Success. Value: " << result << std::endl;
|
||||
} else {
|
||||
std::cout << "Failure. Value: " << result << std::endl;
|
||||
}
|
||||
|
||||
// Step 7: Launch kernels
|
||||
// Launch zeroAddr kernel
|
||||
zeroAddr<<<1, 1>>>((int*)virtualPointer);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Check zeroAddr kernel result
|
||||
result = 1;
|
||||
HIP_CHECK(hipMemcpy(&result, virtualPointer, sizeof(int), hipMemcpyDeviceToHost));
|
||||
if( result == 0) {
|
||||
std::cout << "Success. zeroAddr kernel: " << result << std::endl;
|
||||
} else {
|
||||
std::cout << "Failure. zeroAddr kernel: " << result << std::endl;
|
||||
}
|
||||
|
||||
// Launch fillAddr kernel
|
||||
fillAddr<<<1, 1>>>((int*)virtualPointer);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Check fillAddr kernel result
|
||||
result = 1;
|
||||
HIP_CHECK(hipMemcpy(&result, virtualPointer, sizeof(int), hipMemcpyDeviceToHost));
|
||||
if( result == 42) {
|
||||
std::cout << "Success. fillAddr kernel: " << result << std::endl;
|
||||
} else {
|
||||
std::cout << "Failure. fillAddr kernel: " << result << std::endl;
|
||||
}
|
||||
|
||||
// Step 8: Cleanup
|
||||
HIP_CHECK(hipMemUnmap(virtualPointer, padded_size));
|
||||
HIP_CHECK(hipMemRelease(allocHandle));
|
||||
HIP_CHECK(hipMemAddressFree(virtualPointer, padded_size));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Virtual aliases
|
||||
================================================================================
|
||||
|
||||
Virtual aliases are multiple virtual memory addresses mapping to the same
|
||||
physical memory on the GPU. When this occurs, different threads, processes, or memory
|
||||
allocations to access shared physical memory through different virtual
|
||||
addresses on different devices.
|
||||
|
||||
Multiple virtual memory mappings can be created using multiple calls to
|
||||
:cpp:func:`hipMemMap` on the same memory allocation.
|
||||
|
||||
.. note::
|
||||
|
||||
RDNA cards may not produce correct results, if users access two different
|
||||
virtual addresses that map to the same physical address. In this case, the
|
||||
L1 data caches will be incoherent due to the virtual-to-physical aliasing.
|
||||
These GPUs will produce correct results if users access virtual-to-physical
|
||||
aliases using volatile pointers.
|
||||
|
||||
NVIDIA GPUs require special fences to produce correct results when
|
||||
using virtual aliases.
|
||||
|
||||
In the following code block, the kernels input device pointers are virtual
|
||||
aliases of the same memory allocation:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void updateBoth(int* pointerA, int* pointerB) {
|
||||
// May produce incorrect results on RDNA and NVIDIA cards.
|
||||
*pointerA = 0;
|
||||
*pointerB = 42;
|
||||
}
|
||||
|
||||
__global__ void updateBoth_v2(volatile int* pointerA, volatile int* pointerB) {
|
||||
// May produce incorrect results on NVIDIA cards.
|
||||
*pointerA = 0;
|
||||
*pointerB = 42;
|
||||
}
|
||||
|
||||
The code sample above assumes that :cpp:func:`hipMemAddressReserve` was able to
|
||||
reserve the memory address at the specified location. However, this isn't
|
||||
guaranteed to be true, so you should validate that ``new_ptr`` points to a
|
||||
specific virtual address before using it.
|
||||
|
||||
@@ -0,0 +1,209 @@
|
||||
.. meta::
|
||||
:description: This chapter describes HIP's kernel language's C++ support.
|
||||
:keywords: AMD, ROCm, HIP, C++ support
|
||||
|
||||
################################################################################
|
||||
Kernel language C++ support
|
||||
################################################################################
|
||||
|
||||
The HIP host API can be compiled with any conforming C++ compiler, as long as no
|
||||
kernel launch is present in the code.
|
||||
|
||||
To compile device code and include kernel launches, a compiler with full HIP
|
||||
support is needed, such as ``amdclang++``. For more information, see :doc:`ROCm
|
||||
compilers <llvm-project:index>`.
|
||||
|
||||
In host code all modern C++ standards that are supported by the compiler can be
|
||||
used. Device code compilation has some restrictions on modern C++ standards, but
|
||||
in general also supports all C++ standards. The biggest restriction is the
|
||||
reduced support of the C++ standard library in device code, as functions are
|
||||
only compiled for the host by default. An exception to this are ``constexpr``
|
||||
functions that are resolved at compile time and can be used in device code.
|
||||
There are ongoing efforts to implement C++ standard library functionality with
|
||||
`libhipcxx <https://github.com/ROCm/libhipcxx>`_.
|
||||
|
||||
********************************************************************************
|
||||
Supported kernel language C++ features
|
||||
********************************************************************************
|
||||
|
||||
This section describes HIP's kernel language C++ feature support for the
|
||||
different versions of the standard.
|
||||
|
||||
General C++ features
|
||||
===============================================================================
|
||||
|
||||
Exception handling
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
An important difference between the host and device code C++ support is
|
||||
exception handling. In device code, exceptions aren't available due to
|
||||
the hardware architecture. The device code must use return codes to handle
|
||||
errors.
|
||||
|
||||
Assertions
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The ``assert`` function is supported in device code. Assertions are used for
|
||||
debugging purposes. When the input expression equals zero, the execution will be
|
||||
stopped. HIP provides its own implementation for ``assert`` for usage in device
|
||||
code in ``hip/hip_runtime.h``.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
void assert(int input)
|
||||
|
||||
HIP also provides the function ``abort()`` which can be used to terminate the
|
||||
application when terminal failures are detected. It is implemented using the
|
||||
``__builtin_trap()`` function.
|
||||
|
||||
This function produces a similar effect as using CUDA's ``asm("trap")``.
|
||||
In HIP, ``abort()`` terminates the entire application, while in CUDA,
|
||||
``asm("trap")`` only terminates the current kernel and the application continues
|
||||
to run.
|
||||
|
||||
printf
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
``printf`` is supported in device code, and can be used just like in host code.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
__global__ void run_printf() { printf("Hello World\n"); }
|
||||
|
||||
int main() {
|
||||
run_printf<<<dim3(1), dim3(1), 0, 0>>>();
|
||||
}
|
||||
|
||||
Device-Side Dynamic Global Memory Allocation
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Device code can use ``new`` or ``malloc`` to dynamically allocate global
|
||||
memory on the device, and ``delete`` or ``free`` to deallocate global memory.
|
||||
|
||||
Classes
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Classes work on both host and device side, with some constraints on the device
|
||||
side.
|
||||
|
||||
Member functions with the appropriate qualifiers can be called in host and
|
||||
device code, and the corresponding overload is executed.
|
||||
|
||||
``virtual`` member functions are also supported, however calling these functions
|
||||
from the host if the object was created on the device, or the other way around,
|
||||
is undefined behaviour.
|
||||
|
||||
The ``__host__``, ``__device__``, ``__managed__``, ``__shared__`` and
|
||||
``__constant__`` memory space qualifiers can not be applied to member variables.
|
||||
|
||||
C++11 support
|
||||
===============================================================================
|
||||
|
||||
``constexpr``
|
||||
Full support in device code. ``constexpr`` implicitly defines ``__host__
|
||||
__device__``, so standard library functions that are marked ``constexpr`` can
|
||||
be used in device code.
|
||||
``constexpr`` variables can be used in both host and device code.
|
||||
|
||||
Lambdas
|
||||
Lambdas are implicitly marked with ``__host__ __device__``. To mark them as
|
||||
only executable for the host or the device, they can be explicitly marked like
|
||||
any other function. There are restrictions on variable capture, however. Host
|
||||
and device specific variables can only be accessed on other devices or the
|
||||
host by explicitly copying them. Accessing captured the variables by
|
||||
reference, when the variable is not located on the executing device or host,
|
||||
causes undefined behaviour.
|
||||
|
||||
Polymorphic function wrappers
|
||||
HIP does not support the polymorphic function wrapper ``std::function``
|
||||
|
||||
|
||||
C++14 support
|
||||
===============================================================================
|
||||
|
||||
All `C++14 language features <https://isocpp.org/wiki/faq/cpp14-language>`_ are
|
||||
supported.
|
||||
|
||||
C++17 support
|
||||
===============================================================================
|
||||
|
||||
All `C++17 language features <https://en.cppreference.com/w/cpp/17>`_ are
|
||||
supported.
|
||||
|
||||
C++20 support
|
||||
===============================================================================
|
||||
|
||||
Most `C++20 language features <https://en.cppreference.com/w/cpp/20>`_ are
|
||||
supported, but some restrictions apply. Coroutines are not available in device
|
||||
code.
|
||||
|
||||
********************************************************************************
|
||||
Compiler features
|
||||
********************************************************************************
|
||||
|
||||
Pragma Unroll
|
||||
================================================================================
|
||||
|
||||
The unroll pragma for unrolling loops with a compile-time constant is supported:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#pragma unroll 16 /* hint to compiler to unroll next loop by 16 */
|
||||
for (int i=0; i<16; i++) ...
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#pragma unroll 1 /* tell compiler to never unroll the loop */
|
||||
for (int i=0; i<16; i++) ...
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#pragma unroll /* hint to compiler to completely unroll next loop. */
|
||||
for (int i=0; i<16; i++) ...
|
||||
|
||||
In-Line Assembly
|
||||
================================================================================
|
||||
|
||||
GCN ISA In-line assembly can be included in device code.
|
||||
|
||||
It has to be mentioned however, that in-line assembly should be used carefully.
|
||||
For more information, please refer to the
|
||||
:doc:`Inline ASM statements section of amdclang<llvm-project:reference/rocmcc>`.
|
||||
|
||||
A short example program including inline assembly can be found in
|
||||
`HIP inline_assembly sample
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/inline_assembly>`_.
|
||||
|
||||
For information on what special AMD GPU hardware features are available
|
||||
through assembly, please refer to the `ISA manuals of the corresponding
|
||||
architecture
|
||||
<https://llvm.org/docs/AMDGPUUsage.html#additional-documentation>`_.
|
||||
|
||||
Kernel Compilation
|
||||
================================================================================
|
||||
|
||||
``hipcc`` now supports compiling C++/HIP kernels to binary code objects. The
|
||||
file format for the binary files is usually ``.co`` which means Code Object.
|
||||
The following command builds the code object using ``hipcc``.
|
||||
|
||||
.. code-block:: bash
|
||||
|
||||
hipcc --genco --offload-arch=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]
|
||||
|
||||
[TARGET GPU] = GPU architecture
|
||||
[INPUT FILE] = Name of the file containing source code
|
||||
[OUTPUT FILE] = Name of the generated code object file
|
||||
|
||||
For an example on how to use these object files, refer to the `HIP module_api
|
||||
sample
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/module_api>`_.
|
||||
|
||||
Architecture specific code
|
||||
================================================================================
|
||||
|
||||
``amdclang++`` defines ``__gfx*__`` macros based on the GPU architecture to be
|
||||
compiled for. These macros can be used to include GPU architecture specific
|
||||
code. Refer to the sample in `HIP gpu_arch sample
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/gpu_arch>`_.
|
||||
@@ -3,6 +3,8 @@
|
||||
developers optimize the performance of HIP-capable GPU architectures.
|
||||
:keywords: AMD, ROCm, HIP, CUDA, performance, guidelines
|
||||
|
||||
.. _how_to_performance_guidelines:
|
||||
|
||||
*******************************************************************************
|
||||
Performance guidelines
|
||||
*******************************************************************************
|
||||
@@ -32,12 +34,14 @@ reveal and efficiently provide as much parallelism as possible. The parallelism
|
||||
can be performed at the application level, device level, and multiprocessor
|
||||
level.
|
||||
|
||||
.. _application_parallel_execution:
|
||||
|
||||
Application level
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
To enable parallel execution of the application across the host and devices, use
|
||||
asynchronous calls and streams. Assign workloads based on efficiency: serial to
|
||||
the host or parallel to the devices.
|
||||
:ref:`asynchronous calls and streams <asynchronous_how-to>`. Assign workloads
|
||||
based on efficiency: serial to the host or parallel to the devices.
|
||||
|
||||
For parallel workloads, when threads belonging to the same block need to
|
||||
synchronize to share data, use :cpp:func:`__syncthreads()` (see:
|
||||
|
||||
@@ -30,6 +30,8 @@ The HIP documentation is organized into the following categories:
|
||||
* [Debugging with HIP](./how-to/debugging)
|
||||
* {doc}`./how-to/logging`
|
||||
* {doc}`./how-to/hip_runtime_api`
|
||||
* {doc}`./how-to/hip_cpp_language_extensions`
|
||||
* {doc}`./how-to/kernel_language_cpp_support`
|
||||
* [HIP porting guide](./how-to/hip_porting_guide)
|
||||
* [HIP porting: driver API guide](./how-to/hip_porting_driver_api)
|
||||
* {doc}`./how-to/hip_rtc`
|
||||
@@ -41,11 +43,9 @@ The HIP documentation is organized into the following categories:
|
||||
|
||||
* [HIP runtime API](./reference/hip_runtime_api_reference)
|
||||
* [HSA runtime API for ROCm](./reference/virtual_rocr)
|
||||
* [C++ language extensions](./reference/cpp_language_extensions)
|
||||
* [C++ language support](./reference/cpp_language_support)
|
||||
* [HIP math API](./reference/math_api)
|
||||
* [HIP environment variables](./reference/env_variables)
|
||||
* [Comparing syntax for different APIs](./reference/terms)
|
||||
* [CUDA to HIP API Function Comparison](./reference/api_syntax)
|
||||
* [List of deprecated APIs](./reference/deprecated_api_list)
|
||||
* [FP8 numbers in HIP](./reference/fp8_numbers)
|
||||
* {doc}`./reference/hardware_features`
|
||||
@@ -55,8 +55,7 @@ The HIP documentation is organized into the following categories:
|
||||
:::{grid-item-card} Tutorial
|
||||
|
||||
* [HIP basic examples](https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic)
|
||||
* [HIP examples](https://github.com/ROCm/HIP-Examples)
|
||||
* [HIP test samples](https://github.com/ROCm/hip-tests/tree/develop/samples)
|
||||
* [HIP examples](https://github.com/ROCm/rocm-examples)
|
||||
* [SAXPY tutorial](./tutorial/saxpy)
|
||||
* [Reduction tutorial](./tutorial/reduction)
|
||||
* [Cooperative groups tutorial](./tutorial/cooperative_groups_tutorial)
|
||||
|
||||
@@ -238,4 +238,4 @@ Run HIP
|
||||
=================================================
|
||||
|
||||
After installation and building HIP, you can compile your application and run.
|
||||
A simple example is `square sample <https://github.com/ROCm/hip-tests/tree/develop/samples/0_Intro/square>`_.
|
||||
Simple examples can be found in the `ROCm-examples repository <https://github.com/ROCm/rocm-examples>`_.
|
||||
|
||||
@@ -0,0 +1,176 @@
|
||||
.. meta::
|
||||
:description: Maps CUDA API syntax to HIP API syntax with an example
|
||||
:keywords: AMD, ROCm, HIP, CUDA, syntax, HIP syntax
|
||||
|
||||
********************************************************************************
|
||||
CUDA to HIP API Function Comparison
|
||||
********************************************************************************
|
||||
|
||||
This page introduces key syntax differences between CUDA and HIP APIs with a focused code
|
||||
example and comparison table. For a complete list of mappings, visit :ref:`HIPIFY <HIPIFY:index>`.
|
||||
|
||||
The following CUDA code example illustrates several CUDA API syntaxes.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
__global__ void block_reduction(const float* input, float* output, int num_elements)
|
||||
{
|
||||
extern __shared__ float s_data[];
|
||||
|
||||
int tid = threadIdx.x;
|
||||
int global_id = blockDim.x * blockIdx.x + tid;
|
||||
|
||||
if (global_id < num_elements)
|
||||
{
|
||||
s_data[tid] = input[global_id];
|
||||
}
|
||||
else
|
||||
{
|
||||
s_data[tid] = 0.0f;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
|
||||
{
|
||||
if (tid < stride)
|
||||
{
|
||||
s_data[tid] += s_data[tid + stride];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
if (tid == 0)
|
||||
{
|
||||
output[blockIdx.x] = s_data[0];
|
||||
}
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
int threads = 256;
|
||||
const int num_elements = 50000;
|
||||
|
||||
std::vector<float> h_a(num_elements);
|
||||
std::vector<float> h_b((num_elements + threads - 1) / threads);
|
||||
|
||||
for (int i = 0; i < num_elements; ++i)
|
||||
{
|
||||
h_a[i] = rand() / static_cast<float>(RAND_MAX);
|
||||
}
|
||||
|
||||
float *d_a, *d_b;
|
||||
cudaMalloc(&d_a, h_a.size() * sizeof(float));
|
||||
cudaMalloc(&d_b, h_b.size() * sizeof(float));
|
||||
|
||||
cudaStream_t stream;
|
||||
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
|
||||
|
||||
cudaEvent_t start_event, stop_event;
|
||||
cudaEventCreate(&start_event);
|
||||
cudaEventCreate(&stop_event);
|
||||
|
||||
cudaMemcpyAsync(d_a, h_a.data(), h_a.size() * sizeof(float), cudaMemcpyHostToDevice, stream);
|
||||
|
||||
cudaEventRecord(start_event, stream);
|
||||
|
||||
int blocks = (num_elements + threads - 1) / threads;
|
||||
block_reduction<<<blocks, threads, threads * sizeof(float), stream>>>(d_a, d_b, num_elements);
|
||||
|
||||
cudaMemcpyAsync(h_b.data(), d_b, h_b.size() * sizeof(float), cudaMemcpyDeviceToHost, stream);
|
||||
|
||||
cudaEventRecord(stop_event, stream);
|
||||
cudaEventSynchronize(stop_event);
|
||||
|
||||
cudaEventElapsedTime(&milliseconds, start_event, stop_event);
|
||||
std::cout << "Kernel execution time: " << milliseconds << " ms\n";
|
||||
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_b);
|
||||
|
||||
cudaEventDestroy(start_event);
|
||||
cudaEventDestroy(stop_event);
|
||||
cudaStreamDestroy(stream);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
The following table maps CUDA API functions to corresponding HIP API functions, as demonstrated in the
|
||||
preceding code examples.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:name: syntax-mapping-table
|
||||
|
||||
*
|
||||
- CUDA
|
||||
- HIP
|
||||
|
||||
*
|
||||
- ``#include <cuda_runtime.h>``
|
||||
- ``#include <hip/hip_runtime.h>``
|
||||
|
||||
*
|
||||
- ``cudaError_t``
|
||||
- ``hipError_t``
|
||||
|
||||
*
|
||||
- ``cudaEvent_t``
|
||||
- ``hipEvent_t``
|
||||
|
||||
*
|
||||
- ``cudaStream_t``
|
||||
- ``hipStream_t``
|
||||
|
||||
*
|
||||
- ``cudaMalloc``
|
||||
- ``hipMalloc``
|
||||
|
||||
*
|
||||
- ``cudaStreamCreateWithFlags``
|
||||
- ``hipStreamCreateWithFlags``
|
||||
|
||||
*
|
||||
- ``cudaStreamNonBlocking``
|
||||
- ``hipStreamNonBlocking``
|
||||
|
||||
*
|
||||
- ``cudaEventCreate``
|
||||
- ``hipEventCreate``
|
||||
|
||||
*
|
||||
- ``cudaMemcpyAsync``
|
||||
- ``hipMemcpyAsync``
|
||||
|
||||
*
|
||||
- ``cudaMemcpyHostToDevice``
|
||||
- ``hipMemcpyHostToDevice``
|
||||
|
||||
*
|
||||
- ``cudaEventRecord``
|
||||
- ``hipEventRecord``
|
||||
|
||||
*
|
||||
- ``cudaEventSynchronize``
|
||||
- ``hipEventSynchronize``
|
||||
|
||||
*
|
||||
- ``cudaEventElapsedTime``
|
||||
- ``hipEventElapsedTime``
|
||||
|
||||
*
|
||||
- ``cudaFree``
|
||||
- ``hipFree``
|
||||
|
||||
*
|
||||
- ``cudaEventDestroy``
|
||||
- ``hipEventDestroy``
|
||||
|
||||
*
|
||||
- ``cudaStreamDestroy``
|
||||
- ``hipStreamDestroy``
|
||||
|
||||
In summary, this comparison highlights the primary differences between CUDA and HIP APIs.
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,171 +0,0 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the C++ support of the HIP ecosystem
|
||||
ROCm software.
|
||||
:keywords: AMD, ROCm, HIP, C++
|
||||
|
||||
*******************************************************************************
|
||||
C++ language support
|
||||
*******************************************************************************
|
||||
|
||||
The ROCm platform enables the power of combined C++ and HIP (Heterogeneous-computing
|
||||
Interface for Portability) code. This code is compiled with a ``clang`` or ``clang++``
|
||||
compiler. The official compilers support the HIP platform, or you can use the
|
||||
``amdclang`` or ``amdclang++`` included in the ROCm installation, which are a wrapper for
|
||||
the official versions.
|
||||
|
||||
The source code is compiled according to the ``C++03``, ``C++11``, ``C++14``, ``C++17``,
|
||||
and ``C++20`` standards, along with HIP-specific extensions, but is subject to
|
||||
restrictions. The key restriction is the reduced support of standard library in device
|
||||
code. This is due to the fact that by default a function is considered to run on host,
|
||||
except for ``constexpr`` functions, which can run on host and device as well.
|
||||
|
||||
.. _language_modern_cpp_support:
|
||||
|
||||
Modern C++ support
|
||||
===============================================================================
|
||||
|
||||
C++ is considered a modern programming language as of C++11. This section describes how
|
||||
HIP supports these new C++ features.
|
||||
|
||||
C++11 support
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
The C++11 standard introduced many new features. These features are supported in HIP host
|
||||
code, with some notable omissions on the device side. The rule of thumb here is that
|
||||
``constexpr`` functions work on device, the rest doesn't. This means that some important
|
||||
functionality like ``std::function`` is missing on the device, but unfortunately the
|
||||
standard library wasn't designed with HIP in mind, which means that the support is in a
|
||||
state of "works as-is".
|
||||
|
||||
Certain features have restrictions and clarifications. For example, any functions using
|
||||
the ``constexpr`` qualifier or the new ``initializer lists``, ``std::move`` or
|
||||
``std::forward`` features are implicitly considered to have the ``__host__`` and
|
||||
``__device__`` execution space specifier. Also, ``constexpr`` variables that are static
|
||||
members or namespace scoped can be used from both host and device, but only for read
|
||||
access. Dereferencing a static ``constexpr`` outside its specified execution space causes
|
||||
an error.
|
||||
|
||||
Lambdas are supported, but there are some extensions and restrictions on their usage. For
|
||||
more information, see the `Extended lambdas`_ section below.
|
||||
|
||||
C++14 support
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
The C++14 language features are supported.
|
||||
|
||||
C++17 support
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
All C++17 language features are supported.
|
||||
|
||||
C++20 support
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
All C++20 language features are supported, but extensions and restrictions apply. C++20
|
||||
introduced coroutines and modules, which fundamentally changed how programs are written.
|
||||
HIP doesn't support these features. However, ``consteval`` functions can be called from
|
||||
host and device, even if specified for host use only.
|
||||
|
||||
The three-way comparison operator (spaceship operator ``<=>``) works with host and device
|
||||
code.
|
||||
|
||||
.. _language_restrictions:
|
||||
|
||||
Extensions and restrictions
|
||||
===============================================================================
|
||||
|
||||
In addition to the deviations from the standard, there are some general extensions and
|
||||
restrictions to consider.
|
||||
|
||||
Global functions
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Functions that serve as an entry point for device execution are called kernels and are
|
||||
specified with the ``__global__`` qualifier. To call a kernel function, use the triple
|
||||
chevron operator: ``<<< >>>``. Kernel functions must have a ``void`` return type. These
|
||||
functions can't:
|
||||
|
||||
* have a ``constexpr`` specifier
|
||||
* have a parameter of type ``std::initializer_list`` or ``va_list``
|
||||
* use an rvalue reference as a parameter.
|
||||
* use parameters having different sizes in host and device code, e.g. long double arguments, or structs containing long double members.
|
||||
* use struct-type arguments which have different layout in host and device code.
|
||||
|
||||
Kernels can have variadic template parameters, but only one parameter pack, which must be
|
||||
the last item in the template parameter list.
|
||||
|
||||
Device space memory specifiers
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
HIP includes device space memory specifiers to indicate whether a variable is allocated
|
||||
in host or device memory and how its memory should be allocated. HIP supports the
|
||||
``__device__``, ``__shared__``, ``__managed__``, and ``__constant__`` specifiers.
|
||||
|
||||
The ``__device__`` and ``__constant__`` specifiers define global variables, which are
|
||||
allocated within global memory on the HIP devices. The only difference is that
|
||||
``__constant__`` variables can't be changed after allocation. The ``__shared__``
|
||||
specifier allocates the variable within shared memory, which is available for all threads
|
||||
in a block.
|
||||
|
||||
The ``__managed__`` variable specifier creates global variables that are initially
|
||||
undefined and unaddressed within the global symbol table. The HIP runtime allocates
|
||||
managed memory and defines the symbol when it loads the device binary. A managed variable
|
||||
can be accessed in both device and host code.
|
||||
|
||||
It's important to know where a variable is stored because it is only available from
|
||||
certain locations. Generally, variables allocated in the host memory are not accessible
|
||||
from the device code, while variables allocated in the device memory are not directly
|
||||
accessible from the host code. Dereferencing a pointer to device memory on the host
|
||||
results in a segmentation fault. Accessing device variables in host code should be done
|
||||
through kernel execution or HIP functions like ``hipMemCpyToSymbol``.
|
||||
|
||||
Exception handling
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
An important difference between the host and device code is exception handling. In device
|
||||
code, this control flow isn't available due to the hardware architecture. The device
|
||||
code must use return codes to handle errors.
|
||||
|
||||
Kernel parameters
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
There are some restrictions on kernel function parameters. They cannot be passed by
|
||||
reference, because these functions are called from the host but run on the device. Also,
|
||||
a variable number of arguments is not allowed.
|
||||
|
||||
Classes
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Classes work on both the host and device side, but there are some constraints. The
|
||||
``static`` member functions can't be ``__global__``. ``Virtual`` member functions work,
|
||||
but a ``virtual`` function must not be called from the host if the parent object was
|
||||
created on the device, or the other way around, because this behavior is undefined.
|
||||
Another minor restriction is that ``__device__`` variables, that are global scoped must
|
||||
have trivial constructors.
|
||||
|
||||
Polymorphic function wrappers
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
HIP doesn't support the polymorphic function wrapper ``std::function``, which was
|
||||
introduced in C++11.
|
||||
|
||||
Extended lambdas
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
HIP supports Lambdas, which by default work as expected.
|
||||
|
||||
Lambdas have implicit host device attributes. This means that they can be executed by
|
||||
both host and device code, and works the way you would expect. To make a lambda callable
|
||||
only by host or device code, users can add ``__host__`` or ``__device__`` attribute. The
|
||||
only restriction is that host variables can only be accessed through copy on the device.
|
||||
Accessing through reference will cause undefined behavior.
|
||||
|
||||
Inline namespaces
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Inline namespaces are supported, but with a few exceptions. The following entities can't
|
||||
be declared in namespace scope within an inline unnamed namespace:
|
||||
|
||||
* ``__managed__``, ``__device__``, ``__shared__`` and ``__constant__`` variables
|
||||
* ``__global__`` function and function templates
|
||||
* variables with surface or texture type
|
||||
@@ -12,162 +12,38 @@ on AMD platform, which are grouped by functionality.
|
||||
GPU isolation variables
|
||||
================================================================================
|
||||
|
||||
The GPU isolation environment variables in HIP are collected in the next table.
|
||||
The GPU isolation environment variables in HIP are collected in the following table.
|
||||
For more information, check :doc:`GPU isolation page <rocm:conceptual/gpu-isolation>`.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 70,30
|
||||
|
||||
* - **Environment variable**
|
||||
- **Value**
|
||||
|
||||
* - | ``ROCR_VISIBLE_DEVICES``
|
||||
| A list of device indices or UUIDs that will be exposed to applications.
|
||||
- Example: ``0,GPU-DEADBEEFDEADBEEF``
|
||||
|
||||
* - | ``GPU_DEVICE_ORDINAL``
|
||||
| Devices indices exposed to OpenCL and HIP applications.
|
||||
- Example: ``0,2``
|
||||
|
||||
* - | ``HIP_VISIBLE_DEVICES`` or ``CUDA_VISIBLE_DEVICES``
|
||||
| Device indices exposed to HIP applications.
|
||||
- Example: ``0,2``
|
||||
.. include-table:: data/env_variables_hip.rst
|
||||
:table: hip-env-isolation
|
||||
|
||||
Profiling variables
|
||||
================================================================================
|
||||
|
||||
The profiling environment variables in HIP are collected in the next table. For
|
||||
The profiling environment variables in HIP are collected in the following table. For
|
||||
more information, check :doc:`setting the number of CUs page <rocm:how-to/setting-cus>`.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 70,30
|
||||
|
||||
* - **Environment variable**
|
||||
- **Value**
|
||||
|
||||
* - | ``HSA_CU_MASK``
|
||||
| Sets the mask on a lower level of queue creation in the driver,
|
||||
| this mask will also be set for queues being profiled.
|
||||
- Example: ``1:0-8``
|
||||
|
||||
* - | ``ROC_GLOBAL_CU_MASK``
|
||||
| Sets the mask on queues created by the HIP or the OpenCL runtimes,
|
||||
| this mask will also be set for queues being profiled.
|
||||
- Example: ``0xf``, enables only 4 CUs
|
||||
|
||||
* - | ``HIP_FORCE_QUEUE_PROFILING``
|
||||
| Used to run the app as if it were run in rocprof. Forces command queue
|
||||
| profiling on by default.
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
.. include-table:: data/env_variables_hip.rst
|
||||
:table: hip-env-prof
|
||||
|
||||
Debug variables
|
||||
================================================================================
|
||||
|
||||
The debugging environment variables in HIP are collected in the next table. For
|
||||
The debugging environment variables in HIP are collected in the following table. For
|
||||
more information, check :ref:`debugging_with_hip`.
|
||||
|
||||
.. include:: ../how-to/debugging_env.rst
|
||||
.. include-table:: data/env_variables_hip.rst
|
||||
:table: hip-env-debug
|
||||
|
||||
Memory management related variables
|
||||
================================================================================
|
||||
|
||||
The memory management related environment variables in HIP are collected in the
|
||||
next table.
|
||||
following table.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 35,14,51
|
||||
|
||||
* - **Environment variable**
|
||||
- **Default value**
|
||||
- **Value**
|
||||
|
||||
* - | ``HIP_HIDDEN_FREE_MEM``
|
||||
| Amount of memory to hide from the free memory reported by hipMemGetInfo.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| Unit: megabyte (MB)
|
||||
|
||||
* - | ``HIP_HOST_COHERENT``
|
||||
| Specifies if the memory is coherent between the host and GPU in ``hipHostMalloc``.
|
||||
- ``0``
|
||||
- | 0: Memory is not coherent.
|
||||
| 1: Memory is coherent.
|
||||
| Environment variable has effect, if the following conditions are statisfied:
|
||||
| - One of the ``hipHostMallocDefault``, ``hipHostMallocPortable``, ``hipHostMallocWriteCombined`` or ``hipHostMallocNumaUser`` flag set to 1.
|
||||
| - ``hipHostMallocCoherent``, ``hipHostMallocNonCoherent`` and ``hipHostMallocMapped`` flags set to 0.
|
||||
|
||||
* - | ``HIP_INITIAL_DM_SIZE``
|
||||
| Set initial heap size for device malloc.
|
||||
- ``8388608``
|
||||
- | Unit: Byte
|
||||
| The default value corresponds to 8 MB.
|
||||
|
||||
* - | ``HIP_MEM_POOL_SUPPORT``
|
||||
| Enables memory pool support in HIP.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``HIP_MEM_POOL_USE_VM``
|
||||
| Enables memory pool support in HIP.
|
||||
- | ``0``: other OS
|
||||
| ``1``: Windows
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``HIP_VMEM_MANAGE_SUPPORT``
|
||||
| Virtual Memory Management Support.
|
||||
- ``1``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``GPU_MAX_HEAP_SIZE``
|
||||
| Set maximum size of the GPU heap to % of board memory.
|
||||
- ``100``
|
||||
- | Unit: Percentage
|
||||
|
||||
* - | ``GPU_MAX_REMOTE_MEM_SIZE``
|
||||
| Maximum size that allows device memory substitution with system.
|
||||
- ``2``
|
||||
- | Unit: kilobyte (KB)
|
||||
|
||||
* - | ``GPU_NUM_MEM_DEPENDENCY``
|
||||
| Number of memory objects for dependency tracking.
|
||||
- ``256``
|
||||
-
|
||||
|
||||
* - | ``GPU_STREAMOPS_CP_WAIT``
|
||||
| Force the stream memory operation to wait on CP.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``HSA_LOCAL_MEMORY_ENABLE``
|
||||
| Enable HSA device local memory usage.
|
||||
- ``1``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``PAL_ALWAYS_RESIDENT``
|
||||
| Force memory resources to become resident at allocation time.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
|
||||
* - | ``PAL_PREPINNED_MEMORY_SIZE``
|
||||
| Size of prepinned memory.
|
||||
- ``64``
|
||||
- | Unit: kilobyte (KB)
|
||||
|
||||
* - | ``REMOTE_ALLOC``
|
||||
| Use remote memory for the global heap allocation.
|
||||
- ``0``
|
||||
- | 0: Disable
|
||||
| 1: Enable
|
||||
.. include-table:: data/env_variables_hip.rst
|
||||
:table: hip-env-memory
|
||||
|
||||
Other useful variables
|
||||
================================================================================
|
||||
@@ -175,15 +51,5 @@ Other useful variables
|
||||
The following table lists environment variables that are useful but relate to
|
||||
different features.
|
||||
|
||||
.. list-table::
|
||||
:header-rows: 1
|
||||
:widths: 35,14,51
|
||||
|
||||
* - **Environment variable**
|
||||
- **Default value**
|
||||
- **Value**
|
||||
|
||||
* - | ``HIPRTC_COMPILE_OPTIONS_APPEND``
|
||||
| Sets compile options needed for ``hiprtc`` compilation.
|
||||
- None
|
||||
- ``--gpu-architecture=gfx906:sramecc+:xnack``, ``-fgpu-rdc``
|
||||
.. include-table:: data/env_variables_hip.rst
|
||||
:table: hip-env-other
|
||||
|
||||
+3
-3
@@ -1,6 +1,6 @@
|
||||
.. meta::
|
||||
:description: This chapter lists types and device API wrappers related to the
|
||||
Cooperative Group feature. Programmers can directly use these
|
||||
:description: This chapter lists types and device API wrappers related to the
|
||||
Cooperative Group feature. Programmers can directly use these
|
||||
API features in their kernels.
|
||||
:keywords: AMD, ROCm, HIP, cooperative groups
|
||||
|
||||
@@ -42,7 +42,7 @@ The following cooperative groups classes can be used on the device side.
|
||||
|
||||
.. doxygenclass:: cooperative_groups::multi_grid_group
|
||||
:members:
|
||||
|
||||
|
||||
.. _thread_block_tile_ref:
|
||||
|
||||
.. doxygenclass:: cooperative_groups::thread_block_tile
|
||||
|
||||
@@ -1,38 +0,0 @@
|
||||
# Table comparing syntax for different compute APIs
|
||||
|
||||
|Term|CUDA|HIP|OpenCL|
|
||||
|---|---|---|---|
|
||||
|Device|`int deviceId`|`int deviceId`|`cl_device`|
|
||||
|Queue|`cudaStream_t`|`hipStream_t`|`cl_command_queue`|
|
||||
|Event|`cudaEvent_t`|`hipEvent_t`|`cl_event`|
|
||||
|Memory|`void *`|`void *`|`cl_mem`|
|
||||
|||||
|
||||
| |grid|grid|NDRange|
|
||||
| |block|block|work-group|
|
||||
| |thread|thread|work-item|
|
||||
| |warp|warp|sub-group|
|
||||
|||||
|
||||
|Thread-<br>index | `threadIdx.x` | `threadIdx.x` | `get_local_id(0)` |
|
||||
|Block-<br>index | `blockIdx.x` | `blockIdx.x` | `get_group_id(0)` |
|
||||
|Block-<br>dim | `blockDim.x` | `blockDim.x` | `get_local_size(0)` |
|
||||
|Grid-dim | `gridDim.x` | `gridDim.x` | `get_num_groups(0)` |
|
||||
|||||
|
||||
|Device Kernel|`__global__`|`__global__`|`__kernel`|
|
||||
|Device Function|`__device__`|`__device__`|Implied in device compilation|
|
||||
|Host Function|`__host_` (default)|`__host_` (default)|Implied in host compilation|
|
||||
|Host + Device Function|`__host__` `__device__`|`__host__` `__device__`| No equivalent|
|
||||
|Kernel Launch|`<<< >>>`|`hipLaunchKernel`/`hipLaunchKernelGGL`/`<<< >>>`|`clEnqueueNDRangeKernel`|
|
||||
||||||
|
||||
|Global Memory|`__global__`|`__global__`|`__global`|
|
||||
|Group Memory|`__shared__`|`__shared__`|`__local`|
|
||||
|Constant|`__constant__`|`__constant__`|`__constant`|
|
||||
||||||
|
||||
||`__syncthreads`|`__syncthreads`|`barrier(CLK_LOCAL_MEMFENCE)`|
|
||||
|Atomic Builtins|`atomicAdd`|`atomicAdd`|`atomic_add`|
|
||||
|Precise Math|`cos(f)`|`cos(f)`|`cos(f)`|
|
||||
|Fast Math|`__cos(f)`|`__cos(f)`|`native_cos(f)`|
|
||||
|Vector|`float4`|`float4`|`float4`|
|
||||
|
||||
## Notes
|
||||
|
||||
The indexing functions (starting with `thread-index`) show the terminology for a 1D grid. Some APIs use reverse order of `xyz` / 012 indexing for 3D grids.
|
||||
@@ -49,12 +49,15 @@ subtrees:
|
||||
- file: how-to/hip_runtime_api/memory_management/virtual_memory
|
||||
- file: how-to/hip_runtime_api/memory_management/stream_ordered_allocator
|
||||
- file: how-to/hip_runtime_api/error_handling
|
||||
- file: how-to/hip_runtime_api/cooperative_groups
|
||||
- file: how-to/hip_runtime_api/hipgraph
|
||||
- file: how-to/hip_runtime_api/call_stack
|
||||
- file: how-to/hip_runtime_api/asynchronous
|
||||
- file: how-to/hip_runtime_api/hipgraph
|
||||
- file: how-to/hip_runtime_api/cooperative_groups
|
||||
- file: how-to/hip_runtime_api/multi_device
|
||||
- file: how-to/hip_runtime_api/opengl_interop
|
||||
- file: how-to/hip_runtime_api/external_interop
|
||||
- file: how-to/hip_cpp_language_extensions
|
||||
- file: how-to/kernel_language_cpp_support
|
||||
- file: how-to/hip_porting_guide
|
||||
- file: how-to/hip_porting_driver_api
|
||||
- file: how-to/hip_rtc
|
||||
@@ -106,14 +109,9 @@ subtrees:
|
||||
- file: doxygen/html/annotated
|
||||
- file: doxygen/html/files
|
||||
- file: reference/virtual_rocr
|
||||
- file: reference/cpp_language_extensions
|
||||
title: C++ language extensions
|
||||
- file: reference/cpp_language_support
|
||||
title: C++ language support
|
||||
- file: reference/math_api
|
||||
- file: reference/env_variables
|
||||
- file: reference/terms
|
||||
title: Comparing syntax for different APIs
|
||||
- file: reference/api_syntax
|
||||
- file: reference/deprecated_api_list
|
||||
title: List of deprecated APIs
|
||||
- file: reference/fp8_numbers
|
||||
@@ -124,10 +122,8 @@ subtrees:
|
||||
entries:
|
||||
- url: https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic
|
||||
title: HIP basic examples
|
||||
- url: https://github.com/ROCm/HIP-Examples
|
||||
- url: https://github.com/ROCm/rocm-examples
|
||||
title: HIP examples
|
||||
- url: https://github.com/ROCm/hip-tests/tree/develop/samples
|
||||
title: HIP test samples
|
||||
- file: tutorial/saxpy
|
||||
- file: tutorial/reduction
|
||||
- file: tutorial/cooperative_groups_tutorial
|
||||
|
||||
@@ -1,2 +1,2 @@
|
||||
rocm-docs-core[api_reference]==1.10.0
|
||||
rocm-docs-core[api_reference]==1.15.0
|
||||
sphinxcontrib.doxylink
|
||||
|
||||
@@ -8,6 +8,13 @@ accessible-pygments==0.0.5
|
||||
# via pydata-sphinx-theme
|
||||
alabaster==1.0.0
|
||||
# via sphinx
|
||||
asttokens==3.0.0
|
||||
# via stack-data
|
||||
attrs==24.3.0
|
||||
# via
|
||||
# jsonschema
|
||||
# jupyter-cache
|
||||
# referencing
|
||||
babel==2.16.0
|
||||
# via
|
||||
# pydata-sphinx-theme
|
||||
@@ -28,11 +35,18 @@ click==8.1.7
|
||||
# via
|
||||
# click-log
|
||||
# doxysphinx
|
||||
# jupyter-cache
|
||||
# sphinx-external-toc
|
||||
click-log==0.4.0
|
||||
# via doxysphinx
|
||||
comm==0.2.2
|
||||
# via ipykernel
|
||||
cryptography==43.0.1
|
||||
# via pyjwt
|
||||
debugpy==1.8.12
|
||||
# via ipykernel
|
||||
decorator==5.1.1
|
||||
# via ipython
|
||||
deprecated==1.2.14
|
||||
# via pygithub
|
||||
docutils==0.21.2
|
||||
@@ -43,20 +57,56 @@ docutils==0.21.2
|
||||
# sphinx
|
||||
doxysphinx==3.3.10
|
||||
# via rocm-docs-core
|
||||
exceptiongroup==1.2.2
|
||||
# via ipython
|
||||
executing==2.1.0
|
||||
# via stack-data
|
||||
fastjsonschema==2.20.0
|
||||
# via rocm-docs-core
|
||||
# via
|
||||
# nbformat
|
||||
# rocm-docs-core
|
||||
gitdb==4.0.11
|
||||
# via gitpython
|
||||
gitpython==3.1.43
|
||||
# via rocm-docs-core
|
||||
greenlet==3.1.1
|
||||
# via sqlalchemy
|
||||
idna==3.8
|
||||
# via requests
|
||||
imagesize==1.4.1
|
||||
# via sphinx
|
||||
importlib-metadata==8.6.1
|
||||
# via
|
||||
# jupyter-cache
|
||||
# myst-nb
|
||||
ipykernel==6.29.5
|
||||
# via myst-nb
|
||||
ipython==8.31.0
|
||||
# via
|
||||
# ipykernel
|
||||
# myst-nb
|
||||
jedi==0.19.2
|
||||
# via ipython
|
||||
jinja2==3.1.4
|
||||
# via
|
||||
# myst-parser
|
||||
# sphinx
|
||||
jsonschema==4.23.0
|
||||
# via nbformat
|
||||
jsonschema-specifications==2024.10.1
|
||||
# via jsonschema
|
||||
jupyter-cache==1.0.1
|
||||
# via myst-nb
|
||||
jupyter-client==8.6.3
|
||||
# via
|
||||
# ipykernel
|
||||
# nbclient
|
||||
jupyter-core==5.7.2
|
||||
# via
|
||||
# ipykernel
|
||||
# jupyter-client
|
||||
# nbclient
|
||||
# nbformat
|
||||
libsass==0.22.0
|
||||
# via doxysphinx
|
||||
lxml==4.9.4
|
||||
@@ -67,20 +117,52 @@ markdown-it-py==3.0.0
|
||||
# myst-parser
|
||||
markupsafe==2.1.5
|
||||
# via jinja2
|
||||
matplotlib-inline==0.1.7
|
||||
# via
|
||||
# ipykernel
|
||||
# ipython
|
||||
mdit-py-plugins==0.4.1
|
||||
# via myst-parser
|
||||
mdurl==0.1.2
|
||||
# via markdown-it-py
|
||||
mpire==2.10.2
|
||||
# via doxysphinx
|
||||
myst-parser==4.0.0
|
||||
myst-nb==1.1.2
|
||||
# via rocm-docs-core
|
||||
myst-parser==4.0.0
|
||||
# via myst-nb
|
||||
nbclient==0.10.2
|
||||
# via
|
||||
# jupyter-cache
|
||||
# myst-nb
|
||||
nbformat==5.10.4
|
||||
# via
|
||||
# jupyter-cache
|
||||
# myst-nb
|
||||
# nbclient
|
||||
nest-asyncio==1.6.0
|
||||
# via ipykernel
|
||||
numpy==1.26.4
|
||||
# via doxysphinx
|
||||
packaging==24.1
|
||||
# via
|
||||
# ipykernel
|
||||
# pydata-sphinx-theme
|
||||
# sphinx
|
||||
parso==0.8.4
|
||||
# via jedi
|
||||
pexpect==4.9.0
|
||||
# via ipython
|
||||
platformdirs==4.3.6
|
||||
# via jupyter-core
|
||||
prompt-toolkit==3.0.50
|
||||
# via ipython
|
||||
psutil==6.1.1
|
||||
# via ipykernel
|
||||
ptyprocess==0.7.0
|
||||
# via pexpect
|
||||
pure-eval==0.2.3
|
||||
# via stack-data
|
||||
pycparser==2.22
|
||||
# via cffi
|
||||
pydata-sphinx-theme==0.15.4
|
||||
@@ -92,6 +174,7 @@ pygithub==2.4.0
|
||||
pygments==2.18.0
|
||||
# via
|
||||
# accessible-pygments
|
||||
# ipython
|
||||
# mpire
|
||||
# pydata-sphinx-theme
|
||||
# sphinx
|
||||
@@ -106,18 +189,34 @@ pyparsing==3.1.4
|
||||
# doxysphinx
|
||||
# sphinxcontrib-doxylink
|
||||
python-dateutil==2.9.0.post0
|
||||
# via sphinxcontrib-doxylink
|
||||
# via
|
||||
# jupyter-client
|
||||
# sphinxcontrib-doxylink
|
||||
pyyaml==6.0.2
|
||||
# via
|
||||
# jupyter-cache
|
||||
# myst-nb
|
||||
# myst-parser
|
||||
# rocm-docs-core
|
||||
# sphinx-external-toc
|
||||
pyzmq==26.2.0
|
||||
# via
|
||||
# ipykernel
|
||||
# jupyter-client
|
||||
referencing==0.36.1
|
||||
# via
|
||||
# jsonschema
|
||||
# jsonschema-specifications
|
||||
requests==2.32.3
|
||||
# via
|
||||
# pygithub
|
||||
# sphinx
|
||||
rocm-docs-core[api-reference]==1.10.0
|
||||
rocm-docs-core[api-reference]==1.15.0
|
||||
# via -r requirements.in
|
||||
rpds-py==0.22.3
|
||||
# via
|
||||
# jsonschema
|
||||
# referencing
|
||||
six==1.16.0
|
||||
# via python-dateutil
|
||||
smmap==5.0.1
|
||||
@@ -129,6 +228,7 @@ soupsieve==2.6
|
||||
sphinx==8.0.2
|
||||
# via
|
||||
# breathe
|
||||
# myst-nb
|
||||
# myst-parser
|
||||
# pydata-sphinx-theme
|
||||
# rocm-docs-core
|
||||
@@ -152,7 +252,7 @@ sphinxcontrib-applehelp==2.0.0
|
||||
# via sphinx
|
||||
sphinxcontrib-devhelp==2.0.0
|
||||
# via sphinx
|
||||
sphinxcontrib-doxylink==1.12.3
|
||||
sphinxcontrib-doxylink==1.12.4
|
||||
# via -r requirements.in
|
||||
sphinxcontrib-htmlhelp==2.1.0
|
||||
# via sphinx
|
||||
@@ -162,17 +262,45 @@ sphinxcontrib-qthelp==2.0.0
|
||||
# via sphinx
|
||||
sphinxcontrib-serializinghtml==2.0.0
|
||||
# via sphinx
|
||||
sqlalchemy==2.0.37
|
||||
# via jupyter-cache
|
||||
stack-data==0.6.3
|
||||
# via ipython
|
||||
tabulate==0.9.0
|
||||
# via jupyter-cache
|
||||
tomli==2.0.1
|
||||
# via sphinx
|
||||
tornado==6.4.2
|
||||
# via
|
||||
# ipykernel
|
||||
# jupyter-client
|
||||
tqdm==4.66.5
|
||||
# via mpire
|
||||
traitlets==5.14.3
|
||||
# via
|
||||
# comm
|
||||
# ipykernel
|
||||
# ipython
|
||||
# jupyter-client
|
||||
# jupyter-core
|
||||
# matplotlib-inline
|
||||
# nbclient
|
||||
# nbformat
|
||||
typing-extensions==4.12.2
|
||||
# via
|
||||
# ipython
|
||||
# myst-nb
|
||||
# pydata-sphinx-theme
|
||||
# pygithub
|
||||
# referencing
|
||||
# sqlalchemy
|
||||
urllib3==2.2.2
|
||||
# via
|
||||
# pygithub
|
||||
# requests
|
||||
wcwidth==0.2.13
|
||||
# via prompt-toolkit
|
||||
wrapt==1.16.0
|
||||
# via deprecated
|
||||
zipp==3.21.0
|
||||
# via importlib-metadata
|
||||
|
||||
@@ -96,5 +96,6 @@ Static libraries
|
||||
ar rcsD libHipDevice.a hipDevice.o
|
||||
hipcc libHipDevice.a test.cpp -fgpu-rdc -o test.out
|
||||
|
||||
For more information, see `HIP samples host functions <https://github.com/ROCm/hip-tests/tree/develop/samples/2_Cookbook/15_static_library/host_functions>`_
|
||||
and `device functions <https://github.com/ROCm/hip-tests/tree/develop/samples/2_Cookbook/15_static_library/device_functions>`_.
|
||||
A full example for this can be found in the ROCm-examples, see the examples for
|
||||
`static host libraries <https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/static_host_library>`_
|
||||
or `static device libraries <https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/static_device_library>`_.
|
||||
|
||||
@@ -95,5 +95,5 @@ language features that are designed to target accelerators, such as:
|
||||
* Math functions that resemble those in ``math.h``, which is included with standard C++ compilers
|
||||
* Built-in functions for accessing specific GPU hardware capabilities
|
||||
|
||||
For further details, check :doc:`C++ language extensions <reference/cpp_language_extensions>`
|
||||
and :doc:`C++ language support <reference/cpp_language_support>`.
|
||||
For further details, check :doc:`HIP C++ language extensions <how-to/hip_cpp_language_extensions>`
|
||||
and :doc:`Kernel language C++ support <how-to/kernel_language_cpp_support>`.
|
||||
|
||||
@@ -401,7 +401,7 @@ typedef struct hipExtent {
|
||||
size_t depth;
|
||||
}hipExtent;
|
||||
/**
|
||||
* HIP position
|
||||
* HIP position
|
||||
*/
|
||||
typedef struct hipPos {
|
||||
size_t x; ///< X coordinate
|
||||
|
||||
@@ -59,7 +59,7 @@ THE SOFTWARE.
|
||||
#define HIP_INTERNAL_EXPORTED_API __attribute__ ((visibility ("default")))
|
||||
#else
|
||||
#define HIP_PUBLIC_API
|
||||
#define HIP_INTERNAL_EXPORTED_API
|
||||
#define HIP_INTERNAL_EXPORTED_API
|
||||
#endif
|
||||
|
||||
#if __HIP_DEVICE_COMPILE__ == 0
|
||||
|
||||
@@ -114,12 +114,15 @@ typedef struct hipDeviceProp_t {
|
||||
int clockRate; ///< Max clock frequency of the multiProcessors in khz.
|
||||
size_t totalConstMem; ///< Size of shared constant memory region on the device
|
||||
///< (in bytes).
|
||||
int major; ///< Major compute capability. On HCC, this is an approximation and features may
|
||||
///< differ from CUDA CC. See the arch feature flags for portable ways to query
|
||||
///< feature caps.
|
||||
int minor; ///< Minor compute capability. On HCC, this is an approximation and features may
|
||||
///< differ from CUDA CC. See the arch feature flags for portable ways to query
|
||||
int major; ///< Major compute capability version. This indicates the core instruction set
|
||||
///< of the GPU architecture. For example, a value of 11 would correspond to
|
||||
///< Navi III (RDNA3). See the arch feature flags for portable ways to query
|
||||
///< feature caps.
|
||||
int minor; ///< Minor compute capability version. This indicates a particular configuration,
|
||||
///< feature set, or variation within the group represented by the major compute
|
||||
///< capability version. For example, different models within the same major version
|
||||
///< might have varying levels of support for certain features or optimizations.
|
||||
///< See the arch feature flags for portable ways to query feature caps.
|
||||
size_t textureAlignment; ///< Alignment requirement for textures
|
||||
size_t texturePitchAlignment; ///< Pitch alignment requirement for texture references bound to
|
||||
int deviceOverlap; ///< Deprecated. Use asyncEngineCount instead
|
||||
@@ -1092,7 +1095,10 @@ typedef enum hipMemAccessFlags {
|
||||
hipMemAccessFlagsProtReadWrite = 3 ///< Set the address range read-write accessible
|
||||
} hipMemAccessFlags;
|
||||
/**
|
||||
* Memory access descriptor
|
||||
* Memory access descriptor structure is used to specify memory access
|
||||
* permissions for a virtual memory region in Virtual Memory Management API.
|
||||
* This structure changes read, and write permissions for
|
||||
* specific memory regions.
|
||||
*/
|
||||
typedef struct hipMemAccessDesc {
|
||||
hipMemLocation location; ///< Location on which the accessibility has to change
|
||||
@@ -6445,7 +6451,7 @@ hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported, #hipErrorOutOfMemory
|
||||
*
|
||||
* @note 3D liner filter isn't supported on GFX90A boards, on which the API @p hipCreateTextureObject will
|
||||
* @note 3D linear filter isn't supported on GFX90A boards, on which the API @p hipCreateTextureObject will
|
||||
* return hipErrorNotSupported.
|
||||
*
|
||||
*/
|
||||
|
||||
Viittaa uudesa ongelmassa
Block a user