diff --git a/projects/hip/.readthedocs.yaml b/projects/hip/.readthedocs.yaml index 4c7088debb..69e5a60f99 100644 --- a/projects/hip/.readthedocs.yaml +++ b/projects/hip/.readthedocs.yaml @@ -15,7 +15,7 @@ python: build: os: ubuntu-22.04 tools: - python: "3.8" + python: "3.10" apt_packages: - "doxygen" - "graphviz" # For dot graphs in doxygen @@ -23,3 +23,8 @@ build: post_checkout: - git clone --depth=1 --single-branch --branch rocdoc-195 https://github.com/StreamHPC/llvm-project.git ../llvm-project - git clone --depth=1 --single-branch --branch develop https://github.com/ROCm/clr.git ../clr + - git clone --depth=1 --single-branch --branch master https://github.com/ROCm/ROCR-Runtime.git ../ROCR-Runtime + post_build: + - rm -rf ../clr + - rm -rf ../llvm-project + - rm -rf ../ROCR-Runtime diff --git a/projects/hip/LICENSE.txt b/projects/hip/LICENSE.txt index 57378c6698..797310b44b 100644 --- a/projects/hip/LICENSE.txt +++ b/projects/hip/LICENSE.txt @@ -1,4 +1,4 @@ -Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +Copyright (c) 2008 - 2024 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal diff --git a/projects/hip/README.md b/projects/hip/README.md index 077ab96d76..a6e887d478 100644 --- a/projects/hip/README.md +++ b/projects/hip/README.md @@ -35,16 +35,16 @@ HIP releases are typically naming convention for each ROCM release to help diffe ## More Info: - [Installation](docs/install/install.rst) -- [HIP FAQ](docs/old/user_guide/faq.md) -- [HIP Kernel Language](docs/old/reference/kernel_language.rst) -- [HIP Porting Guide](docs/old/user_guide/hip_porting_guide.md) -- [HIP Porting Driver Guide](docs/old/user_guide/hip_porting_driver_api.md) -- [HIP Programming Guide](docs/old/user_guide/programming_manual.md) +- [HIP FAQ](docs/how-to/faq.md) +- [HIP Kernel Language](docs/reference/kernel_language.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 Logging ](docs/how-to/logging.rst) - [Building HIP From Source](docs/install/build.rst) - [HIP Debugging ](docs/how-to/debugging.rst) -- [HIP RTC](docs/old/user_guide/hip_rtc.md) -- [HIP Terminology](docs/old/reference/terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/OpenCL) +- [HIP RTC](docs/how-to/hip_rtc.md) +- [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) @@ -56,12 +56,12 @@ HIP releases are typically naming convention for each ROCM release to help diffe * [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) -- [Developer/CONTRIBUTING Info](contributing.md) +- [Developer/CONTRIBUTING Info](CONTRIBUTING.md) - [Release Notes](RELEASE.md) ## How do I get set up? -See the [Installation](docs/how_to_guides/install.md) notes. +See the [Installation](docs/install/install.rst) notes. ## Simple Example The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree. @@ -87,7 +87,7 @@ hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost); The HIP kernel language defines builtins for determining grid and block coordinates, math functions, short vectors, atomics, and timer functions. -It also specifies additional defines and keywords for function types, address spaces, and optimization controls (See the [HIP Kernel Language](docs/old/reference/kernel_language.rst) for a full description). +It also specifies additional defines and keywords for function types, address spaces, and optimization controls (See the [HIP Kernel Language](docs/reference/kernel_language.rst) for a full description). Here's an example of defining a simple 'vector_square' kernel. @@ -131,7 +131,7 @@ 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/old/user_guide/hip_porting_guide.html#porting-a-new-cuda-project) +* 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 diff --git a/projects/hip/docs/doxygen/Doxyfile b/projects/hip/docs/doxygen/Doxyfile index 4d868978a4..d4bb54bd5d 100644 --- a/projects/hip/docs/doxygen/Doxyfile +++ b/projects/hip/docs/doxygen/Doxyfile @@ -832,7 +832,8 @@ WARN_LOGFILE = INPUT = mainpage.md \ ../../include/hip \ ../../../clr/hipamd/include/hip/amd_detail/amd_hip_gl_interop.h \ - ../../../llvm-project/clang/lib/Headers/__clang_hip_math.h + ../../../llvm-project/clang/lib/Headers/__clang_hip_math.h \ + ../../../ROCR-Runtime/src/inc/hsa_ext_amd.h # This tag can be used to specify the character encoding of the source files # that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses diff --git a/projects/hip/docs/how-to/debugging.rst b/projects/hip/docs/how-to/debugging.rst index 5e6405a844..340c4c8df6 100644 --- a/projects/hip/docs/how-to/debugging.rst +++ b/projects/hip/docs/how-to/debugging.rst @@ -268,85 +268,96 @@ HSA provides environment variables that help analyze issues in drivers or hardwa polling, rather than interrupts. HIP environment variable summary ------------------------------------------------------------------------------------------------ +-------------------------------- Here are some of the more commonly used environment variables: +.. # COMMENT: The following lines define a break for use in the table below. +.. |br| raw:: html + +
+ .. list-table:: * - **Environment variable** - **Default value** - **Usage** - * - AMD_LOG_LEVEL - | Enable HIP log on different Level + * - AMD_LOG_LEVEL + |br| Enable HIP log on different Level - 0 - - 0: Disable log. - | 1: Enable log on error level - | 2: Enable log on warning and below levels - | 0x3: Enable log on information and below levels - | 0x4: Decode and display AQL packets + - 0: Disable log. + |br| 1: Enable log on error level + |br| 2: Enable log on warning and below levels + |br| 0x3: Enable log on information and below levels + |br| 0x4: Decode and display AQL packets * - AMD_LOG_MASK - | Enable HIP log on different Level + |br| Enable HIP log on different Level - 0x7FFFFFFF - 0x1: Log API calls - | 0x02: Kernel and Copy Commands and Barriers - | 0x4: Synchronization and waiting for commands to finish - | 0x8: Enable log on information and below levels - | 0x20: Queue commands and queue contents - | 0x40: Signal creation, allocation, pool - | 0x80: Locks and thread-safety code - | 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 - | 0xFFFFFFFF: Log always even mask flag is zero + |br| 0x02: Kernel and Copy Commands and Barriers + |br| 0x4: Synchronization and waiting for commands to finish + |br| 0x8: Enable log on information and below levels + |br| 0x20: Queue commands and queue contents + |br| 0x40: Signal creation, allocation, pool + |br| 0x80: Locks and thread-safety code + |br| 0x100: Copy debug + |br| 0x200: Detailed copy debug + |br| 0x400: Resource allocation, performance-impacting events + |br| 0x800: Initialization and shutdown + |br| 0x1000: Misc debug, not yet classified + |br| 0x2000: Show raw bytes of AQL packet + |br| 0x4000: Show code creation debug + |br| 0x8000: More detailed command info, including barrier commands + |br| 0x10000: Log message location + |br| 0xFFFFFFFF: Log always even mask flag is zero + * - HIP_LAUNCH_BLOCKING + |br| Used for serialization on kernel execution. + - 0 + - 0: Disable. Kernel executes normally. + |br| 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 + |br| Only devices whose index is present in the sequence are visible to HIP - - 0,1,2: Depending on the number of devices on the system * - GPU_DUMP_CODE_OBJECT - | Dump code object + |br| Dump code object - 0 - 0: Disable - | 1: Enable + |br| 1: Enable * - AMD_SERIALIZE_KERNEL - | Serialize kernel enqueue + |br| Serialize kernel enqueue - 0 - 1: Wait for completion before enqueue - | 2: Wait for completion after enqueue - | 3: Both + |br| 2: Wait for completion after enqueue + |br| 3: Both * - AMD_SERIALIZE_COPY - | Serialize copies + |br| Serialize copies - 0 - 1: Wait for completion before enqueue - | 2: Wait for completion after enqueue - | 3: Both + |br| 2: Wait for completion after enqueue + |br| 3: Both * - HIP_HOST_COHERENT - | Coherent memory in hipHostMalloc + |br| Coherent memory in hipHostMalloc - 0 - 0: memory is not coherent between host and GPU - | 1: memory is coherent with host + |br| 1: memory is coherent with host * - AMD_DIRECT_DISPATCH - | Enable direct kernel dispatch (Currently for Linux; under development for Windows) + |br| Enable direct kernel dispatch (Currently for Linux; under development for Windows) - 1 - 0: Disable - | 1: Enable + |br| 1: Enable * - GPU_MAX_HW_QUEUES - | The maximum number of hardware queues allocated per device + |br| 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 @@ -357,26 +368,22 @@ Here are some of the more commonly used environment variables: General debugging tips ====================================================== -* ``gdb --args`` can be used to pass the executable and arguments to gdb. -* Uou can set environment variables (``set env``) from within GDB on Linux (note that this command - doesn't use an equal (=) sign: +* ``gdb --args`` can be used to pass the executable and arguments to ``gdb``. +* You can set environment variables (``set env``) from within GDB on Linux: .. code:: bash (gdb) set env AMD_SERIALIZE_KERNEL 3 -* The GDB backtrace shows a path in the runtime. This is because a fault is caught by the runtime, but - it is generated by an asynchronous command running on the GPU. -* To determine the true location of a fault, you can force the kernels to run synchronously by setting - the environment variables ``AMD_SERIALIZE_KERNEL=3`` and ``AMD_SERIALIZE_COPY=3``. This - forces HIP runtime to wait for the kernel to finish running before retuning. If the fault occurs when - a kernel is running, you can see the code that launched the kernel inside the backtrace. The thread - that's causing the issue is typically the one inside ``libhsa-runtime64.so``. + .. note:: + This ``gdb`` command does not use an equal (=) sign. + +* The GDB backtrace shows a path in the runtime. This is because a fault is caught by the runtime, but it is generated by an asynchronous command running on the GPU. +* To determine the true location of a fault, you can force the kernels to run synchronously by setting the environment variables ``AMD_SERIALIZE_KERNEL=3`` and ``AMD_SERIALIZE_COPY=3``. This forces HIP runtime to wait for the kernel to finish running before retuning. If the fault occurs when a kernel is running, you can see the code that launched the kernel inside the backtrace. The thread that's causing the issue is typically the one inside ``libhsa-runtime64.so``. * VM faults inside kernels can be caused by: * Incorrect code (e.g., a for loop that extends past array boundaries) - * Memory issues, such as invalid kernel arguments (null pointers, unregistered host pointers, bad - pointers) + * Memory issues, such as invalid kernel arguments (null pointers, unregistered host pointers, bad pointers) * Synchronization issues * Compiler issues (incorrect code generation from the compiler) * Runtime issues diff --git a/projects/hip/docs/old/user_guide/faq.md b/projects/hip/docs/how-to/faq.md similarity index 94% rename from projects/hip/docs/old/user_guide/faq.md rename to projects/hip/docs/how-to/faq.md index 7b4d741250..2c6a11c7d8 100644 --- a/projects/hip/docs/old/user_guide/faq.md +++ b/projects/hip/docs/how-to/faq.md @@ -33,7 +33,7 @@ See the [API Support Table](https://github.com/ROCm/HIPIFY/blob/amd-staging/docs - Virtual functions, indirect functions and try/catch (CUDA 4.0) - `__prof_trigger` - PTX assembly (CUDA 4.0). HIP-Clang supports inline GCN assembly. -- Several kernel features are under development. See the {doc}`/old/reference/kernel_language` for more information. +- Several kernel features are under development. See the {doc}`/reference/kernel_language` for more information. ## Is HIP a drop-in replacement for CUDA? @@ -54,7 +54,7 @@ However, we can provide a rough summary of the features included in each CUDA SD - CUDA 6.0 : - Managed memory (under development) - CUDA 6.5 : - - __shfl intriniscs (supported) + - __shfl intrinsic (supported) - CUDA 7.0 : - Per-thread default streams (supported) - C++11 (Hip-Clang supports all of C++11, all of C++14 and some C++17 features) @@ -102,7 +102,7 @@ The tools also struggle with more complex CUDA applications, in particular, thos ## What hardware does HIP support? - For AMD platforms, see the [ROCm documentation](https://github.com/RadeonOpenCompute/ROCm#supported-gpus) for the list of supported platforms. -- For Nvidia platforms, HIP requires Unified Memory and should run on any device supporting CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and Tesla K40. +- For Nvidia platforms, HIP requires unified memory and should run on any device supporting CUDA SDK 6.0 or newer. We have tested the Nvidia Titan and Tesla K40. ## Do HIPIFY tools automatically convert all source code? Typically, HIPIFY tools can automatically convert almost all run-time code. @@ -159,14 +159,14 @@ HIP is a source-portable language that can be compiled to run on either AMD or N ## On HIP-Clang, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ? Yes. HIP generates the object code which conforms to the GCC ABI, and also links with libstdc++. This means you can compile host code with the compiler of your choice and link the generated object code -with GPU code compiled with HIP. Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) and host code (compiled with gcc, icc, or clang). These projects +with GPU code compiled with HIP. Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) and host code (compiled with gcc, icc, or clang). These projects can convert the accelerator code to HIP, compile that code with hipcc, and link with object code from their preferred compiler. ## Can HIP API support C style application? What is the difference between C and C++ ? HIP is C++ runtime API that supports C style applications as well. -Some C style applications (and interfaces to other languages (Fortran, Python)) would call certain HIP APIs but not use kernel programming. -They can be compiled with a C compiler and run correctly, however, small details must be considered in the code. For example, initializtion, as shown in the simple application below, uses HIP structs dim3 with the file name "test.hip.cpp" +Some C style applications (and interfaces to other languages (FORTRAN, Python)) would call certain HIP APIs but not use kernel programming. +They can be compiled with a C compiler and run correctly, however, small details must be considered in the code. For example, initialization, as shown in the simple application below, uses HIP structs dim3 with the file name "test.hip.cpp" ``` #include "hip/hip_runtime_api.h" #include "stdio.h" @@ -187,8 +187,8 @@ $ ./test dim3 grid1; x=1, y=1, z=1 dim3 grid2 = {1,1,1}; x=1, y=1, z=1 ``` -In which "dim3 grid1;" will yield a dim3 grid with all dimentional members x,y,z initalized to 1, as the default constructor behaves that way. -Further, if write, +In which "dim3 grid1;" will yield a dim3 grid with all dimensional members x,y,z initalized to 1, as the default constructor behaves that way. +Further, if written: ``` dim3 grid(2); // yields {2,1,1} dim3 grid(2,3); yields {2,3,1} @@ -201,7 +201,7 @@ $ ./test dim3 grid1; x=646881376, y=21975, z=1517277280 dim3 grid2 = {1,1,1}; x=1, y=1, z=1 ``` -In which "dim3 grid;" does not imply any initialization, no constructor is called, and dimentional values x,y,z of grid are undefined. +In which "dim3 grid;" does not imply any initialization, no constructor is called, and dimensional values x,y,z of grid are undefined. NOTE: To get the C++ default behavior, C programmers must additionally specify the right-hand side as shown below, ``` dim3 grid = {1,1,1}; // initialized as in C++ @@ -245,7 +245,7 @@ If platform portability is important, use #ifdef __HIP_PLATFORM_NVIDIA__ to guar ## How do I trace HIP application flow? See {doc}`/how-to/logging` for more information. -## What is maximum limit of kernel launching parameter? +## What are the maximum limits of kernel launch parameters? Product of block.x, block.y, and block.z should be less than 1024. Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32, so gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32. @@ -256,7 +256,7 @@ __shfl_*_sync is not supported on HIP but for nvcc path CUDA 9.0 and above all s The compiler defines the `__HIP_DEVICE_COMPILE__` macro only when compiling the code for the GPU. It could be used to guard code that is specific to the host or the GPU. ## Why _OpenMP is undefined when compiling with -fopenmp? -When compiling an OpenMP source file with `hipcc -fopenmp`, the compiler may generate error if there is a reference to the `_OPENMP` macro. This is due to a limitation in hipcc that treats any source file type (e.g., `.cpp`) as an HIP translation unit leading to some conflicts with the OpenMP language switch. If the OpenMP source file doesn't contain any HIP language construct, you could workaround this issue by adding the `-x c++` switch to force the compiler to treat the file as regular C++. Another approach would be to guard the OpenMP code with `#ifdef _OPENMP` so that the code block is disabled when compiling for the GPU. The `__HIP_DEVICE_COMPILE__` macro defined by the HIP compiler when compiling GPU code could also be used for guarding code paths specific to the host or the GPU. +When compiling an OpenMP source file with `hipcc -fopenmp`, the compiler may generate error if there is a reference to the `_OPENMP` macro. This is due to a limitation in hipcc that treats any source file type (for example `.cpp`) as an HIP translation unit leading to some conflicts with the OpenMP language switch. If the OpenMP source file doesn't contain any HIP language constructs you could work around this issue by adding the `-x c++` switch to force the compiler to treat the file as regular C++. Another approach would be to guard the OpenMP code with `#ifdef _OPENMP` so that the code block is disabled when compiling for the GPU. The `__HIP_DEVICE_COMPILE__` macro defined by the HIP compiler when compiling GPU code could also be used for guarding code paths specific to the host or the GPU. ## Does the HIP-Clang compiler support extern shared declarations? diff --git a/projects/hip/docs/old/user_guide/hip_porting_driver_api.md b/projects/hip/docs/how-to/hip_porting_driver_api.md similarity index 100% rename from projects/hip/docs/old/user_guide/hip_porting_driver_api.md rename to projects/hip/docs/how-to/hip_porting_driver_api.md diff --git a/projects/hip/docs/old/user_guide/hip_porting_guide.md b/projects/hip/docs/how-to/hip_porting_guide.md similarity index 100% rename from projects/hip/docs/old/user_guide/hip_porting_guide.md rename to projects/hip/docs/how-to/hip_porting_guide.md diff --git a/projects/hip/docs/old/user_guide/hip_rtc.md b/projects/hip/docs/how-to/hip_rtc.md similarity index 95% rename from projects/hip/docs/old/user_guide/hip_rtc.md rename to projects/hip/docs/how-to/hip_rtc.md index 64d05ae759..bd22beeebf 100644 --- a/projects/hip/docs/old/user_guide/hip_rtc.md +++ b/projects/hip/docs/how-to/hip_rtc.md @@ -1,18 +1,15 @@ -# HIP RTC Programming Guide +# Programming for HIP Runtime Compiler (RTC) -## HIP RTC lib -HIP allows you to compile kernels at runtime with its ```hiprtc*``` APIs. -Kernels can be stored as a text string and can be passed on to HIPRTC APIs alongside options to guide the compilation. +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: - - This library can be used on systems without HIP install nor AMD GPU driver installed at all (offline compilation). Therefore, it does not depend on any HIP runtime library. + - 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 COMGr. You may try to statically link COMGr into HIPRTC to avoid any ambiguity. - Developers can decide to bundle this library with their application. -## Compile APIs - -#### Example +## Example To use HIPRTC functionality, HIPRTC header needs to be included first. ```#include ``` @@ -34,12 +31,12 @@ R"( Now to compile this kernel, it needs to be associated with hiprtcProgram type, which is done by declaring ```hiprtcProgram prog;``` and associating the string of kernel with this program: ```cpp -hiprtcCreateProgram(&prog, // HIPRTC program handle - kernel_source, // HIP kernel source string - "vector_add.cpp", // Name of the HIP program, can be null or an empty string - 0, // Number of headers - NULL, // Header sources - NULL); // Name of header files +hiprtcCreateProgram(&prog, // HIPRTC program + kernel, // kernel string + "gpu_kernel.cu", // Name of the file + num_headers, // Number of headers + &header_sources[0], // Header sources + &header_names[0]); // Name of header files ``` hiprtcCreateProgram API also allows you to add headers which can be included in your rtc program. @@ -222,14 +219,14 @@ int main() { } ``` -#### HIPRTC specific options +## HIPRTC specific options HIPRTC provides a few HIPRTC specific flags - ```--gpu-architecture``` : This flag can guide the code object generation for a specific gpu arch. Example: ```--gpu-architecture=gfx906:sramecc+:xnack-```, its equivalent to ```--offload-arch```. - This option is compulsory if compilation is done on a system without AMD GPUs supported by HIP runtime. - Otherwise, HIPRTC will load the hip runtime and gather the current device and its architecture info and use it as option. - ```-fgpu-rdc``` : This flag when provided during the hiprtcCompileProgram generates the bitcode (HIPRTC doesn't convert this bitcode into ISA and binary). This bitcode can later be fetched using hiprtcGetBitcode and hiprtcGetBitcodeSize APIs. -#### Bitcode +### Bitcode In the usual scenario, the kernel associated with hiprtcProgram is compiled into the binary which can be loaded and run. However, if -fpu-rdc option is provided in the compile options, HIPRTC calls comgr and generates only the LLVM bitcode. It doesn't convert this bitcode to ISA and generate the final binary. ```cpp std::string sarg = std::string("-fgpu-rdc"); @@ -249,7 +246,7 @@ vector kernel_bitcode(bitCodeSize); hiprtcGetBitcode(prog, kernel_bitcode.data()); ``` -#### CU Mode vs WGP mode +### CU Mode vs WGP mode AMD GPUs consist of an array of workgroup processors, each built with 2 compute units (CUs) capable of executing SIMD32. All the CUs inside a workgroup processor use local data share (LDS). @@ -263,10 +260,9 @@ HIPRTC assumes **WGP mode by default** for gfx10+. This can be overridden by pas ## Linker APIs -#### Introduction The bitcode generated using the HIPRTC Bitcode APIs can be loaded using hipModule APIs and also can be linked with other generated bitcodes with appropriate linker flags using the HIPRTC linker APIs. This also provides more flexibility and optimizations to the applications who want to generate the binary dynamically according to their needs. The input bitcodes can be generated only for a specific architecture or it can be a bundled bitcode which is generated for multiple architectures. -#### Example +### Example Firstly, HIPRTC link instance or a pending linker invocation must be created using hiprtcLinkCreate, with the appropriate linker options provided. ```cpp hiprtcLinkCreate( num_options, // number of options @@ -314,7 +310,7 @@ hiprtcLinkDestroy(rtc_link_state); ``` - The correct sequence of calls is : hiprtcLinkCreate, hiprtcLinkAddData or hiprtcLinkAddFile, hiprtcLinkComplete, hiprtcModuleLoadData, hiprtcLinkDestroy. -#### Input Types +### Input Types HIPRTC provides hiprtcJITInputType enumeration type which defines the input types accepted by the Linker APIs. Here are the enum values of hiprtcJITInputType. However only the input types HIPRTC_JIT_INPUT_LLVM_BITCODE, HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE and HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE are supported currently. HIPRTC_JIT_INPUT_LLVM_BITCODE can be used to load both LLVM bitcode or LLVM IR assembly code. However, HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE and HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE are only for bundled bitcode and archive of bundled bitcode. @@ -333,7 +329,7 @@ HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE = 102, HIPRTC_JIT_NUM_INPUT_TYPES = (HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES + 3) ``` -#### Backward Compatibility of LLVM Bitcode/IR +### Backward Compatibility of LLVM Bitcode/IR For HIP applications utilizing HIPRTC to compile LLVM bitcode/IR, compatibility is assured only when the ROCm or HIP SDK version used for generating the LLVM bitcode/IR matches the version used during the runtime compilation. When an application requires the ingestion of bitcode/IR not derived from the currently installed AMD compiler, it must run with HIPRTC and COMgr dynamic libraries that are compatible with the version of the bitcode/IR. @@ -341,7 +337,7 @@ COMgr, a shared library, incorporates the LLVM/Clang compiler that HIPRTC relies To ensure smooth operation and compatibility, an application may choose to ship the specific versions of HIPRTC and COMgr dynamic libraries, or it may opt to clearly specify the version requirements and dependencies. This approach guarantees that the application can correctly compile the specified version of bitcode/IR. -#### Link Options +### Link Options - `HIPRTC_JIT_IR_TO_ISA_OPT_EXT` - AMD Only. Options to be passed on to link step of compiler by `hiprtcLinkCreate`. - `HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT` - AMD Only. Count of options passed on to link step of compiler. @@ -396,12 +392,12 @@ HIPRTC mangles the ```__global__``` function names and names of ```__device__``` The two APIs hiprtcAddNameExpression and hiprtcGetLoweredName provide this functionality. First, a 'name expression' string denoting the address for the ```__global__``` function or ```__device__/__constant__``` variable is provided to hiprtcAddNameExpression. Then, the program is compiled with hiprtcCompileProgram. During compilation, HIPRTC will parse the name expression string as a C++ constant expression at the end of the user program. Finally, the function hiprtcGetLoweredName is called with the original name expression and it returns a pointer to the lowered name. The lowered name can be used to refer to the kernel or variable in the HIP Runtime API. -#### Note +### Note - The identical name expression string must be provided on a subsequent call to hiprtcGetLoweredName to extract the lowered name. - The correct sequence of calls is : hiprtcAddNameExpression, hiprtcCompileProgram, hiprtcGetLoweredName, hiprtcDestroyProgram. - The lowered names must be fetched using hiprtcGetLoweredName only after the HIPRTC program has been compiled, and before it has been destroyed. -#### Example +### Example kernel containing various definitions ```__global__``` functions/function templates and ```__device__/__constant__``` variables can be stored in a string. ```cpp diff --git a/projects/hip/docs/how-to/logging.rst b/projects/hip/docs/how-to/logging.rst index 2081fb3cf5..d996243000 100644 --- a/projects/hip/docs/how-to/logging.rst +++ b/projects/hip/docs/how-to/logging.rst @@ -4,7 +4,7 @@ :keywords: AMD, ROCm, HIP, logging ********************************************************** -Generate logs +Logging HIP activity ********************************************************** HIP provides a logging mechanism that allows you to trace HIP API and runtime codes when running a diff --git a/projects/hip/docs/how-to/performance_guidelines.rst b/projects/hip/docs/how-to/performance_guidelines.rst new file mode 100644 index 0000000000..aa8bcb1fce --- /dev/null +++ b/projects/hip/docs/how-to/performance_guidelines.rst @@ -0,0 +1,301 @@ +.. meta:: + :description: This chapter describes a set of best practices designed to help developers optimize the performance of HIP-capable GPU architectures. + :keywords: AMD, ROCm, HIP, CUDA, performance, guidelines + +******************************************************************************* +Performance Guidelines +******************************************************************************* + +The AMD HIP Performance Guidelines are a set of best practices designed to help +developers optimize the performance of AMD GPUs. They cover established +parallelization and optimization techniques, coding metaphors, and idioms that +can greatly simplify programming for HIP-capable GPU architectures. + +By following four main cornerstones, we can exploit the performance +optimization potential of HIP. + +- parallel execution +- memory usage optimization +- optimization for maximum throughput +- minimizing memory thrashing + +In the following chapters, we will show you their benefits and how to use them +effectively. + +.. _parallel execution: + +Parallel execution +================== + +For optimal use, the application should reveal and efficiently imply as much +parallelism as possible to keep all system components active. + +Application level +----------------- + +The application should optimize parallel execution across the host and devices +using asynchronous calls and streams. Workloads should be assigned based on +efficiency: serial to the host, parallel to the devices. + +For parallel workloads, when threads need to synchronize to share data, if they +belong to the same block, they should use ``__syncthreads()`` (see: +:ref:`synchronization functions`) within the same kernel invocation. If they +belong to different blocks, they must use global memory with two separate +kernel invocations. The latter should be minimized as it adds overhead. + +Device level +------------ + +Device-level optimization primarily involves maximizing parallel execution +across the multiprocessors of the device. This can be achieved by executing +multiple kernels concurrently on a device. The management of these kernels is +facilitated by streams, which allow for the overlapping of computation and data +transfers, enhancing performance. The aim is to keep all multiprocessors busy +by executing enough kernels concurrently. However, launching too many kernels +can lead to resource contention, so a balance must be found for optimal +performance. This approach helps in achieving maximum utilization of the +resources of the device. + +Multiprocessor level +-------------------- + +Multiprocessor-level optimization involves maximizing parallel execution within +each multiprocessor on a device. Each multiprocessor can execute a number of +threads concurrently, and the total number of threads that can run in parallel +is determined by the number of concurrent threads each multiprocessor can +handle. + +The key to multiprocessor-level optimization is to efficiently utilize the +various functional units within a multiprocessor. This can be achieved by +ensuring a sufficient number of resident warps, as at every instruction issue +time, a warp scheduler selects an instruction that is ready to execute. This +instruction can be another independent instruction of the same warp, exploiting +:ref:`instruction optimization`, or more commonly an instruction of another warp, +exploiting thread-level parallelism. + +In comparison, device-level optimization focuses on the device as a whole, +aiming to keep all multiprocessors busy by executing enough kernels +concurrently. Both levels of optimization are crucial for achieving maximum +performance. They work together to ensure efficient utilization of the +resources of the GPU, from the individual multiprocessors to the device as a +whole. + +.. _memory optimization: + +Memory optimization +=================== + +The first step in maximizing memory throughput is to minimize low-bandwidth +data transfers. This involves reducing data transfers between the host and the +device, as these have lower bandwidth than transfers between global memory and +the device. + +Additionally, data transfers between global memory and the device should be +minimized by maximizing the use of on-chip memory: shared memory and caches. +Shared memory acts as a user-managed cache, where the application explicitly +allocates and accesses it. A common programming pattern is to stage data from +device memory into shared memory. This involves each thread of a block loading +data from device memory to shared memory, synchronizing with all other threads +of the block, processing the data in shared memory, synchronizing again if +necessary, and writing the results back to device global memory. + +For some applications, a traditional hardware-managed cache is more appropriate +to exploit data locality. On devices of certain compute capabilities, the same +on-chip memory is used for both L1 and shared memory, and the amount dedicated +to each is configurable for each kernel call. + +Finally, the throughput of memory accesses by a kernel can vary significantly +depending on the access pattern for each type of memory. Therefore, the next +step in maximizing memory throughput is to organize memory accesses as +optimally as possible. This is especially important for global memory accesses, +as global memory bandwidth is low compared to available on-chip bandwidths and +arithmetic instruction throughput. Thus, non-optimal global memory accesses +generally have a high impact on performance. + +Data Transfer +------------- + +Applications should aim to minimize data transfers between the host and the +device. This can be achieved by moving more computations from the host to the +device, even if it means running kernels that do not fully utilize the +parallelism for device. Intermediate data structures can be created, used, +and discarded in device memory without being mapped or copied to host memory. + +Batching small transfers into a single large transfer can improve performance +due to the overhead associated with each transfer. On systems with a front-side +bus, using page-locked host memory can enhance data transfer performance. + +When using mapped page-locked memory, there is no need to allocate device +memory or explicitly copy data between device and host memory. Data transfers +occur implicitly each time the kernel accesses the mapped memory. For optimal +performance, these memory accesses should be coalesced, similar to global +memory accesses. + +On integrated systems where device and host memory are physically the same, +any copy operation between host and device memory is unnecessary, and mapped +page-locked memory should be used instead. Applications can check if a device +is integrated by querying the integrated device property. + + +Device Memory Access +-------------------- + +Memory access instructions may be repeated due to the spread of memory +addresses across warp threads. The impact on throughput varies with memory type +and is generally reduced when addresses are more scattered, especially in +global memory. + +Device memory is accessed via 32-, 64-, or 128-byte transactions that must be +naturally aligned. Maximizing memory throughput involves coalescing memory +accesses of threads within a warp into minimal transactions, following optimal +access patterns, using properly sized and aligned data types, and padding data +when necessary. + +Global memory instructions support reading or writing data of specific sizes +(1, 2, 4, 8, or 16 bytes) that are naturally aligned. If the size and alignment +requirements are not met, it leads to multiple instructions, reducing +performance. Therefore, using data types that meet these requirements, ensuring +alignment for structures, and maintaining alignment for all values or arrays is +crucial for correct results and optimal performance. + +Threads often access 2D arrays at an address calculated as +``BaseAddress + xIndex + width * yIndex``. For efficient memory access, the +array and thread block widths should be multiples of the warp size. If the +array width is not a multiple of the warp size, it is usually more efficient to +allocate it with a width rounded up to the nearest multiple and pad the rows +accordingly. + +Local memory is used for certain automatic variables, such as arrays with +non-constant indices, large structures or arrays, and any variable when the +kernel uses more registers than available. Local memory resides in device +memory, leading to high latency and low bandwidth similar to global memory +accesses. However, it is organized for consecutive 32-bit words to be accessed +by consecutive thread IDs, allowing full coalescing when all threads in a warp +access the same relative address. + +Shared memory, located on-chip, provides higher bandwidth and lower latency +than local or global memory. It is divided into banks that can be +simultaneously accessed, boosting bandwidth. However, bank conflicts, where two +addresses fall in the same bank, lead to serialized access and decreased +throughput. Therefore, understanding how memory addresses map to banks and +scheduling requests to minimize conflicts is crucial for optimal performance. + +Constant memory is in device memory and cached in the constant cache. Requests +are split based on different memory addresses, affecting throughput, and are +serviced at the throughput of the constant cache for cache hits, or the +throughput of the device memory otherwise. + +Texture and surface memory are stored in device memory and cached in texture +cache. This setup optimizes 2D spatial locality, leading to better performance +for threads reading close 2D addresses. Reading device memory through texture +or surface fetching can be advantageous, offering higher bandwidth for local +texture fetches or surface reads, offloading addressing calculations, +allowing data broadcasting, and optional conversion of 8-bit and 16-bit integer +input data to 32-bit floating-point values on-the-fly. + +.. _instruction optimization: + +Optimization for maximum instruction throughput +=============================================== + +To maximize instruction throughput: + +- minimize low throughput arithmetic instructions +- minimize divergent warps inflicted by control flow instructions +- minimize the number of instruction as possible +- maximize instruction parallelism + +Arithmetic instructions +----------------------- + +The type and complexity of arithmetic operations can significantly impact the +performance of your application. We are highlighting some hints how to maximize +it. + +Using efficient operations: Some arithmetic operations are more costly than +others. For example, multiplication is typically faster than division, and +integer operations are usually faster than floating-point operations, +especially with double-precision. + +Minimizing low-throughput instructions: This might involve trading precision +for speed when it does not affect the final result. For instance, consider +using single-precision arithmetic instead of double-precision. + +Leverage intrinsic functions: Intrinsic functions are pre-defined functions +available in HIP that can often be executed faster than equivalent arithmetic +operations (subject to some input or accuracy restrictions). They can help +optimize performance by replacing more complex arithmetic operations. + +Avoiding divergent warps: Divergent warps occur when threads within the same +warp follow different execution paths. This can happen due to conditional +statements that lead to different arithmetic operations being performed by +different threads. Divergent warps can significantly reduce instruction +throughput, so try to structure your code to minimize divergence. + +Optimizing memory access: The efficiency of memory access can impact the speed +of arithmetic operations. Coalesced memory access, where threads in a warp +access consecutive memory locations, can improve memory throughput and thus +the speed of arithmetic operations. + +Maximizing instruction parallelism: Some GPU architectures could issue parallel +independent instructions simultaneously, for example integer and floating +point, or two operations with independent inputs and outputs. Mostly this is a +work for compiler, but expressing parallelism in the code explicitly can +improve instructions throughput. + +Control flow instructions +------------------------- + +Flow control instructions (``if``, ``else``, ``for``, ``do``, ``while``, +``break``, ``continue``, ``switch``) can impact instruction throughput by +causing threads within a warp to diverge and follow different execution paths. +To optimize performance, control conditions should be written to minimize +divergent warps. For example, when the control condition depends on +(``threadIdx`` / ``warpSize``), no warp diverges. The compiler may optimize +loops or short if or switch blocks using branch predication, preventing warp +divergence. With branch predication, instructions associated with a false +predicate are scheduled but not executed, avoiding unnecessary operations. + +Synchronization +--------------- + +Synchronization ensures that all threads within a block have completed their +computations and memory accesses before moving forward, which is critical when +threads are dependent on the results of other threads. However, +synchronization can also lead to performance overhead, as it requires threads +to wait, potentially leading to idle GPU resources. + +``__syncthreads()`` is used to synchronize all threads in a block, ensuring +that all threads have reached the same point in the code and that shared memory +is visible to all threads after the point of synchronization. + +An alternative way to synchronize is using streams. Different streams can +execute commands out of order with respect to one another or concurrently. This +allows for more fine-grained control over the execution order of commands, +which can be beneficial in certain scenarios. + +Minimizing memory thrashing +=========================== + +Applications frequently allocating and freeing memory may experience slower +allocation calls over time. This is expected as memory is released back to the +operating system. To optimize performance in such scenarios, consider some +recommendations: + +- avoid allocating all available memory with ``hipMalloc`` / ``hipHostMalloc``, + as this immediately reserves memory and can block other applications from + using it. This could strain the operating system schedulers or even prevent + other applications from running on the same GPU. +- aim to allocate memory in suitably sized blocks early in the lifecycle of the + application and deallocate only when the application no longer needs it. + Minimize the number of ``hipMalloc`` and ``hipFree`` calls in your + application, particularly in areas critical to performance. +- if an application is unable to allocate sufficient device memory, consider + resorting to other memory types such as ``hipHostMalloc`` or + ``hipMallocManaged``. While these may not offer the same performance, they + can allow the application to continue running. +- For supported platforms, ``hipMallocManaged`` allows for oversubscription. + With the right memory advise policies, it can maintain most, if not all, of + the performance of ``hipMalloc``. ``hipMallocManaged`` does not require an + allocation to be resident until it is needed or prefetched, easing the load + on the operating system schedulers and facilitating multi-tenant scenarios. diff --git a/projects/hip/docs/old/user_guide/programming_manual.md b/projects/hip/docs/how-to/programming_manual.md similarity index 91% rename from projects/hip/docs/old/user_guide/programming_manual.md rename to projects/hip/docs/how-to/programming_manual.md index 809d1ddc22..5fe5c18ac7 100644 --- a/projects/hip/docs/old/user_guide/programming_manual.md +++ b/projects/hip/docs/how-to/programming_manual.md @@ -63,7 +63,7 @@ Developers can control the release scope for hipEvents: A stronger system-level fence can be specified when the event is created with hipEventCreateWithFlags: - hipEventReleaseToSystem : Perform a system-scope release operation when the event is recorded. This will make both Coherent and Non-Coherent host memory visible to other agents in the system, but may involve heavyweight operations such as cache flushing. Coherent memory will typically use lighter-weight in-kernel synchronization mechanisms such as an atomic operation and thus does not need to use hipEventReleaseToSystem. -- hipEventDisableTiming: Events created with this flag would not record profiling data and provide best performance if used for synchronization. +- hipEventDisableTiming: Events created with this flag will not record profiling data and provide the best performance if used for synchronization. ### Summary and Recommendations: @@ -74,9 +74,9 @@ A stronger system-level fence can be specified when the event is created with hi Managed memory, including the `__managed__` keyword, is supported in HIP combined host/device compilation, on Linux, not on Windows (under development). Managed memory, via unified memory allocation, allows data be shared and accessible to both the CPU and GPU using a single pointer. -The allocation will be managed by AMD GPU driver using the linux HMM (Heterogeneous Memory Management) mechanism, the user can call managed memory API hipMallocManaged to allocate a large chuch of HMM memory, execute kernels on device and fetch data between the host and device as needed. +The allocation will be managed by AMD GPU driver using the Linux HMM (Heterogeneous Memory Management) mechanism, the user can call managed memory API hipMallocManaged to allocate a large chunk of HMM memory, execute kernels on device and fetch data between the host and device as needed. -In HIP application, It is recommend to do the capability check before calling the managed memory APIs. For example: +In HIP application, it is recommended to do the capability check before calling the managed memory APIs. For example: ``` int managed_memory = 0; @@ -107,7 +107,7 @@ HIP supports Stream Memory Operations to enable direct synchronization between N Note, CPU access to the semaphore's memory requires volatile keyword to disable CPU compiler's optimizations on memory access. For more details, please check the documentation HIP-API.pdf. -Please note, HIP stream does not gurantee concurrency on AMD hardware for the case of multiple (at least 6) long running streams executing concurrently, using hipStreamSynchronize(nullptr) for synchronization. +Please note, HIP stream does not guarantee concurrency on AMD hardware for the case of multiple (at least 6) long-running streams executing concurrently, using hipStreamSynchronize(nullptr) for synchronization. ## Direct Dispatch HIP runtime has Direct Dispatch enabled by default in ROCM 4.4 on Linux. @@ -129,7 +129,7 @@ HIPRTC APIs accept HIP source files in character string format as input paramete For more details on HIPRTC APIs, refer to [HIP Runtime API Reference](https://rocm.docs.amd.com/projects/HIP/en/latest/doxygen/html/index.html). -For Linux developers, the link [here](https://github.com/ROCm/hip-tests/blob/develop/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail HIPRTC programming guide is also available in Github (https://github.com/ROCm/HIP/blob/develop/docs/user_guide/hip_rtc.md). +For Linux developers, the link [here](https://github.com/ROCm/hip-tests/blob/develop/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and a detailed [HIPRTC programming guide](./hip_rtc) is also available. ## HIP Graph HIP graph is supported. For more details, refer to the HIP API Guide. @@ -170,11 +170,11 @@ can be contracted. Tolerance should be used for floating point comparsions. ## Math functions with special rounding modes -Note: Currently, HIP only supports basic math functions with rounding mode rn (round to nearest). HIP does not support basic math functions with rounding modes ru (round up), rd (round down), and rz (round towards zero). +Note: Currently, HIP only supports basic math functions with rounding modern (round to nearest). HIP does not support basic math functions with rounding modes ru (round up), rd (round down), and rz (round towards zero). ## Creating Static Libraries -HIP-Clang supports generating two types of static libraries. The first type of static library does not export device functions, and only exports and launches host functions within the same library. The advantage of this type is the ability to link with a non-hipcc compiler such as gcc. The second type exports device functions to be linked by other code objects. However this requires using hipcc as the linker. +HIP-Clang supports generating two types of static libraries. The first type of static library does not export device functions, and only exports and launches host functions within the same library. The advantage of this type is the ability to link with a non-hipcc compiler such as gcc. The second type exports device functions to be linked by other code objects. However, this requires using hipcc as the linker. In addition, the first type of library contains host objects with device code embedded as fat binaries. It is generated using the flag --emit-static-lib. The second type of library contains relocatable device objects and is generated using ar. @@ -191,4 +191,4 @@ Here is an example to create and use static libraries: hipcc libHipDevice.a test.cpp -fgpu-rdc -o test.out ``` -For more information, please see [HIP samples](https://github.com/ROCm/hip-tests/tree/develop/samples/2_Cookbook/15_static_library/host_functions) and [samples](https://github.com/ROCm/hip-tests/tree/rocm-5.5.x/samples/2_Cookbook/15_static_library/device_functions). +For more information, please 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/rocm-5.5.x/samples/2_Cookbook/15_static_library/device_functions). diff --git a/projects/hip/docs/index.md b/projects/hip/docs/index.md index 571f9c08ad..8f13c546ca 100644 --- a/projects/hip/docs/index.md +++ b/projects/hip/docs/index.md @@ -1,71 +1,63 @@ # HIP documentation -HIP is a C++ runtime API and kernel language that allows developers to create +HIP is a C++ runtime API and kernel language that lets developers create portable applications for AMD and NVIDIA GPUs from single source code. ## Overview ::::{grid} 1 1 2 2 -:gutter: 1 +:gutter: 3 -:::{grid-item-card} Installation +:::{grid-item-card} Install -* [Install HIP](./install/install) -* [Build HIP from source](./install/build) - -::: - -:::{grid-item-card} Reference - -* {doc}`/reference/programming_model` -* {doc}`/doxygen/html/index` -* [Deprecated APIs](./reference/deprecated_api_list) - -::: - -:::{grid-item-card} How-to - -* [Debug with HIP](./how-to/debugging) -* [Generate logs](./how-to/logging) +* [Installing HIP](./install/install) +* [Building HIP from source](./install/build) ::: :::{grid-item-card} Conceptual -* {doc}`/understand/programming_model` +* {doc}`./understand/programming_model` +* {doc}`./understand/programming_model_reference` +* {doc}`./understand/hardware_implementation` ::: -:::: +:::{grid-item-card} How to -## Legacy documentation +* [Programming Manual](./how-to/programming_manual) +* [HIP Porting Guide](./how-to/hip_porting_guide) +* [HIP Porting: Driver API Guide](./how-to/hip_porting_driver_api) +* {doc}`./how-to/hip_rtc` +* {doc}`./how-to/performance_guidelines` +* [Debugging with HIP](./how-to/debugging) +* {doc}`./how-to/logging` +* {doc}`./how-to/faq` -These documents have not yet been ported over to the Diátaxis framework. - -::::{grid} 1 1 2 2 -:gutter: 1 +::: :::{grid-item-card} Reference -* [C++ kernel language](./old/reference/kernel_language) -* [Table Comparing Syntax for Different Compute APIs](./old/reference/terms) +* {doc}`/doxygen/html/index` +* [C++ language extensions](./reference/kernel_language) +* [Comparing Syntax for different APIs](./reference/terms) +* [HSA Runtime API for ROCm](./reference/virtual_rocr) +* [List of deprecated APIs](./reference/deprecated_api_list) ::: -:::{grid-item-card} User Guide +:::{grid-item-card} Tutorial -* [HIP Porting Guide](./old/user_guide/hip_porting_guide) -* [HIP Porting Driver API Guide](./old/user_guide/hip_porting_driver_api) -* [HIP RTC Programming Guide](./old/user_guide/hip_rtc.md) -* [HIP Programming Manual](./old/user_guide/programming_manual.md) -* [Frequently asked questions](./old/user_guide/faq.md) +* [HIP examples](https://github.com/ROCm/HIP-Examples) +* [HIP test samples](https://github.com/ROCm/hip-tests/tree/develop/samples) ::: :::: -We welcome collaboration! If you’d like to contribute to our documentation, you can find instructions -on our {doc}`Contribute to ROCm docs ` page. Known issues are listed on -[GitHub](https://github.com/ROCm/HIP/issues). +Known issues are listed on the [HIP GitHub repository](https://github.com/ROCm/HIP/issues). -If you want to contribute to the HIP project, refer to our [Contributor guidelines](https://github.com/ROCm/HIP/CONTRIBUTING.md). +To contribute features or functions to the HIP project, refer to [Contributing to HIP](https://github.com/ROCm/HIP/blob/develop/CONTRIBUTING.md). +To contribute to the documentation, refer to {doc}`Contributing to ROCm docs ` page. + +You can find licensing information on the [Licensing](https://rocm.docs.amd.com/en/latest/about/license.html) page. diff --git a/projects/hip/docs/install/build.rst b/projects/hip/docs/install/build.rst index d0ce971ad5..9ef966ff60 100644 --- a/projects/hip/docs/install/build.rst +++ b/projects/hip/docs/install/build.rst @@ -5,22 +5,25 @@ Build HIP from source Prerequisites ================================================= -HIP code can be developed either on AMD ROCm platform using HIP-Clang compiler, or a CUDA platform with nvcc installed. -Before build and run HIP, make sure drivers and pre-build packages are installed properly on the platform. +HIP code can be developed either on AMD ROCm platform using HIP-Clang compiler, or a CUDA platform with ``nvcc`` installed. +Before building and running HIP, make sure drivers and prebuilt packages are installed properly on the platform. -You also need to install Python 3, which includes the CppHeaderParser package. -Install Python 3 using the command: +You also need to install Python 3, which includes the ``CppHeaderParser`` package. +Install Python 3 using the following command: .. code:: shell apt-get install python3 -Check and install CppHeaderParser package using the command: +Check and install ``CppHeaderParser`` package using the command: .. code:: shell pip3 install CppHeaderParser + +.. _Building the HIP runtime: + Building the HIP runtime ========================================================== diff --git a/projects/hip/docs/install/install.rst b/projects/hip/docs/install/install.rst index db70407d20..2ef868f10e 100644 --- a/projects/hip/docs/install/install.rst +++ b/projects/hip/docs/install/install.rst @@ -38,8 +38,8 @@ Installation HIP is automatically installed during the ROCm installation. If you haven't yet installed ROCm, you can find installation instructions here: - * :doc:`rocm-install-on-linux:reference/system-requirements` - * :doc:`rocm-install-on-windows:reference/system-requirements` + * :doc:`rocm-install-on-linux:index` + * :doc:`rocm-install-on-windows:index` By default, HIP is installed into ``/opt/rocm/hip``. diff --git a/projects/hip/docs/old/glossary.md b/projects/hip/docs/old/glossary.md deleted file mode 100644 index ac9d24cf1c..0000000000 --- a/projects/hip/docs/old/glossary.md +++ /dev/null @@ -1,23 +0,0 @@ -# Terms used in HIP Documentation - -- host, host cpu : Executes the HIP runtime API and is capable of initiating kernel launches to one or more devices. -- default device : Each host thread maintains a "default device". -Most HIP runtime APIs (including memory allocation, copy commands, kernel launches) do not use accept an explicit device -argument but instead implicitly use the default device. -The default device can be set with hipSetDevice. - -- "active host thread" - the thread which is running the HIP APIs. - -- HIP-Clang - Heterogeneous AMDGPU Compiler, with its capability to compile HIP programs on AMD platform (https://github.com/RadeonOpenCompute/llvm-project). - -- clr - a repository for AMD Common Language Runtime, contains source codes for AMD's compute languages runtimes: HIP and OpenCL�. -clr (https://github.com/ROCm/clr) contains the following three parts, -hipamd: contains implementation of HIP language on AMD platform. -rocclr: contains common runtime used in HIP and OpenCL�, which provides virtual device interfaces that compute runtimes interact with different backends such as ROCr on Linux or PAL on Windows. -opencl: contains implementation of OpenCL� on AMD platform. - -- hipify tools - tools to convert CUDA code to portable C++ code (https://github.com/ROCm/HIPIFY). - -- hipconfig - tool to report various configuration properties of the target platform. - -- nvcc = nvcc compiler, do not capitalize. diff --git a/projects/hip/docs/old/reference/kernel_language.rst b/projects/hip/docs/reference/kernel_language.rst similarity index 99% rename from projects/hip/docs/old/reference/kernel_language.rst rename to projects/hip/docs/reference/kernel_language.rst index ba667ab458..328d517042 100644 --- a/projects/hip/docs/old/reference/kernel_language.rst +++ b/projects/hip/docs/reference/kernel_language.rst @@ -314,8 +314,8 @@ The ``__syncthreads()`` built-in function is supported in HIP. The ``__syncthrea Math functions ==================================================== -HIP-Clang supports a set of math operations that are callable from the device. These are described in -the following sections. +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 in the following sections. Single precision mathematical functions -------------------------------------------------------------------------------------------- @@ -422,7 +422,7 @@ Following is the list of supported single precision mathematical functions. - ✓ * - | float fabsf ( float x ) - | Calculate the absolute value of its argument. + | Calculate the absolute value of its argument. - ✓ - ✓ @@ -2180,8 +2180,8 @@ The file format for binary is `.co` which means Code Object. The following comma .. note:: - When using binary code objects is that the number of arguments to the kernel is different on HIP-Clang and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel. + When using binary code objects is that the number of arguments to the kernel is different on HIP-Clang and NVCC path. Refer to the `HIP module_api sample `_ for differences in the arguments to be passed to the kernel. gfx-arch-specific-kernel ============================================================ -Clang defined '__gfx*__' macros can be used to execute gfx arch specific codes inside the kernel. Refer to the sample ``14_gpu_arch`` in ``samples/2_Cookbook``. +Clang defined '__gfx*__' macros can be used to execute gfx arch specific codes inside the kernel. Refer to the sample in `HIP 14_gpu_arch sample `_. diff --git a/projects/hip/docs/reference/programming_model.rst b/projects/hip/docs/reference/programming_model.rst deleted file mode 100644 index 6d9fff1d09..0000000000 --- a/projects/hip/docs/reference/programming_model.rst +++ /dev/null @@ -1,282 +0,0 @@ -.. meta:: - :description: This chapter describes the HIP programming model, the contract - between the programmer and the compiler/runtime executing the - code. - :keywords: AMD, ROCm, HIP, CUDA, C++ language extensions - -******************************************************************************* -Programming Model -******************************************************************************* - -HIP defines a model of mapping SIMT programs (Single Instruction, Multiple -Threads) onto various architectures, primarily GPUs. While the model may be -expressed in most imperative languages, (eg. Python via PyHIP) this document -will focus on the original C/C++ API of HIP. - -Threading Model -=============================================================================== - -The SIMT nature of HIP is captured by the ability to execute user-provided -device programs, expressed as single-source C/C++ functions or sources compiled -online/offline to binaries in bulk. - -Multiple instances of the device program (aka. kernel), so called threads, may execute in parallel, -all uniquely identified by a set of integral values. The set of integers identifying a thread relate to the hierarchy in -which threads execute. - -.. _inherent_thread_model: - -Inherent Thread Model -------------------------------------------------------------------------------- - -The thread hierarchy inherent to how AMD GPUs operate is depicted in -:numref:`inherent_thread_hierarchy`. - -.. _inherent_thread_hierarchy: - -.. figure:: ../data/reference/programming_model/thread_hierarchy.svg - :alt: Diagram depicting the thread hierarchy in HIP using nested rectangles. - The outermost rectangle is titled Grid, containing uniform rectangles - layered on one another titled Block. Each Block contains sets of uniform - rectangles layered on one another titled Warp. Each of the Warp titled - rectangles is filled with downward pointing arrows, representing single - threads. - - Hierarchy of thread groups. The arrows represent the threads. - -Warp - The most tightly coupled group of threads, both physically and logically. - Also known as a wavefront in ISA (Instruction Set Architecture) terms. - - When referring to threads inside a warp, they may be called lanes, and the - integral value identifying them the lane ID. Lane IDs aren't queried like - other thread IDs, but are user-calculated. As a consequence they are only as - multi-dimensional as the user interprets the calculated values to be. - - The size of a warp is architecture dependent and always fixed. Warps are - signified by the set of communication primitives at their disposal, detailed - under :ref:`warp_cross_lane_functions`. - -Block - The defining feature of a block (or thread block) is that all threads in the - same block are executed on the same Compute Unit, share an instance of - memory which they may use to share data, and can synchronize with one another. - - The size of a block is user-configurable but is maximized by the queryable - capabilities of the executing hardware. The unique ID of the thread within a - block is 3-dimensional as provided by the `threadIdx` built-in. When - linearizing thread IDs within a block, assume the "fast index" being - dimension ``x``, followed by the ``y`` and ``z`` dimensions. - -Grid - The outermost grouping is called a grid. A grid manifests as a single - dispatch of a kernel for execution. The unique ID of each block within a grid - is 3-dimensional, as provided by the `blockIdx` built-in and is queryable - by every thread within the block. - -Cooperative Groups Thread Model -------------------------------------------------------------------------------- - -The Cooperative Groups API introduces new APIs to launch, group, subdivide, -synchronize and identify threads, as well as some predefined group-collective -algorithms, but most importantly a matching threading model to think in terms -of. It relaxes some of the restrictions of the :ref:`inherent_thread_model` -imposed by the strict 1:1 mapping of architectural details to the programming -model. - -The rich set of APIs introduced by Cooperative Groups allow the programmer -to define their own set of thread groups which may fit their user-cases better -than those defined by the hardware. The set of implicit groups by kernel launch -parameters are still available. - -The thread hierarchy abstraction of Cooperative Groups manifest as depicted in -:numref:`coop_thread_hierarchy`. - -.. _coop_thread_hierarchy: - -.. figure:: ../data/reference/programming_model/thread_hierarchy_coop.svg - :alt: Diagram depicting the structure of Cooperative Groups using nested - rectangles. The outermost rectangle is titled Multi Grid, containing - sets of different shaped rectangles titled Grid. Each Grid contains sets - of uniform rectangles layered on oneanother titled Cluster. The Clusters - have different shapes in different Grids. Inside of the Clusters are - uniform rectangles layered on each other titled Block, which include - arrows that represent threads. - - Cooperative group thread hierarchy. - -Multi Grid - An abstraction of potentially multiple simultaneous launches of the same - kernel over multiple devices. Grids inside a multi device kernel launch need - not be of uniform size, thus allowing taking into account different device - capabilities and preferences. - - .. deprecated:: 5.0 - -.. note:: - - The performance hit of implementing cooperative groups spanning multiple - devices proved to outweigh the provided convenience, hence the use of - Multi Grid is discouraged. - -Grid - Same as the :ref:`inherent_thread_model` Grid entity. The ability to - synchronize over a grid requires the kernel to be launched using the - Cooperative Groups API. - -Cluster block - The defining feature of a cluster or block cluster is that all threads in a - cluster will use a common set of distributed shared memory which they may - use to share data or synchronize with one another. - -Block - Same as the :ref:`inherent_thread_model` Block entity. - -.. note:: - - Explicit warp-level thread handling is absent from the Cooperative Groups API. - In order to exploit the known hardware SIMD width on which built-in - functionality translates to simpler logic, one may use the group partitioning - part of the API, for instance, ``tiled_partition``. - -Memory Model -=============================================================================== - -The hierarchy of threads introduced by the :ref:`inherent_thread_model` is -induced by the memory subsystem of GPUs. :numref:`memory_hierarchy` summarizes -the memory namespaces and how they relate to the various levels of the threading -model. - -.. _memory_hierarchy: - -.. figure:: ../data/reference/programming_model/memory_hierarchy.svg - :alt: Diagram depicting the memory hierarchy using nested rectangles. The - outermost is title Grid, containing two rectangles, one titled Cluster - and the other titled Global. Cluster contains two identical rectangles - titled Block, which are partly overlaid and connected by a rectangle - titled Cluster Shared. The Block rectangles each contain a rectangle - titled Local, which in turn contain rectangles titled Warp that include - arrows representing the threads. Cluster shared contains two rectangles - titled Shared, each located within one of the Blocks. Global contains - three rectangles, titled Constant, Texture and Surface. - - Memory hierarchy. - -Local - Read-write storage only visible to the threads defining the given variables, - also called per-threas memory. The size of a block for a given kernel, - the number of concurrent warps are limited by local memory usage. - This relates to an important aspect: occupancy. This is the default memory - namespace. - -Shared - Read-write storage visible to all the threads in a given block. - -Cluster shared - Read-write storage visible to all the threads in a given cluster. - -Global - Read-write storage visible to all threads in a given grid. There are - specialized versions of global memory with different usage semantics which - are typically backed by the same hardware. - -Constant - Read-only storage visible to all threads in a given grid. It is a limited - segment of global with queryable size. - -Texture - Read-only storage visible to all threads in a given grid and accessible - through additional APIs. - -Surface - Read-write version of texture memory. - -Execution Model -=============================================================================== - -HIP programs consist of two distinct scopes: - -* The host-side running on the host processor. - - There are two types of APIs available: - - * The HIP runtime API which enables use of the single-source programming - model. - - * The HIP driver API which sits at a lower level and most importantly differs - by removing some of the facilities provided by the runtime API, most - importantly around kernel launching and argument setting. It is geared - towards implementing abstractions atop, such as the runtime API itself. - -* The device-side kernels running on GPUs. - -Both the host and the device-side APIs have synchronous and asynchronous functions in them. - -Host-side execution -------------------------------------------------------------------------------- - -The part of the host-side API which deals with device management and their -queries are synchronous. 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 on a given device. -Commands which enqueue tasks on a stream all return promptly and the command is -executed asynchronously. All side-effects of a command on a stream are visible -to all subsequent commands on the same stream. Multiple streams may point to -the same device and those streams may be fed from multiple concurrent host-side -threads. Execution on multiple streams may be concurrent but isn't required to -be. - -Asynchronous APIs involving a stream all return a stream event which may be -used to synchronize the execution of multiple streams. A user may enqueue a -barrier onto a stream referencing an event. The barrier will block until all -commands related to the event complete, at which point all side-effects of -the commands are visible to commands following the barrier, even if those -side-effects manifest on different devices. - -Streams also support executing user-defined functions as callbacks on the host. -The stream will not launch subsequent commands until the callback completes. - -Device-side execution -------------------------------------------------------------------------------- - -The SIMT programming model behind the HIP device-side execution is a -middle-ground between SMT (Simultaneous Multi-Threading) programming known from -multi-core CPUs, and SIMD (Single Instruction, Multiple Data) programming -mostly known from exploiting relevant instruction sets on CPUs (eg. -SSE/AVX/Neon). - -A HIP device compiler maps our SIMT code written in HIP C++ to an inherently -SIMD architecture (like GPUs) not by exploiting data parallelism within a -single instance of a kernel and spreading identical instructions over the SIMD -engines at hand, but by scalarizing the entire kernel and issuing the scalar -instructions of multiple kernel instances to each of the SIMD engine lanes. - -Kernel launch -------------------------------------------------------------------------------- - -Kernels may be launched in multiple ways all with different syntaxes and -intended use-cases. - -* Using the triple-chevron ``<<<...>>>`` operator on a ``__global__`` annotated - function. - -* Using ``hipLaunchKernelGGL()`` on a ``__global__`` annotated function. - - .. tip:: - - ``hipLaunchKernelGGL()`` by default is a macro expanding to triple-chevron. In cases where - language syntax extensions are undesirable, or where launching templated - and/or overloaded kernel functions define the - ``HIP_TEMPLATE_KERNEL_LAUNCH`` preprocessor macro before including the HIP - headers to turn it into a templated function. - -* Using the launch APIs supporting the triple-chevron syntax directly. - - .. caution:: - - These APIs are intended to be used/generated by tools such as the HIP - compiler itself and not intended for end-user code. Should you be - writing a tool having to launch device code using HIP, consider using these - over the alternatives. diff --git a/projects/hip/docs/old/reference/terms.md b/projects/hip/docs/reference/terms.md similarity index 99% rename from projects/hip/docs/old/reference/terms.md rename to projects/hip/docs/reference/terms.md index db9a90eee4..d79d955603 100644 --- a/projects/hip/docs/old/reference/terms.md +++ b/projects/hip/docs/reference/terms.md @@ -33,6 +33,6 @@ |Fast Math|`__cos(f)`|`__cos(f)`|`native_cos(f)`| |Vector|`float4`|`float4`|`float4`| -### Notes +## 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. diff --git a/projects/hip/docs/reference/virtual_rocr.rst b/projects/hip/docs/reference/virtual_rocr.rst new file mode 100644 index 0000000000..8241fa07ef --- /dev/null +++ b/projects/hip/docs/reference/virtual_rocr.rst @@ -0,0 +1,35 @@ +.. meta:: + :description: This chapter lists user-mode API interfaces and libraries + necessary for host applications to launch compute kernels to + available HSA ROCm kernel agents. + :keywords: AMD, ROCm, HIP, HSA, ROCR runtime, virtual memory management + +******************************************************************************* +HSA Runtime API for ROCm +******************************************************************************* + +The following functions are located in the https://github.com/ROCm/ROCR-Runtime repository. + +.. doxygenfunction:: hsa_amd_vmem_address_reserve + +.. doxygenfunction:: hsa_amd_vmem_address_free + +.. doxygenfunction:: hsa_amd_vmem_handle_create + +.. doxygenfunction:: hsa_amd_vmem_handle_release + +.. doxygenfunction:: hsa_amd_vmem_map + +.. doxygenfunction:: hsa_amd_vmem_unmap + +.. doxygenfunction:: hsa_amd_vmem_set_access + +.. doxygenfunction:: hsa_amd_vmem_get_access + +.. doxygenfunction:: hsa_amd_vmem_export_shareable_handle + +.. doxygenfunction:: hsa_amd_vmem_import_shareable_handle + +.. doxygenfunction:: hsa_amd_vmem_retain_alloc_handle + +.. doxygenfunction:: hsa_amd_vmem_get_alloc_properties_from_handle diff --git a/projects/hip/docs/sphinx/_toc.yml.in b/projects/hip/docs/sphinx/_toc.yml.in index 7df9f870bf..17af3731fc 100644 --- a/projects/hip/docs/sphinx/_toc.yml.in +++ b/projects/hip/docs/sphinx/_toc.yml.in @@ -6,52 +6,48 @@ defaults: root: index subtrees: -- caption: Installation +- caption: Install entries: - file: install/install - title: Install HIP + title: Installing HIP - file: install/build - title: Build HIP - -- caption: Reference - entries: - - file: reference/programming_model - - file: old/reference/kernel_language - - file: doxygen/mainpage.md - title: API library - - file: reference/deprecated_api_list - title: Deprecated API list - -- caption: How-to - entries: - - file: how-to/logging - title: Logging - - file: how-to/debugging - title: Debugging + title: Building HIP from source - caption: Conceptual entries: - file: understand/programming_model + - file: understand/programming_model_reference + - file: understand/hardware_implementation -- caption: Compatibility & support +- caption: How to entries: - - url: https://rocm.docs.amd.com/projects/install-on-linux/en/${branch}/reference/system-requirements.html - title: Linux (GPU & OS) - - url: https://rocm.docs.amd.com/projects/install-on-windows/en/${branch}/reference/system-requirements.html - title: Windows (GPU & OS) - - url: https://rocm.docs.amd.com/projects/install-on-linux/en/${branch}/reference/3rd-partry-support-matrix.html - title: Third-party - - url: https://rocm.docs.amd.com/projects/install-on-linux/en/${branch}/reference/user-kernel-space-compat-matrix.html - title: User/kernel space support - - url: https://rocm.docs.amd.com/projects/install-on-linux/en/${branch}/reference/docker-image-support-matrix.html - title: Docker - - url: https://rocm.docs.amd.com/en/${branch}/about/compatibility/openmp.html - title: OpenMP + - file: how-to/programming_manual + - file: how-to/hip_porting_guide + - file: how-to/hip_porting_driver_api + - file: how-to/hip_rtc + - file: how-to/performance_guidelines + - file: how-to/debugging + - file: how-to/logging + - file: how-to/faq -- caption: Contribute +- caption: Reference + entries: + - file: doxygen/html/index + - file: reference/kernel_language + title: C++ language extensions + - file: reference/terms + title: Comparing Syntax for different APIs + - file: reference/virtual_rocr + - file: reference/deprecated_api_list + title: List of deprecated APIs + +- caption: Tutorials + entries: + - url: https://github.com/ROCm/HIP-Examples + title: HIP examples + - url: https://github.com/ROCm/hip-tests/tree/develop/samples + title: HIP test samples + +- caption: About entries: - - file: CONTRIBUTING - title: Contribute to HIP - - url: https://rocm.docs.amd.com/en/latest/contribute/contributing.html - title: Contribute to ROCm docs - file: license.md diff --git a/projects/hip/docs/sphinx/requirements.in b/projects/hip/docs/sphinx/requirements.in index 69d77b3814..8d22b2d9da 100644 --- a/projects/hip/docs/sphinx/requirements.in +++ b/projects/hip/docs/sphinx/requirements.in @@ -1,2 +1,2 @@ -rocm-docs-core[api_reference]>=0.36.0 +rocm-docs-core[api_reference]==1.1.1 sphinxcontrib.doxylink diff --git a/projects/hip/docs/sphinx/requirements.txt b/projects/hip/docs/sphinx/requirements.txt index 203f05f30a..dbe8cdca79 100644 --- a/projects/hip/docs/sphinx/requirements.txt +++ b/projects/hip/docs/sphinx/requirements.txt @@ -2,40 +2,40 @@ # This file is autogenerated by pip-compile with Python 3.10 # by the following command: # -# pip-compile ./docs/sphinx/requirements.in +# pip-compile requirements.in # -accessible-pygments==0.0.3 +accessible-pygments==0.0.4 # via pydata-sphinx-theme -alabaster==0.7.12 +alabaster==0.7.16 # via sphinx -babel==2.10.3 +babel==2.14.0 # via # pydata-sphinx-theme # sphinx -beautifulsoup4==4.11.1 +beautifulsoup4==4.12.3 # via pydata-sphinx-theme -breathe==4.34.0 +breathe==4.35.0 # via rocm-docs-core -certifi==2023.7.22 +certifi==2024.2.2 # via requests -cffi==1.15.1 +cffi==1.16.0 # via # cryptography # pynacl -charset-normalizer==2.1.0 +charset-normalizer==3.3.2 # via requests -click==8.1.3 +click==8.1.7 # via # click-log # doxysphinx # sphinx-external-toc click-log==0.4.0 # via doxysphinx -cryptography==41.0.6 +cryptography==42.0.5 # via pyjwt -deprecated==1.2.13 +deprecated==1.2.14 # via pygithub -docutils==0.16 +docutils==0.21.2 # via # breathe # myst-parser @@ -43,13 +43,13 @@ docutils==0.16 # sphinx doxysphinx==3.3.7 # via rocm-docs-core -fastjsonschema==2.16.2 +fastjsonschema==2.19.1 # via rocm-docs-core -gitdb==4.0.10 +gitdb==4.0.11 # via gitpython -gitpython==3.1.41 +gitpython==3.1.43 # via rocm-docs-core -idna==3.3 +idna==3.7 # via requests imagesize==1.4.1 # via sphinx @@ -59,56 +59,53 @@ jinja2==3.1.3 # sphinx libsass==0.22.0 # via doxysphinx -lxml==4.9.2 +lxml==4.9.4 # via doxysphinx -markdown-it-py==2.2.0 +markdown-it-py==3.0.0 # via # mdit-py-plugins # myst-parser -markupsafe==2.1.1 +markupsafe==2.1.5 # via jinja2 -mdit-py-plugins==0.3.5 +mdit-py-plugins==0.4.0 # via myst-parser mdurl==0.1.2 # via markdown-it-py -mpire==2.8.0 +mpire==2.10.1 # via doxysphinx -myst-parser==1.0.0 +myst-parser==3.0.0 # via rocm-docs-core -packaging==21.3 +packaging==24.0 # via # pydata-sphinx-theme # sphinx -pycparser==2.21 +pycparser==2.22 # via cffi -pydata-sphinx-theme==0.13.3 +pydata-sphinx-theme==0.15.2 # via # rocm-docs-core # sphinx-book-theme -pygithub==1.58.2 +pygithub==2.3.0 # via rocm-docs-core -pygments==2.15.0 +pygments==2.17.2 # via # accessible-pygments # mpire # pydata-sphinx-theme # sphinx -pyjson5==1.6.4 +pyjson5==1.6.6 # via doxysphinx -pyjwt[crypto]==2.6.0 +pyjwt[crypto]==2.8.0 # via pygithub pynacl==1.5.0 # via pygithub -pyparsing==3.0.9 +pyparsing==3.1.2 # via # doxysphinx - # packaging # sphinxcontrib-doxylink python-dateutil==2.9.0.post0 # via sphinxcontrib-doxylink -pytz==2022.1 - # via babel -pyyaml==6.0 +pyyaml==6.0.1 # via # myst-parser # rocm-docs-core @@ -117,17 +114,17 @@ requests==2.31.0 # via # pygithub # sphinx -rocm-docs-core[api-reference]==0.36.0 - # via -r ./docs/sphinx/requirements.in +rocm-docs-core[api-reference]==1.1.1 + # via -r requirements.in six==1.16.0 # via python-dateutil -smmap==5.0.0 +smmap==5.0.1 # via gitdb snowballstemmer==2.2.0 # via sphinx -soupsieve==2.3.2.post1 +soupsieve==2.5 # via beautifulsoup4 -sphinx==5.3.0 +sphinx==7.3.7 # via # breathe # myst-parser @@ -139,35 +136,41 @@ sphinx==5.3.0 # sphinx-external-toc # sphinx-notfound-page # sphinxcontrib-doxylink -sphinx-book-theme==1.0.1 +sphinx-book-theme==1.1.2 # via rocm-docs-core -sphinx-copybutton==0.5.1 +sphinx-copybutton==0.5.2 # via rocm-docs-core -sphinx-design==0.3.0 +sphinx-design==0.5.0 # via rocm-docs-core -sphinx-external-toc==0.3.1 +sphinx-external-toc==1.0.1 # via rocm-docs-core -sphinx-notfound-page==0.8.3 +sphinx-notfound-page==1.0.0 # via rocm-docs-core -sphinxcontrib-applehelp==1.0.2 +sphinxcontrib-applehelp==1.0.8 # via sphinx -sphinxcontrib-devhelp==1.0.2 +sphinxcontrib-devhelp==1.0.6 # via sphinx sphinxcontrib-doxylink==1.12.3 - # via -r ./docs/sphinx/requirements.in -sphinxcontrib-htmlhelp==2.0.0 + # via -r requirements.in +sphinxcontrib-htmlhelp==2.0.5 # via sphinx sphinxcontrib-jsmath==1.0.1 # via sphinx -sphinxcontrib-qthelp==1.0.3 +sphinxcontrib-qthelp==1.0.7 # via sphinx -sphinxcontrib-serializinghtml==1.1.5 +sphinxcontrib-serializinghtml==1.1.10 # via sphinx -tqdm==4.66.1 +tomli==2.0.1 + # via sphinx +tqdm==4.66.2 # via mpire -typing-extensions==4.4.0 - # via pydata-sphinx-theme -urllib3==1.26.18 - # via requests -wrapt==1.15.0 +typing-extensions==4.11.0 + # via + # pydata-sphinx-theme + # pygithub +urllib3==2.2.1 + # via + # pygithub + # requests +wrapt==1.16.0 # via deprecated diff --git a/projects/hip/docs/understand/glossary.md b/projects/hip/docs/understand/glossary.md new file mode 100644 index 0000000000..69f697f565 --- /dev/null +++ b/projects/hip/docs/understand/glossary.md @@ -0,0 +1,23 @@ +# Glossary of terms + +- **host**, **host cpu** : Executes the HIP runtime API and is capable of initiating kernel launches to one or more devices. +- **default device** : Each host thread maintains a default device. +Most HIP runtime APIs (including memory allocation, copy commands, kernel launches) do not accept an explicit device +argument but instead implicitly use the default device. +The default device can be set with ```hipSetDevice```. + +- **active host thread** - the thread which is running the HIP APIs. + +- **HIP-Clang** - Heterogeneous AMDGPU Compiler, with its capability to compile HIP programs on AMD platform (https://github.com/RadeonOpenCompute/llvm-project). + +- **clr** - a repository for AMD Common Language Runtime, contains source codes for AMD's compute languages runtimes: HIP and OpenCL. +clr (https://github.com/ROCm/clr) contains the following three parts, + - ```hipamd```: contains implementation of HIP language on AMD platform. + - ```rocclr```: contains common runtime used in HIP and OpenCL, which provides virtual device interfaces that compute runtimes interact with different backends such as ROCr on Linux or PAL on Windows. + - ```opencl```: contains implementation of OpenCL on AMD platform. + +- **hipify tools** - tools to convert CUDA code to portable C++ code (https://github.com/ROCm/HIPIFY). + +- **hipconfig** - tool to report various configuration properties of the target platform. + +- **nvcc** - NVIDIA CUDA ```nvcc``` compiler, do not capitalize. diff --git a/projects/hip/docs/understand/hardware_implementation.rst b/projects/hip/docs/understand/hardware_implementation.rst new file mode 100644 index 0000000000..f95d8fc6b4 --- /dev/null +++ b/projects/hip/docs/understand/hardware_implementation.rst @@ -0,0 +1,163 @@ +.. meta:: + :description: This chapter describes the typical hardware implementation of GPUs supported by HIP. + :keywords: AMD, ROCm, HIP, Hardware, Compute Unit, ALU, VALU, Cache, Registers, LDS + +.. _hardware_implementation: + +******************************************************************************* +Hardware Implementation +******************************************************************************* + +This chapter describes the typical hardware implementation of GPUs supported by +HIP, and how the :ref:`inherent_thread_model` maps to the hardware. + +Compute Units +============= + +The basic building block of a GPU is a compute unit (CU), also known +as streaming multiprocessor (SM) on NVIDIA GPUs. The thread blocks making up a +grid are scheduled for execution on CUs. Each block is assigned to an +individual CU, and a CU can accommodate several blocks. Depending on their +resource usage up to thousands of threads can reside on a CU. + +CUs contain an array of processing elements, referred to as +vector ALU (VALU), that execute the actual instructions of the threads +according to the :ref:`SIMT model`, together with the +necessary registers and caches. + +The threads are executed in groupings called warps. The amount of threads +making up a warp is architecture dependent. +On AMD GPUs the warp size is commonly 64 threads, except in RDNA +architectures which can utilize a warp size of 32 or 64 respectively. +The warp size of supported AMD GPUs is listed in the :doc:`rocm:reference/gpu-arch-specs`. +NVIDIA GPUs have a warp size of 32. + +In contrast to CPUs, GPUs generally do not employ complex cache structures or +control logic, like branch prediction or out-of-order execution, but instead +rely on massive hardware multithreading to hide latency. + +Context switching between warps residing on a CU incurs no overhead, +as the context for the warps is stored on the CU and does not need to +be fetched from memory. If there are not enough free registers to accommodate +all warps of a block, the block can not be scheduled to that CU and it +has to wait until other blocks finish execution. + +The amount of warps that can reside concurrently on a CU, known +as occupancy, is determined by the warp's resource usage of registers and +shared memory. + +.. figure:: ../data/hardware_implementation/compute_unit.svg + :alt: Diagram depicting the general structure of a compute unit of an AMD + GPU. + + An AMD Graphics Core Next (GCN) CU. The CDNA and RDNA CUs are based on + variations of the GCN CU. + +On AMD GCN GPUs the basic structure of a CU is: + +* four Single Instruction Multiple Data units (SIMDs) +* a vector cache +* a local data share +* and a scalar unit + +SIMD +---- + +A SIMD consists of a VALU, that executes the instruction of a warp, together +with a register file, that provides the registers warps. + +The size of the warp is inherently related to the width of the vector ALU of +the SIMD. On GCN compute units the width of the VALU is 16, so a warp can be +issued to a SIMD every 4 cycles. Since a CU has 4 SIMDs it issues one +warp per cycle. The instructions of a warp are effectively executed in +lock-step. + +A SIMD always executes the same instruction for the whole VALU. If the control +flow of a warp diverges, the performance is decreased, as the results for the +threads that do not participate in that branch have to be masked out, and the +instructions of the other branch have to be executed in the same way. The best +performance can therefore be achieved when thread divergence is kept to a warp +level, i.e. when all threads in a warp take the same execution path. + +Vector Cache +------------ + +The usage of cache on a GPU differs from that on a CPU, as there is less cache +available per thread. Its main purpose is to coalesce memory accesses of the +warps in order to reduce the amount of accesses to device memory, and make that +memory available for other warps that currently reside on the compute unit, that +also need to load those values. + +Local Data Share +---------------- + +The local data share is memory that is accessible to all threads within a block. +Its latency and bandwidth is comparable to that of the vector cache. It can be +used to share memory between the threads in a block, or as a software managed +cache. + +Scalar Unit +----------- + +The scalar unit performs instructions that are uniform within a warp. It +thereby improves efficiency and reduces the pressure on the vector ALUs and the +vector register file. + +CDNA Architecture +================= + +The general structure of CUs stays mostly as it is in GCN +architectures. The most prominent change is the addition of matrix ALUs, which +can greatly improve the performance of algorithms involving matrix +multiply-accumulate operations for +:doc:`int8, float16, bfloat16 or float32`. + +.. figure:: ../data/hardware_implementation/cdna3_cu.png + :alt: Block diagram showing the structure of a CDNA3 compute unit. It includes + Shader Cores, the Matrix Core Unit, a Local Data Share used for sharing + memory between threads in a block, an L1 Cache and a Scheduler. The + Shader Cores represent the vector ALUs and the Matrix Core Unit + represents the matrix ALUs. The Local Data Share is used as the shared + memory. + + Block Diagram of a CDNA3 Compute Unit. + +RDNA Architecture +================= + +RDNA makes a fundamental change to CU design, by changing the +size of a warp to 32 threads. This is done by effectively combining two GCN5 +SIMDs, creating a VALU of width 32, so that a whole warp can be issued in one +cycle. The CU is also replaced by the work group processor (WGP), +which encompasses two CUs. For backwards compatibility the WGP can +also run in wave64 mode, in which it issues a warp of size 64 in two cycles. + +It also adds an extra layer of cache to the WGP, shared by the CUs +within it. This cache is referred to as L1 cache, promoting the per-CU cache to +an L0 cache. + +.. figure:: ../data/hardware_implementation/rdna3_cu.png + :alt: Block diagram showing the structure of an RDNA3 Compute Unit. It + consists of four SIMD units, each including a vector and scalar register + file, with the corresponding scalar and vector ALUs. All four SIMDs + share a scalar and instruction cache, as well as the shared memory. Two + of the SIMD units each share an L0 cache. + + Block Diagram of an RDNA3 work group processor. + +Shader Engines +============== + +For hardware implementation's sake, multiple CUs are grouped +together into a Shader Engine or Compute Engine, typically sharing some fixed +function units or memory subsystem resources. + +.. figure:: ../data/hardware_implementation/cdna2_gcd.png + :alt: Block diagram showing four Compute Engines each with 28 Compute Units + inside. These four Compute Engines share one block of L2 Cache. Around + them are four Memory Controllers. To the top and bottom of all these are + eight blocks of Infinity Fabric Links. Two Video Core Next blocks sit in + the top corners. At the very bottom spans a colored section reading + Infinity Fabric. + + Block Diagram of a CDNA2 Graphics Compute Die. diff --git a/projects/hip/docs/understand/programming_model.rst b/projects/hip/docs/understand/programming_model.rst index 6f1d5c32fd..092cf6796c 100644 --- a/projects/hip/docs/understand/programming_model.rst +++ b/projects/hip/docs/understand/programming_model.rst @@ -5,32 +5,32 @@ :keywords: AMD, ROCm, HIP, CUDA, API design ******************************************************************************* -Programming Model +Understanding the HIP programming model ******************************************************************************* The HIP programming model makes it easy to map data-parallel C/C++ algorithms to -massively parallel, wide SIMD architectures, such as GPUs. As a consequence, one -needs a basic understanding of the underlying device architecture to make efficient -use of HIP and GPGPU (General Purpose Graphics Processing Unit) programming in general. +massively parallel, wide single instruction, multiple data (SIMD) architectures, +such as GPUs. A basic understanding of the underlying device architecture helps you +make efficient use of HIP and general purpose graphics processing unit (GPGPU) +programming in general. RDNA & CDNA Architecture Summary -=============================================================================== +================================ Most GPU architectures, like RDNA and CDNA, have a hierarchical structure. -The inner-most piece is a Single Instruction Multiple Data (SIMD) enabled -vector Arithmetic Logical Unit (ALU). In addition to the vector ALUs, most -recent GPUs also house also house matrix ALUs for accelerating algorithms involving -matrix multiply-accumulate operations. AMD GPUs also contain scalar ALUs, that -can be used to reduce the load on the vector ALU by performing operations which -are uniform for all threads of a warp. +The innermost piece is a SIMD-enabled vector Arithmetic Logical Unit (ALU). +In addition to the vector ALUs, most recent GPUs also house matrix ALUs for +accelerating algorithms involving matrix multiply-accumulate operations. +AMD GPUs also contain scalar ALUs, that can be used to reduce the load on the +vector ALU by performing operations which are uniform for all threads of a warp. A set of ALUs, together with register files, caches and shared memory, comprise -a larger block, often referred to as a Compute Unit (CU), e.g. in OpenCL and -AMD block diagrams, or as Streaming Multiprocessor (SM). +a larger block, often referred to as a compute unit (CU), e.g. in OpenCL and +AMD block diagrams, or as streaming multiprocessor (SM). .. _rdna3_cu: -.. figure:: ../data/understand/programming_model/rdna3_cu.png +.. figure:: ../data/programming_model/understand/rdna3_cu.png :alt: Block diagram showing the structure of an RDNA3 Compute Unit. It consists of four SIMD units, each including a vector and scalar register file, with the corresponding scalar and vector ALUs. All four SIMDs @@ -41,7 +41,7 @@ AMD block diagrams, or as Streaming Multiprocessor (SM). .. _cdna3_cu: -.. figure:: ../data/understand/programming_model/cdna3_cu.png +.. figure:: ../data/programming_model/understand/cdna3_cu.png :alt: Block diagram showing the structure of a CDNA3 compute unit. It includes Shader Cores, the Matrix Core Unit, a Local Data Share used for sharing memory between threads in a block, an L1 Cache and a Scheduler. The @@ -50,13 +50,13 @@ AMD block diagrams, or as Streaming Multiprocessor (SM). Block Diagram of a CDNA3 Compute Unit. -FFor implementation in hardware, multiple Compute Units are grouped together into +For implementation in hardware, multiple CUs are grouped together into a Shader Engine or Compute Engine, typically sharing some fixed function units or memory subsystem resources. .. _cdna2_gcd: -.. figure:: ../data/understand/programming_model/cdna2_gcd.png +.. figure:: ../data/programming_model/understand/cdna2_gcd.png :alt: Block diagram showing four Compute Engines each with 28 Compute Units inside. These four Compute Engines share one block of L2 Cache. Around them are four Memory Controllers. To the top and bottom of all these are @@ -66,21 +66,23 @@ memory subsystem resources. Block Diagram of a CDNA2 Graphics Compute Die. -Single Instruction Multiple Threads -=============================================================================== +.. _programming_model_simt: -The SIMT programming model behind the HIP device-side execution is a -middle-ground between SMT (Simultaneous Multi-Threading) programming known from -multi-core CPUs, and SIMD (Single Instruction, Multiple Data) programming -mostly known from exploiting relevant instruction sets on CPUs (eg. SSE/AVX/Neon). +Single Instruction Multiple Threads +=================================== + +The single instruction, multiple threads (SIMT) programming model behind the +HIP device-side execution is a middle-ground between SMT (Simultaneous Multi-Threading) +programming known from multicore CPUs, and SIMD (Single Instruction, Multiple Data) programming +mostly known from exploiting relevant instruction sets on CPUs (for example SSE/AVX/Neon). A HIP device compiler maps SIMT code written in HIP C++ to an inherently SIMD -architecture (like GPUs), by scalarizing the entire kernel, and issuing the scalar +architecture (like GPUs). This is done by scalarizing the entire kernel and issuing the scalar instructions of multiple kernel instances to each of the SIMD engine lanes, rather -than exploiting data parallelism within a single instance of a kernel and spreading +than exploiting data parallelism within a single instance of a kernel and spreading identical instructions over the available SIMD engines. -Consider the following kernel +Consider the following kernel: .. code:: cu @@ -101,7 +103,7 @@ typically look the following: .. _simt: -.. figure:: ../data/understand/programming_model/simt.svg +.. figure:: ../data/programming_model/understand/simt.svg :alt: Image representing the instruction flow of a SIMT program. Two identical arrows pointing downward with blocks representing the instructions inside and ellipsis between the arrows. The instructions represented in @@ -115,8 +117,8 @@ usually isn't exploited from the width of the built-in vector types, but via the thread id constants ``threadIdx.x``, ``blockIdx.x``, etc. For more details, refer to :ref:`inherent_thread_model`. -Heterogenous Programming -=============================================================================== +Heterogeneous Programming +========================= The HIP programming model assumes two execution contexts. One is referred to as *host* while compute kernels execute on a *device*. These contexts have @@ -129,7 +131,7 @@ a few key differences between the two: * The C++ abstract machine assumes a unified memory address space, meaning that one can always access any given address in memory (assuming the absence of data races). HIP however introduces several memory namespaces, an address - from one means nothing in another. Moreover not all address spaces are + from one means nothing in another. Moreover, not all address spaces are accessible from all contexts. If one were to look at :ref:`cdna2_gcd` and inside the :ref:`cdna3_cu`, @@ -148,7 +150,7 @@ a few key differences between the two: architectures. * Asynchrony is at the forefront of the HIP API. Computations launched on the device - execute asynchronously with respect to the host and it is the user's responsibility to + execute asynchronously with respect to the host, and it is the user's responsibility to synchronize their data dispatch/fetch with computations on the device. HIP does perform implicit synchronization on occasions, more advanced than other APIs such as OpenCL or SYCL, in which the responsibility of synchronization mostly depends on the user. diff --git a/projects/hip/docs/understand/programming_model_reference.rst b/projects/hip/docs/understand/programming_model_reference.rst new file mode 100644 index 0000000000..600fcad3da --- /dev/null +++ b/projects/hip/docs/understand/programming_model_reference.rst @@ -0,0 +1,265 @@ +.. meta:: + :description: This chapter describes the HIP programming model, the contract + between the programmer and the compiler/runtime executing the + code. + :keywords: AMD, ROCm, HIP, CUDA, C++ language extensions + +******************************************************************************* +Programming model reference +******************************************************************************* + +HIP defines a model for mapping single instruction, multiple threads (SIMT) programs +onto various architectures, primarily GPUs. While the model may be expressed +in most imperative languages, (eg. Python via PyHIP) this document will focus on +the original C/C++ API of HIP. + +Threading Model +=============== + +The SIMT nature of HIP is captured by the ability to execute user-provided +device programs, expressed as single-source C/C++ functions or sources compiled +online/offline to binaries in bulk. + +Multiple instances of the device program (or kernel) are called threads and may +execute in parallel. All uniquely identified by a set of integral values, or thread IDs. +The set of integers identifying a thread relate to the hierarchy in which threads execute. + +.. _inherent_thread_model: + +Inherent Thread Model +--------------------- + +The thread hierarchy inherent to how AMD GPUs operate is depicted in +:numref:`inherent_thread_hierarchy`. + +.. _inherent_thread_hierarchy: + +.. figure:: ../data/programming_model/reference/thread_hierarchy.svg + :alt: Diagram depicting nested rectangles of varying color. The outermost one + titled "Grid", inside sets of uniform rectangles layered on one another + titled "Block". Each "Block" containing sets of uniform rectangles + layered on one another titled "Warp". Each of the "Warp" titled + rectangles filled with downward pointing arrows inside. + + Hierarchy of thread groups. + +Warp + The innermost grouping is called a warp, or a wavefront in ISA terms. A warp + is the most tightly coupled groups of threads, both physically and logically. + + When referring to threads inside a warp, they may be called lanes, and the + integral value identifying them the lane ID. Lane IDs aren't queried like + other thread IDs, but are user-calculated. As a consequence they are only as + multidimensional as the user interprets the calculated values to be. + + The size of a warp is architecture dependent and always fixed. Warps are + signified by the set of communication primitives at their disposal, as + discussed in :ref:`warp-cross-lane`. + +Block + The middle grouping is called a block or thread block. The defining feature + of a block is that all threads in a block will share an instance of memory + which they may use to share data or synchronize with one another. + + The size of a block is user-configurable but is maximized by the queryable + capabilities of the executing hardware. The unique ID of the thread within a + block is 3-dimensional as provided by the API. When linearizing thread IDs + within a block, assume the "fast index" being dimension ``x``, followed by + the ``y`` and ``z`` dimensions. + +Grid + The outermost grouping is called a grid. A grid manifests as a single + dispatch of kernels for execution. The unique ID of each block within a grid + is 3-dimensional, as provided by the API and is queryable by every thread + within the block. + +Cooperative Groups Thread Model +------------------------------- + +The Cooperative Groups API introduces new APIs to launch, group, subdivide, +synchronize and identify threads, as well as some predefined group-collective +algorithms, but most importantly a matching threading model to think in terms +of. It relaxes some restrictions of the :ref:`inherent_thread_model` +imposed by the strict 1:1 mapping of architectural details to the programming +model. + +The rich set of APIs introduced by Cooperative Groups allow the programmer to +define their own set of thread groups which may fit their user-cases better than +those defined by the hardware. The set of implicit groups by kernel launch +parameters are still available. + +The thread hierarchy abstraction of Cooperative Groups manifest as depicted in +:numref:`coop_thread_hierarchy`. + +.. _coop_thread_hierarchy: + +.. figure:: ../data/programming_model/reference/thread_hierarchy_coop.svg + :alt: Diagram depicting nested rectangles of varying color. The outermost one + titled "Grid", inside sets of different sized rectangles layered on + one another titled "Block". Each "Block" containing sets of uniform + rectangles layered on one another titled "Warp". Each of the "Warp" + titled rectangles filled with downward pointing arrows inside. + + Cooperative group thread hierarchy. + +Multi Grid + An abstraction of potentially multiple simultaneous launches of + the same kernel over multiple devices. Grids inside a multi device kernel + launch need not be of uniform size, thus allowing taking into account + different device capabilities and preferences. + + .. deprecated:: 5.0 + +Grid + Same as the :ref:`inherent_thread_model` Grid entity. The ability to + synchronize over a grid requires the kernel to be launched using the + Cooperative Groups API. + +Block + Same as the :ref:`inherent_thread_model` Block entity. + +.. note:: + + Explicit warp-level thread handling is absent from the Cooperative Groups API. + In order to exploit the known hardware SIMD width on which built-in + functionality translates to simpler logic, one may use the group partitioning + part of the API, such as ``tiled_partition``. + +Memory Model +============ + +The hierarchy of threads introduced by :ref:`inherent_thread_model` is induced +by the memory subsystem of GPUs. :numref:`memory_hierarchy` summarizes that memory namespaces and +how they relate to the various levels of the threading model. + +.. _memory_hierarchy: + +.. figure:: ../data/programming_model/reference/memory_hierarchy.svg + :alt: Diagram depicting nested rectangles of varying color. The outermost one + titled "Grid", inside on the upper half a rectangle titled "Cluster". + Inside it are two identical rectangles titled "Block", inside them are + ones titled "Local" with multiple "Warp" titled rectangles. Blocks have + not just Local inside, but also rectangles titled "Shared". The Shared + rectangles of Blocks in the same Cluster are grouped together with a + translucent halo titled "Cluster shared". Outside the Cluster but + inside the Grid is a rectangle titled "Global" with three others + inside: "Surface", "Texture" (same color) and "Constant" (different + color). + + Memory hierarchy. + +Local or per-thread memory + Read-write storage only visible to the threads defining the given variables, + also called per-thread memory. The size of a block for a given kernel, + the number of concurrent warps are limited by local memory usage. + This relates to an important aspect: occupancy. This is the default memory + namespace. + +Shared memory + Read-write storage visible to all the threads in a given block. + +Distributed shared memory + Read-write storage visible to all the threads in a given block cluster. + +Global + Read-write storage visible to all threads in a given grid. There are + specialized versions of global memory with different usage semantics which + are typically backed by the same hardware storing global. + + Constant + Read-only storage visible to all threads in a given grid. It is a limited + segment of global with queryable size. + + Texture + Read-only storage visible to all threads in a given grid and accessible + through additional APIs. + + Surface + A read-write version of texture memory. + +Execution Model +=============== + +HIP programs consist of two distinct scopes: + +* The host-side API running on the host processor. There are to APIs available: + + * The HIP runtime API which enables use of the single-source programming + model. + + * The HIP driver API which sits at a lower level and most importantly differs + by removing some facilities provided by the runtime API, most + importantly around kernel launching and argument setting. It is geared + towards implementing abstractions atop, such as the runtime API itself. + +* The device-side kernels running on GPUs. Both the host and the device-side + APIs have synchronous and asynchronous functions in them. + +Host-side execution +------------------- + +The part of the host-side API which deals with device management and their +queries are synchronous. 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 relating to a given device. +Commands which enqueue tasks on a stream all return promptly and the command is +executed asynchronously. All side effects of a command on a stream are visible +to all subsequent commands on the same stream. Multiple streams may point to +the same device and those streams may be fed from multiple concurrent host-side +threads. Execution on multiple streams may be concurrent but isn't required to +be. + +Asynchronous APIs involving a stream all return a stream event which may be +used to synchronize the execution of multiple streams. A user may enqueue a +barrier onto a stream referencing an event. The barrier will block until +the command related to the event does not complete, at which point all +side effects of the command shall be visible to commands following the barrier, +even if those side effects manifest on different devices. + +Streams also support executing user-defined functions as callbacks on the host. +The stream will not launch subsequent commands until the callback completes. + +Device-side execution +--------------------- + +The SIMT programming model behind the HIP device-side execution is a +middle-ground between SMT (Simultaneous Multi-Threading) programming known from +multicore CPUs, and SIMD (Single Instruction, Multiple Data) programming +mostly known from exploiting relevant instruction sets on CPUs (for example +SSE/AVX/Neon). + +A HIP device compiler maps our SIMT code written in HIP C++ to an inherently +SIMD architecture (like GPUs) not by exploiting data parallelism within a +single instance of a kernel and spreading identical instructions over the SIMD +engines at hand, but by scalarizing the entire kernel and issuing the scalar +instructions of multiple kernel instances to each of the SIMD engine lanes. + +Kernel launch +------------- + +Kernels may be launched in multiple ways all with different syntaxes and +intended use-cases. + +* Using the triple-chevron ``<<<...>>>`` operator on a ``__global__`` annotated + function. + +* Using ``hipLaunchKernelGGL()`` on a ``__global__`` annotated function. + + .. tip:: + + This name by default is a macro expanding to triple-chevron. In cases where + language syntax extensions are undesirable, or where launching templated + and/or overloaded kernel functions define the + ``HIP_TEMPLATE_KERNEL_LAUNCH`` preprocessor macro before including the HIP + headers to turn it into a templated function. + +* Using the launch APIs supporting the triple-chevron syntax directly. + + .. caution:: + + These APIs are intended to be used/generated by tools such as the HIP + compiler itself and not intended towards end-user code. Should you be + writing a tool having to launch device code using HIP, consider using these + over the alternatives. diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 2d3cc8a93f..2fea639d78 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -4252,8 +4252,10 @@ hipError_t hipGetProcAddress(const char* symbol, void** pfn, int hipVersion, ui * the host side. The symbol can be in __constant or device space. * Note that the symbol name needs to be encased in the HIP_SYMBOL macro. * This also applies to hipMemcpyFromSymbol, hipGetSymbolAddress, and hipGetSymbolSize. - * For detail usage, see the example at - * https://github.com/ROCm/HIP/blob/develop/docs/user_guide/hip_porting_guide.md + * For detailed usage, see the + * memcpyToSymbol example + * in the HIP Porting Guide. + * * * @param[out] symbol pointer to the device symbole * @param[in] src pointer to the source address