diff --git a/README.md b/README.md index f721347b33..39165993cc 100644 --- a/README.md +++ b/README.md @@ -67,7 +67,7 @@ See the [Installation](INSTALL.md) notes. ## Simple Example The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree. Programmers familiar with CUDA will also be able to quickly learn and start coding with the HIP API. -Compute kernels are launched with the "hipLaunchKernel" macro call. Here is simple example showing a +Compute kernels are launched with the "hipLaunchKernelGGL" macro call. Here is simple example showing a snippet of HIP API code: ```cpp @@ -78,7 +78,7 @@ hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice); const unsigned blocks = 512; const unsigned threadsPerBlock = 256; -hipLaunchKernel(vector_square, /* compute kernel*/ +hipLaunchKernelGGL(vector_square, /* compute kernel*/ dim3(blocks), dim3(threadsPerBlock), 0/*dynamic shared*/, 0/*stream*/, /* launch config*/ C_d, A_d, N); /* arguments to the compute kernel */ @@ -142,7 +142,7 @@ The README with the procedures and tips the team used during this porting effort ## Tour of the HIP Directories * **include**: * **hip_runtime_api.h** : Defines HIP runtime APIs and can be compiled with many standard Linux compilers (GCC, ICC, CLANG, etc), in either C or C++ mode. - * **hip_runtime.h** : Includes everything in hip_runtime_api.h PLUS hipLaunchKernel and syntax for writing device kernels and device functions. hip_runtime.h can be compiled using a standard C++ compiler but will expose a subset of the available functions. + * **hip_runtime.h** : Includes everything in hip_runtime_api.h PLUS hipLaunchKernelGGL and syntax for writing device kernels and device functions. hip_runtime.h can be compiled using a standard C++ compiler but will expose a subset of the available functions. * **amd_detail/**** , **nvidia_detail/**** : Implementation details for specific platforms. HIP applications should not include these files directly. * **bin**: Tools and scripts to help with hip porting diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index d1d816123c..a26597ca67 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -42,7 +42,7 @@ HIP provides the following: - Memory management (hipMalloc(), hipMemcpy(), hipFree(), etc.) - Streams (hipStreamCreate(),hipStreamSynchronize(), hipStreamWaitEvent(), etc.) - Events (hipEventRecord(), hipEventElapsedTime(), etc.) -- Kernel launching (hipLaunchKernel is a standard C/C++ function that replaces <<< >>>) +- Kernel launching (hipLaunchKernel/hipLaunchKernelGGL is the preferred way of launching kernels. hipLaunchKernelGGL is a standard C/C++ macro that can serve as an alternative way to launch kernels, replacing the CUDA triple-chevron (<<< >>>) syntax). - HIP Module API to control when adn how code is loaded. - CUDA-style kernel coordinate functions (threadIdx, blockIdx, blockDim, gridDim) - Cross-lane instructions including shfl, ballot, any, all diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 70a02acb29..04c40ec39e 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -31,7 +31,6 @@ - [Surface Functions](#surface-functions) - [Timer Functions](#timer-functions) - [Atomic Functions](#atomic-functions) - * [Caveats and Features Under-Development:](#caveats-and-features-under-development) - [Warp Cross-Lane Functions](#warp-cross-lane-functions) * [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions) * [Warp Shuffle Functions](#warp-shuffle-functions) @@ -101,8 +100,8 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t ## Calling `__global__` Functions `__global__` functions are often referred to as *kernels,* and calling one is termed *launching the kernel.* These functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute. HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel in addition to the Cuda <<< >>> syntax. In HIP, -- Kernels launch with either <<< >>> syntax or the "hipLaunchKernel" function -- The first five parameters to hipLaunchKernel are the following: +- Kernels launch with either <<< >>> syntax or the "hipLaunchKernelGGL" function +- The first five parameters to hipLaunchKernelGGL are the following: - **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. The hipify tools insert this automatically. - **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. - **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block. @@ -112,7 +111,7 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t ``` -// Example pseudo code introducing hipLaunchKernel: +// Example pseudo code introducing hipLaunchKernelGGL: __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N) { ... @@ -120,11 +119,11 @@ __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N) MyKernel<<>> (a,b,c,n); // Alternatively, kernel can be launched by -// hipLaunchKernel(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n); +// hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n); ``` -The hipLaunchKernel macro always starts with the five parameters specified above, followed by the kernel arguments. HIPIFY tools optionally convert Cuda launch syntax to hipLaunchKernel, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernel parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (thread*, block*, grid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins). +The hipLaunchKernelGGL macro always starts with the five parameters specified above, followed by the kernel arguments. HIPIFY tools optionally convert Cuda launch syntax to hipLaunchKernelGGL, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernelGGL parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (thread*, block*, grid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins). ## Kernel-Launch Example @@ -154,7 +153,7 @@ void callMyKernel() MyKernel<<>> (a,b,c,n); // Alternatively, kernel can be launched by - // hipLaunchKernel(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N); + // hipLaunchKernelGGL(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N); } ``` @@ -172,8 +171,7 @@ Previously, it was essential to declare dynamic shared memory using the HIP_DYNA Now, the HIP-Clang compiler provides support for extern shared declarations, and the HIP_DYNAMIC_SHARED option is no longer required.. ### `__managed__` -Managed memory, except the `__managed__` keyword, are supported in HIP combined host/device compilation. -Support of `__managed__` keyword is under development. +Managed memory, including the `__managed__` keyword, are supported in HIP combined host/device compilation. ### `__restrict__` The `__restrict__` keyword tells the compiler that the associated memory pointer will not alias with any other pointer in the kernel or function. This feature can help the compiler generate better code. In most cases, all pointer arguments must use this keyword to realize the benefit. @@ -182,25 +180,9 @@ The `__restrict__` keyword tells the compiler that the associated memory pointer ## Built-In Variables ### Coordinate Built-Ins -These built-ins determine the coordinate of the active work item in the execution grid. They are defined in hip_runtime.h (rather than being implicitly defined by the compiler). - -| **HIP Syntax** | **Cuda Syntax** | -| --- | --- | -| threadIdx.x | threadIdx.x | -| threadIdx.y | threadIdx.y | -| threadIdx.z | threadIdx.z | -| | | -| blockIdx.x | blockIdx.x | -| blockIdx.y | blockIdx.y | -| blockIdx.z | blockIdx.z | -| | | -| blockDim.x | blockDim.x | -| blockDim.y | blockDim.y | -| blockDim.z | blockDim.z | -| | | -| gridDim.x | gridDim.x | -| gridDim.y | gridDim.y | -| gridDim.z | gridDim.z | +Built-ins determine the coordinate of the active work item in the execution grid. They are defined in amd_hip_runtime.h (rather than being implicitly defined by the compiler). +In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance: +threadIdx.x, blockIdx.y, gridDim.y, etc. ### warpSize The warpSize variable is of type int and contains the warp size (in threads) for the target device. Note that all current Nvidia devices return 32 for this variable, and all current AMD devices return 64. Device code should use the warpSize built-in to develop portable wave-aware code. @@ -208,7 +190,7 @@ The warpSize variable is of type int and contains the warp size (in threads) for ## Vector Types -Note that these types are defined in hip_runtime.h and are not automatically provided by the compiler. +Note that these types are defined in hip_runtime.h and are not automatically provided by the compiler. ### Short Vector Types @@ -588,10 +570,6 @@ By default, this compilation flag is not set("0"), so hip runtime will use curre If this compilation flag is set to "1", that is, with the cmake option "-D__HIP_USE_CMPXCHG_FOR_FP_ATOMICS=1", the old float/double atomicAdd functions will be used instead, for compatibility with compilers not supporting floating point atomics. For details steps how to build hip runtime, please refer to the section "build HIPAMD" (https://github.com/ROCm-Developer-Tools/hipamd/blob/develop/INSTALL.md). -### Caveats and Features Under-Development: - -- HIP enables atomic operations on 32-bit integers. Additionally, it supports an atomic float add. AMD hardware, however, implements the float add using a CAS loop, so this function may not perform efficiently. - ## Warp Cross-Lane Functions Warp cross-lane functions operate across all lanes in a warp. The hardware guarantees that all warp lanes will execute in lockstep, so additional synchronization is unnecessary, and the instructions use no shared memory. @@ -708,12 +686,31 @@ The Cuda `__prof_trigger()` instruction is not supported. ## Assert -The assert function is under development. -HIP does support an "abort" call which will terminate the process execution from inside the kernel. +The assert function is supported in HIP. +Assert function is used for debugging purpose, when the input expression equals to zero, the execution will be stopped. +``` +void assert(int input) +``` + +There are two kinds of implementations for assert functions depending on the use sceneries, +- One is for the host version of assert, which is defined in assert.h, +- Another is the device version of assert, which is implemented in hip/hip_runtime.h. +Users need to include assert.h to use assert. For assert to work in both device and host functions, users need to include "hip/hip_runtime.h". ## Printf -The printf function is under development. +Printf function is supported in HIP. +The following is a simple example to print information in the kernel. + +``` +#include + +__global__ void run_printf() { printf("Hello World\n"); } + +int main() { + run_printf<<>>(); +} +``` ## Device-Side Dynamic Global Memory Allocation @@ -728,15 +725,15 @@ GPU multiprocessors have a fixed pool of resources (primarily registers and shar __launch_bounds__ allows the application to provide usage hints that influence the resources (primarily registers) used by the generated code. It is a function attribute that must be attached to a __global__ function: ``` -__global__ void `__launch_bounds__`(MAX_THREADS_PER_BLOCK, MIN_WARPS_PER_EU) MyKernel(...) ... +__global__ void `__launch_bounds__`(MAX_THREADS_PER_BLOCK, MIN_WARPS_PER_EXECUTION_UNIT) MyKernel(hipGridLaunch lp, ...) ... ``` __launch_bounds__ supports two parameters: - MAX_THREADS_PER_BLOCK - The programmers guarantees that kernel will be launched with threads less than MAX_THREADS_PER_BLOCK. (On NVCC this maps to the .maxntid PTX directive). If no launch_bounds is specified, MAX_THREADS_PER_BLOCK is the maximum block size supported by the device (typically 1024 or larger). Specifying MAX_THREADS_PER_BLOCK less than the maximum effectively allows the compiler to use more resources than a default unconstrained compilation that supports all possible block sizes at launch time. -The threads-per-block is the product of (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z). -- MIN_WARPS_PER_EU - directs the compiler to minimize resource usage so that the requested number of warps can be simultaneously active on a multi-processor. Since active warps compete for the same fixed pool of resources, the compiler must reduce resources required by each warp(primarily registers). MIN_WARPS_PER_EU is optional and defaults to 1 if not specified. Specifying a MIN_WARPS_PER_EU greater than the default 1 effectively constrains the compiler's resource usage. +The threads-per-block is the product of (blockDim.x * blockDim.y * blockDim.z). +- MIN_WARPS_PER_EXECUTION_UNIT - directs the compiler to minimize resource usage so that the requested number of warps can be simultaneously active on a multi-processor. Since active warps compete for the same fixed pool of resources, the compiler must reduce resources required by each warp(primarily registers). MIN_WARPS_PER_EXECUTION_UNIT is optional and defaults to 1 if not specified. Specifying a MIN_WARPS_PER_EXECUTION_UNIT greater than the default 1 effectively constrains the compiler's resource usage. When launch kernel with HIP APIs, for example, hipModuleLaunchKernel(), HIP will do validation to make sure input kernel dimension size is not larger than specified launch_bounds. In case exceeded, HIP would return launch failure, if AMD_LOG_LEVEL is set with proper value (for details, please refer to docs/markdown/hip_logging.md), detail information will be shown in the error log message, including @@ -748,15 +745,15 @@ The compiler uses these parameters as follows: - Compilation fails if compiler cannot generate a kernel which meets the requirements of the specified launch bounds. - From MAX_THREADS_PER_BLOCK, the compiler derives the maximum number of warps/block that can be used at launch time. Values of MAX_THREADS_PER_BLOCK less than the default allows the compiler to use a larger pool of registers : each warp uses registers, and this hint constains the launch to a warps/block size which is less than maximum. -- From MIN_WARPS_PER_EU, the compiler derives a maximum number of registers that can be used by the kernel (to meet the required #simultaneous active blocks). -If MIN_WARPS_PER_EU is 1, then the kernel can use all registers supported by the multiprocessor. +- From MIN_WARPS_PER_EXECUTION_UNIT, the compiler derives a maximum number of registers that can be used by the kernel (to meet the required #simultaneous active blocks). +If MIN_WARPS_PER_EXECUTION_UNIT is 1, then the kernel can use all registers supported by the multiprocessor. - The compiler ensures that the registers used in the kernel is less than both allowed maximums, typically by spilling registers (to shared or global memory), or by using more instructions. - The compiler may use hueristics to increase register usage, or may simply be able to avoid spilling. The MAX_THREADS_PER_BLOCK is particularly useful in this cases, since it allows the compiler to use more registers and avoid situations where the compiler constrains the register usage (potentially spilling) to meet the requirements of a large block size that is never used at launch time. ### CU and EU Definitions A compute unit (CU) is responsible for executing the waves of a work-group. It is composed of one or more execution units (EU) which are responsible for executing waves. An EU can have enough resources to maintain the state of more than one executing wave. This allows an EU to hide latency by switching between waves in a similar way to symmetric multithreading on a CPU. In order to allow the state for multiple waves to fit on an EU, the resources used by a single wave have to be limited. Limiting such resources can allow greater latency hiding, but can result in having to spill some register state to memory. This attribute allows an advanced developer to tune the number of waves that are capable of fitting within the resources of an EU. It can be used to ensure at least a certain number will fit to help hide latency, and can also be used to ensure no more than a certain number will fit to limit cache thrashing. - + ### Porting from CUDA __launch_bounds CUDA defines a __launch_bounds which is also designed to control occupancy: ``` diff --git a/docs/markdown/hip_porting_driver_api.md b/docs/markdown/hip_porting_driver_api.md index 90ff2174ce..08df9aaa8a 100644 --- a/docs/markdown/hip_porting_driver_api.md +++ b/docs/markdown/hip_porting_driver_api.md @@ -97,7 +97,7 @@ hip-clang generates initializatiion and termination functions for each translati hip-clang emits a global variable `__hip_gpubin_handle` of void** type with linkonce linkage and inital value 0 for each host translation unit. Each initialization function checks `__hip_gpubin_handle` and register the fatbinary only if `__hip_gpubin_handle` is 0 and saves the return value of `__hip_gpubin_handle` to `__hip_gpubin_handle`. This is to guarantee that the fatbinary is only registered once. Similar check is done in the termination functions. #### Kernel Launching -hip-clang supports kernel launching by CUDA `<<<>>>` syntax, hipLaunchKernel, and hipLaunchKernelGGL. The latter two are macros which expand to CUDA `<<<>>>` syntax. +hip-clang supports kernel launching by CUDA `<<<>>>` syntax, hipLaunchKernelGGL. The latter one is macro which expand to CUDA `<<<>>>` syntax. When the executable or shared library is loaded by the dynamic linker, the initilization functions are called. In the initialization functions, when `__hipRegisterFatBinary` is called, the code objects containing all kernels are loaded; when `__hipRegisterFunction` is called, the stub functions are associated with the corresponding kernels in code objects. @@ -105,8 +105,6 @@ hip-clang implements two sets of kernel launching APIs. By default, in the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function. In hipLaunchByPtr, the real kernel associated with the stub function is launched. -If HIP program is compiled with -fhip-new-launch-api, in the host code, for the `<<<>>>` statement, hip-clang first emits call of `__hipPushCallConfiguration` to save the grid dimension, block dimension, shared memory usage and stream to a stack, then emits call of the stub function with the given arguments. In the stub function, `__hipPopCallConfiguration` is called to get the saved grid dimension, block dimension, shared memory usage and stream, then hipLaunchKernel is called with a function pointer to the stub function. In hipLaunchKernel, the real kernel associated with the stub function is launched. - ### NVCC Implementation Notes #### Interoperation between HIP and CUDA Driver diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index 4d9075f702..33f6847f75 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -23,7 +23,7 @@ and provides practical suggestions on how to port CUDA code and work through com * [Table of Architecture Properties](#table-of-architecture-properties) - [Finding HIP](#finding-hip) - [Identifying HIP Runtime](#identifying-hip-runtime) -- [hipLaunchKernel](#hiplaunchkernel) +- [hipLaunchKernelGGL](#hiplaunchkernelGGL) - [Compiler Options](#compiler-options) - [Linking Issues](#linking-issues) * [Linking With hipcc](#linking-with-hipcc) @@ -294,37 +294,10 @@ On Nvidia platform, HIP is just a thin layer on top of CUDA. On non-AMD platform, HIP runtime determines if cuda is available and can be used. If available, HIP_PLATFORM is set to nvidia and underneath CUDA path is used. -## hipLaunchKernel - -hipLaunchKernel is a variadic macro which accepts as parameters the launch configurations (grid dims, group dims, stream, dynamic shared size) followed by a variable number of kernel arguments. -This sequence is then expanded into the appropriate kernel launch syntax depending on the platform. -While this can be a convenient single-line kernel launch syntax, the macro implementation can cause issues when nested inside other macros. For example, consider the following: - -``` -// Will cause compile error: -#define MY_LAUNCH(command, doTrace) \ -{\ - if (doTrace) printf ("TRACE: %s\n", #command); \ - (command); /* The nested ( ) will cause compile error */\ -} - -MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); -``` - -Avoid nesting macro parameters inside parenthesis - here's an alternative that will work: - -``` -#define MY_LAUNCH(command, doTrace) \ -{\ - if (doTrace) printf ("TRACE: %s\n", #command); \ - command;\ -} - -MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); -``` - - +## hipLaunchKernelGGL +hipLaunchKernelGGL is a macro that can serve as an alternative way to launch kernel, which accepts parameters of launch configurations (grid dims, group dims, stream, dynamic shared size) followed by a variable number of kernel arguments. +It can replace <<< >>>, if the user so desires. ## Compiler Options @@ -482,7 +455,7 @@ int main() HIP_ASSERT(hipMalloc((void**)&Ad, SIZE)); HIP_ASSERT(hipMemcpyToSymbol(HIP_SYMBOL(Value), A, SIZE, 0, hipMemcpyHostToDevice)); - hipLaunchKernel(Get, dim3(1,1,1), dim3(LEN,1,1), 0, 0, Ad); + hipLaunchKernelGGL(Get, dim3(1,1,1), dim3(LEN,1,1), 0, 0, Ad); HIP_ASSERT(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); for(unsigned i=0;i inline void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, std::uint32_t sharedMemBytes, hipStream_t stream, diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index f2a24e6dbe..10d08268e8 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -160,7 +160,7 @@ typedef enum hipMemoryType { hipMemoryTypeArray, ///< Array memory, physically located on device. (see deviceId for specific ///< device) hipMemoryTypeUnified ///< Not used currently -}hipMemoryType; +} hipMemoryType; /** @@ -460,7 +460,7 @@ enum hipComputeMode { }; /** - * @} + * @} */ #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__)) @@ -544,67 +544,109 @@ enum hipLimit_t { * @addtogroup GlobalDefs More * @{ */ -//! Flags that can be used with hipStreamCreateWithFlags -#define hipStreamDefault \ - 0x00 ///< Default stream creation flags. These are used with hipStreamCreate(). -#define hipStreamNonBlocking 0x01 ///< Stream does not implicitly synchronize with null stream -//! Flags that can be used with hipEventCreateWithFlags: -#define hipEventDefault 0x0 ///< Default flags -#define hipEventBlockingSync \ - 0x1 ///< Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency. -#define hipEventDisableTiming \ - 0x2 ///< Disable event's capability to record timing information. May improve performance. -#define hipEventInterprocess 0x4 ///< Event can support IPC. @warning - not supported in HIP. -#define hipEventReleaseToDevice \ - 0x40000000 /// < Use a device-scope release when recording this event. This flag is useful to - /// obtain more precise timings of commands between events. The flag is a no-op on - /// CUDA platforms. -#define hipEventReleaseToSystem \ - 0x80000000 /// < Use a system-scope release when recording this event. This flag is - /// useful to make non-coherent host memory visible to the host. The flag is a - /// no-op on CUDA platforms. -//! Flags that can be used with hipHostMalloc +//Flags that can be used with hipStreamCreateWithFlags. +/** Default stream creation flags. These are used with hipStreamCreate().*/ +#define hipStreamDefault 0x00 + +/** Stream does not implicitly synchronize with null stream.*/ +#define hipStreamNonBlocking 0x01 + +//Flags that can be used with hipEventCreateWithFlags. +/** Default flags.*/ +#define hipEventDefault 0x0 + +/** Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency.*/ +#define hipEventBlockingSync 0x1 + +/** Disable event's capability to record timing information. May improve performance.*/ +#define hipEventDisableTiming 0x2 + +/** Event can support IPC. Warnig: It is not supported in HIP.*/ +#define hipEventInterprocess 0x4 + +/** Use a device-scope release when recording this event. This flag is useful to obtain more + * precise timings of commands between events. The flag is a no-op on CUDA platforms.*/ +#define hipEventReleaseToDevice 0x40000000 + +/** Use a system-scope release when recording this event. This flag is useful to make + * non-coherent host memory visible to the host. The flag is a no-op on CUDA platforms.*/ +#define hipEventReleaseToSystem 0x80000000 + +//Flags that can be used with hipHostMalloc. +/** Default pinned memory allocation on the host.*/ #define hipHostMallocDefault 0x0 -#define hipHostMallocPortable 0x1 ///< Memory is considered allocated by all contexts. -#define hipHostMallocMapped \ - 0x2 ///< Map the allocation into the address space for the current device. The device pointer - ///< can be obtained with #hipHostGetDevicePointer. + +/** Memory is considered allocated by all contexts.*/ +#define hipHostMallocPortable 0x1 + +/** Map the allocation into the address space for the current device. The device pointer + * can be obtained with #hipHostGetDevicePointer.*/ +#define hipHostMallocMapped 0x2 + +/** Allocates the memory as write-combined. On some system configurations, write-combined allocation + * may be transferred faster across the PCI Express bus, however, could have low read efficiency by + * most CPUs. It's a good option for data tranfer from host to device via mapped pinned memory.*/ #define hipHostMallocWriteCombined 0x4 -#define hipHostMallocNumaUser \ - 0x20000000 ///< Host memory allocation will follow numa policy set by user -#define hipHostMallocCoherent \ - 0x40000000 ///< Allocate coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific - ///< allocation. -#define hipHostMallocNonCoherent \ - 0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific - ///< allocation. -#define hipMemAttachGlobal 0x01 ///< Memory can be accessed by any stream on any device -#define hipMemAttachHost 0x02 ///< Memory cannot be accessed by any stream on any device -#define hipMemAttachSingle 0x04 ///< Memory can only be accessed by a single stream on - ///< the associated device + +/** Host memory allocation will follow numa policy set by user.*/ +#define hipHostMallocNumaUser 0x20000000 + +/** Allocate coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific allocation.*/ +#define hipHostMallocCoherent 0x40000000 + +/** Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific allocation.*/ +#define hipHostMallocNonCoherent 0x80000000 + +/** Memory can be accessed by any stream on any device*/ +#define hipMemAttachGlobal 0x01 + +/** Memory cannot be accessed by any stream on any device.*/ +#define hipMemAttachHost 0x02 + +/** Memory can only be accessed by a single stream on the associated device.*/ +#define hipMemAttachSingle 0x04 + #define hipDeviceMallocDefault 0x0 -#define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device. -#define hipMallocSignalMemory 0x2 ///< Memory represents a HSA signal. -//! Flags that can be used with hipHostRegister -#define hipHostRegisterDefault 0x0 ///< Memory is Mapped and Portable -#define hipHostRegisterPortable 0x1 ///< Memory is considered registered by all contexts. -#define hipHostRegisterMapped \ - 0x2 ///< Map the allocation into the address space for the current device. The device pointer - ///< can be obtained with #hipHostGetDevicePointer. -#define hipHostRegisterIoMemory 0x4 ///< Not supported. -#define hipExtHostRegisterCoarseGrained 0x8 ///< Coarse Grained host memory lock -#define hipDeviceScheduleAuto 0x0 ///< Automatically select between Spin and Yield -#define hipDeviceScheduleSpin \ - 0x1 ///< Dedicate a CPU core to spin-wait. Provides lowest latency, but burns a CPU core and - ///< may consume more power. -#define hipDeviceScheduleYield \ - 0x2 ///< Yield the CPU to the operating system when waiting. May increase latency, but lowers - ///< power and is friendlier to other threads in the system. + +/** Memory is allocated in fine grained region of device.*/ +#define hipDeviceMallocFinegrained 0x1 + +/** Memory represents a HSA signal.*/ +#define hipMallocSignalMemory 0x2 + +//Flags that can be used with hipHostRegister. +/** Memory is Mapped and Portable.*/ +#define hipHostRegisterDefault 0x0 + +/** Memory is considered registered by all contexts.*/ +#define hipHostRegisterPortable 0x1 + +/** Map the allocation into the address space for the current device. The device pointer + * can be obtained with #hipHostGetDevicePointer.*/ +#define hipHostRegisterMapped 0x2 + +/** Not supported.*/ +#define hipHostRegisterIoMemory 0x4 + +/** Coarse Grained host memory lock.*/ +#define hipExtHostRegisterCoarseGrained 0x8 + +/** Automatically select between Spin and Yield.*/ +#define hipDeviceScheduleAuto 0x0 + +/** Dedicate a CPU core to spin-wait. Provides lowest latency, but burns a CPU core and may + * consume more power.*/ +#define hipDeviceScheduleSpin 0x1 + +/** Yield the CPU to the operating system when waiting. May increase latency, but lowers power + * and is friendlier to other threads in the system.*/ +#define hipDeviceScheduleYield 0x2 #define hipDeviceScheduleBlockingSync 0x4 #define hipDeviceScheduleMask 0x7 #define hipDeviceMapHost 0x8 #define hipDeviceLmemResizeToMax 0x16 -#define hipArrayDefault 0x00 ///< Default HIP array allocation flag +/** Default HIP array allocation flag.*/ +#define hipArrayDefault 0x00 #define hipArrayLayered 0x01 #define hipArraySurfaceLoadStore 0x02 #define hipArrayCubemap 0x04 @@ -614,15 +656,17 @@ enum hipLimit_t { #define hipCooperativeLaunchMultiDeviceNoPostSync 0x02 #define hipCpuDeviceId ((int)-1) #define hipInvalidDeviceId ((int)-2) -// Flags that can be used with hipExtLaunch Set of APIs -#define hipExtAnyOrderLaunch 0x01 ///< AnyOrderLaunch of kernels -// Flags to be used with hipStreamWaitValue32 and hipStreamWaitValue64 +//Flags that can be used with hipExtLaunch Set of APIs. +/** AnyOrderLaunch of kernels.*/ +#define hipExtAnyOrderLaunch 0x01 +// Flags to be used with hipStreamWaitValue32 and hipStreamWaitValue64. #define hipStreamWaitValueGte 0x0 #define hipStreamWaitValueEq 0x1 #define hipStreamWaitValueAnd 0x2 #define hipStreamWaitValueNor 0x3 // Stream per thread -#define hipStreamPerThread ((hipStream_t)2) ///< Implicit stream per application thread +/** Implicit stream per application thread.*/ +#define hipStreamPerThread ((hipStream_t)2) /* * @brief HIP Memory Advise values * @enum @@ -3953,6 +3997,26 @@ hipError_t hipLaunchKernel(const void* function_address, */ hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy); //TODO: Move this to hip_ext.h +/** + * @brief Launches kernel from the pointer address, with arguments and shared memory on stream. + * + * @param [in] function_address pointer to the Kernel to launch. + * @param [in] numBlocks number of blocks. + * @param [in] dimBlocks dimension of a block. + * @param [in] args pointer to kernel arguments. + * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. + * HIP-Clang compiler provides support for extern shared declarations. + * @param [in] stream Stream where the kernel should be dispatched. + * @param [in] startEvent If non-null, specified event will be updated to track the start time of + * the kernel launch. The event must be created before calling this API. + * @param [in] stopEvent If non-null, specified event will be updated to track the stop time of + * the kernel launch. The event must be created before calling this API. + * May be 0, in which case the default stream is used with associated synchronization rules. + * @param [in] flags. The value of hipExtAnyOrderLaunch, signifies if kernel can be + * launched in any order. + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue. + * + */ hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks, void** args, size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent, hipEvent_t stopEvent, int flags); @@ -3967,69 +4031,9 @@ hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 * @{ * This section describes the texture management functions of HIP runtime API. */ -hipError_t hipBindTextureToMipmappedArray( - const textureReference* tex, - hipMipmappedArray_const_t mipmappedArray, - const hipChannelFormatDesc* desc); -hipError_t hipGetTextureReference( - const textureReference** texref, - const void* symbol); -hipError_t hipCreateTextureObject( - hipTextureObject_t* pTexObject, - const hipResourceDesc* pResDesc, - const hipTextureDesc* pTexDesc, - const struct hipResourceViewDesc* pResViewDesc); -hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject); -hipError_t hipGetChannelDesc( - hipChannelFormatDesc* desc, - hipArray_const_t array); -hipError_t hipGetTextureObjectResourceDesc( - hipResourceDesc* pResDesc, - hipTextureObject_t textureObject); -hipError_t hipGetTextureObjectResourceViewDesc( - struct hipResourceViewDesc* pResViewDesc, - hipTextureObject_t textureObject); -hipError_t hipGetTextureObjectTextureDesc( - hipTextureDesc* pTexDesc, - hipTextureObject_t textureObject); -hipError_t hipTexRefSetAddressMode( - textureReference* texRef, - int dim, - enum hipTextureAddressMode am); -hipError_t hipTexRefSetArray( - textureReference* tex, - hipArray_const_t array, - unsigned int flags); -hipError_t hipTexRefSetFilterMode( - textureReference* texRef, - enum hipTextureFilterMode fm); -hipError_t hipTexRefSetFlags( - textureReference* texRef, - unsigned int Flags); -hipError_t hipTexRefSetFormat( - textureReference* texRef, - hipArray_Format fmt, - int NumPackedComponents); -hipError_t hipTexObjectCreate( - hipTextureObject_t* pTexObject, - const HIP_RESOURCE_DESC* pResDesc, - const HIP_TEXTURE_DESC* pTexDesc, - const HIP_RESOURCE_VIEW_DESC* pResViewDesc); -hipError_t hipTexObjectDestroy( - hipTextureObject_t texObject); -hipError_t hipTexObjectGetResourceDesc( - HIP_RESOURCE_DESC* pResDesc, - hipTextureObject_t texObject); -hipError_t hipTexObjectGetResourceViewDesc( - HIP_RESOURCE_VIEW_DESC* pResViewDesc, - hipTextureObject_t texObject); -hipError_t hipTexObjectGetTextureDesc( - HIP_TEXTURE_DESC* pTexDesc, - hipTextureObject_t texObject); - /** * - * @addtogroup TexturD Texture Management [Deprecated] + * @addtogroup TextureD Texture Management [Deprecated] * @{ * @ingroup Texture * This section describes the deprecated texture management functions of HIP runtime API. @@ -4124,6 +4128,73 @@ hipError_t hipTexRefSetMaxAnisotropy( /** * @} */ +hipError_t hipBindTextureToMipmappedArray( + const textureReference* tex, + hipMipmappedArray_const_t mipmappedArray, + const hipChannelFormatDesc* desc); + hipError_t hipGetTextureReference( + const textureReference** texref, + const void* symbol); +hipError_t hipCreateTextureObject( + hipTextureObject_t* pTexObject, + const hipResourceDesc* pResDesc, + const hipTextureDesc* pTexDesc, + const struct hipResourceViewDesc* pResViewDesc); +hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject); +hipError_t hipGetChannelDesc( + hipChannelFormatDesc* desc, + hipArray_const_t array); +hipError_t hipGetTextureObjectResourceDesc( + hipResourceDesc* pResDesc, + hipTextureObject_t textureObject); +hipError_t hipGetTextureObjectResourceViewDesc( + struct hipResourceViewDesc* pResViewDesc, + hipTextureObject_t textureObject); +hipError_t hipGetTextureObjectTextureDesc( + hipTextureDesc* pTexDesc, + hipTextureObject_t textureObject); + +hipError_t hipTexRefSetAddressMode( + textureReference* texRef, + int dim, + enum hipTextureAddressMode am); +hipError_t hipTexRefSetArray( + textureReference* tex, + hipArray_const_t array, + unsigned int flags); +hipError_t hipTexRefSetFilterMode( + textureReference* texRef, + enum hipTextureFilterMode fm); +hipError_t hipTexRefSetFlags( + textureReference* texRef, + unsigned int Flags); +hipError_t hipTexRefSetFormat( + textureReference* texRef, + hipArray_Format fmt, + int NumPackedComponents); +hipError_t hipTexObjectCreate( + hipTextureObject_t* pTexObject, + const HIP_RESOURCE_DESC* pResDesc, + const HIP_TEXTURE_DESC* pTexDesc, + const HIP_RESOURCE_VIEW_DESC* pResViewDesc); +hipError_t hipTexObjectDestroy( + hipTextureObject_t texObject); +hipError_t hipTexObjectGetResourceDesc( + HIP_RESOURCE_DESC* pResDesc, + hipTextureObject_t texObject); +hipError_t hipTexObjectGetResourceViewDesc( + HIP_RESOURCE_VIEW_DESC* pResViewDesc, + hipTextureObject_t texObject); +hipError_t hipTexObjectGetTextureDesc( + HIP_TEXTURE_DESC* pTexDesc, + hipTextureObject_t texObject); +// doxygen end Texture management +/** + * @addtogroup TextureU Texture Management [Unsupported] + * @{ + * @ingroup Texture + * This section describes the unsupported texture management functions of HIP runtime API. + */ // The following are not supported. /** @@ -4161,11 +4232,10 @@ hipError_t hipMipmappedArrayGetLevel( hipArray_t* pLevelArray, hipMipmappedArray_t hMipMappedArray, unsigned int level); -// doxygen end Texture management unsupported +// doxygen end unsuppported texture management /** * @} */ - // doxygen end Texture management /** * @} @@ -4207,7 +4277,7 @@ int hipGetStreamDeviceId(hipStream_t stream); /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- - * @defgroup Graph Management + * @defgroup Graph Graph Management * @{ * This section describes the graph management types & functions of HIP runtime API. */ @@ -5247,6 +5317,7 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph * @} */ + /** *------------------------------------------------------------------------------------------------- *-------------------------------------------------------------------------------------------------