diff --git a/.wordlist.txt b/.wordlist.txt index 3bdda5f630..2033214571 100644 --- a/.wordlist.txt +++ b/.wordlist.txt @@ -87,6 +87,7 @@ ltrace makefile Malloc malloc +memset multicore multigrid multithreading @@ -105,7 +106,9 @@ overindexing oversubscription pixelated pragmas +preallocated preconditioners +predefining prefetched preprocessor PTX diff --git a/docs/data/how-to/hipgraph/hip_graph.drawio b/docs/data/how-to/hipgraph/hip_graph.drawio new file mode 100644 index 0000000000..03569ac734 --- /dev/null +++ b/docs/data/how-to/hipgraph/hip_graph.drawio @@ -0,0 +1,76 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/docs/data/how-to/hipgraph/hip_graph.svg b/docs/data/how-to/hipgraph/hip_graph.svg new file mode 100644 index 0000000000..6eed6b92e5 --- /dev/null +++ b/docs/data/how-to/hipgraph/hip_graph.svg @@ -0,0 +1,4 @@ + + + +Stream 1
Kernel B
Kernel BStream 2
Kernel A
Kernel A
hipDeviceSynchronize
hipDeviceSynchronize
Kernel C
Kernel C
hipDeviceSynchronize
hipDeviceSynchronize
Kernel D
Kernel D
Kernel A
Kernel A
Kernel B
Kernel B
Kernel C
Kernel C
Kernel D
Kernel DText is not SVG - cannot display \ No newline at end of file diff --git a/docs/data/how-to/hipgraph/hip_graph_speedup.drawio b/docs/data/how-to/hipgraph/hip_graph_speedup.drawio new file mode 100644 index 0000000000..7802785f6b --- /dev/null +++ b/docs/data/how-to/hipgraph/hip_graph_speedup.drawio @@ -0,0 +1,162 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/docs/data/how-to/hipgraph/hip_graph_speedup.svg b/docs/data/how-to/hipgraph/hip_graph_speedup.svg new file mode 100644 index 0000000000..f16123b9e2 --- /dev/null +++ b/docs/data/how-to/hipgraph/hip_graph_speedup.svg @@ -0,0 +1,4 @@ + + + +Streams
kernel A
kernel launch A
kernel B
kernel C
kernel launch B
kernel launch C
host activity
device activity
time
kernel launch D
kernel D
device idling due to kernel launch congestion
kernel A
kernel B
kernel C
graph launch
host activity
device activity
kernel D
Graph
speedup
 \ No newline at end of file diff --git a/docs/how-to/hipgraph.rst b/docs/how-to/hipgraph.rst new file mode 100644 index 0000000000..83d6fa0f61 --- /dev/null +++ b/docs/how-to/hipgraph.rst @@ -0,0 +1,517 @@ +.. meta:: + :description: This chapter describes how to use HIP graphs and highlights their use cases. + :keywords: ROCm, HIP, graph, stream + +.. _how_to_HIP_graph: + +******************************************************************************** +HIP graphs +******************************************************************************** + +.. note:: + The HIP graph API is currently in Beta. Some features can change and might + have outstanding issues. Not all features supported by CUDA graphs are yet + supported. For a list of all currently supported functions see the + :doc:`HIP graph API documentation<../doxygen/html/group___graph>`. + +HIP graphs are an alternative way of executing tasks on a GPU that can provide +performance benefits over launching kernels using the standard +method via streams. A HIP graph is made up of nodes and edges. The nodes of a +HIP graph represent the operations performed, while the edges mark dependencies +between those operations. + +The nodes can be one of the following: + +- empty nodes +- nested graphs +- kernel launches +- host-side function calls +- HIP memory functions (copy, memset, ...) +- HIP events +- signalling or waiting on external semaphores + +.. note:: + The available node types are specified by :cpp:enumerator:`hipGraphNodeType`. + +The following figure visualizes the concept of graphs, compared to using streams. + +.. figure:: ../data/how-to/hipgraph/hip_graph.svg + :alt: Diagram depicting the difference between using streams to execute + kernels with dependencies, resolved by explicitly synchronizing, + or using graphs, where the edges denote the dependencies. + +The standard method of launching kernels incurs a small overhead for each +iteration of the operation involved. That overhead is negligible, when the +kernel is launched directly with the HIP C/C++ API, but depending on the +framework used, there can be several levels of redirection, until the actual +kernel is launched by the HIP runtime, leading to significant overhead. +Especially for some AI frameworks, a GPU kernel might run faster than the time +it takes for the framework to set up and launch the kernel, and so the overhead +of repeatedly launching kernels can have a significant impact on performance. + +HIP graphs are designed to address this issue, by predefining the HIP API calls +and their dependencies with a graph, and performing most of the initialization +beforehand. Launching a graph only requires a single call, after which the +HIP runtime takes care of executing the operations within the graph. +Graphs can provide additional performance benefits, by enabling optimizations +that are only possible when knowing the dependencies between the operations. + +.. figure:: ../data/how-to/hipgraph/hip_graph_speedup.svg + :alt: Diagram depicting the speed up achievable with HIP graphs compared to + HIP streams when launching many short-running kernels. + + Qualitative presentation of the execution time of many short-running kernels + when launched using HIP stream versus HIP graph. This does not include the + time needed to set up the graph. + + +******************************************************************************** +Using HIP graphs +******************************************************************************** + +There are two different ways of creating graphs: Capturing kernel launches from +a stream, or explicitly creating graphs. The difference between the two +approaches is explained later in this chapter. + +The general flow for using HIP graphs includes the following steps. + +#. Create a :cpp:type:`hipGraph_t` graph template using one of the two approaches described in this chapter +#. Create a :cpp:type:`hipGraphExec_t` executable instance of the graph template using :cpp:func:`hipGraphInstantiate` +#. Use :cpp:func:`hipGraphLaunch` to launch the executable graph to a stream +#. After execution completes free and destroy graph resources + +The first two steps are the initial setup and only need to be executed once. First +step is the definition of the operations (nodes) and the dependencies (edges) +between them. The second step is the instantiation of the graph. This takes care +of validating and initializing the graph, to reduce the overhead when executing +the graph. The third step is the execution of the graph, which takes care of +launching all the kernels and executing the operations while respecting their +dependencies and necessary synchronizations as specified. + +Because HIP graphs require some setup and initialization overhead before their +first execution, graphs only provide a benefit for workloads that require +many iterations to complete. + +In both methods the :cpp:type:`hipGraph_t` template for a graph is used to define the graph. +In order to actually launch a graph, the template needs to be instantiated using +:cpp:func:`hipGraphInstantiate`, which results in an executable graph of type :cpp:type:`hipGraphExec_t`. +This executable graph can then be launched with :cpp:func:`hipGraphLaunch`, replaying the +operations within the graph. Note, that launching graphs is fundamentally no +different to executing other HIP functions on a stream, except for the fact, +that scheduling the operations within the graph encompasses less overhead and +can enable some optimizations, but they still need to be associated with a stream for execution. + +Memory management +----------------- + +Memory that is used by operations in graphs can either be pre-allocated or +managed within the graph. Graphs can contain nodes that take care of allocating +memory on the device or copying memory between the host and the device. +Whether you want to pre-allocate the memory or manage it within the graph +depends on the use-case. If the graph is executed in a tight loop the +performance is usually better when the memory is preallocated, so that it +does not need to be reallocated in every iteration. + +The same rules as for normal memory allocations apply for memory allocated and +freed by nodes, meaning that the nodes that access memory allocated in a graph +must be ordered after allocation and before freeing. + +Memory management within the graph enables the runtime to take care of memory reuse and optimizations. +The lifetime of memory managed in a graph begins when the execution reaches the +node allocating the memory, and ends when either reaching the corresponding +free node within the graph, or after graph execution when a corresponding +:cpp:func:`hipFreeAsync` or :cpp:func:`hipFree` call is reached. +The memory can also be freed with a free node in a different graph that is +associated with the same memory address. + +Unlike device memory that is not associated with a graph, this does not necessarily +mean that the freed memory is returned back to the operating system immediately. +Graphs can retain a memory pool for quickly reusing memory within the graph. +This can be especially useful when memory is freed and reallocated later on +within a graph, as that memory doesn't have to be requested from the operating system. +It also potentially reduces the total memory footprint of the graph, by reusing the same memory. + +The amount of memory allocated for graph memory pools on a specific device can +be queried using :cpp:func:`hipDeviceGetGraphMemAttribute`. +In order to return the freed memory :cpp:func:`hipDeviceGraphMemTrim` can be used. +This will return any memory that is not in active use by graphs. + +These memory allocations can also be set up to allow access from multiple GPUs, +just like normal allocations. HIP then takes care of allocating and mapping the +memory to the GPUs. When capturing a graph from a stream, the node sets the +accessibility according to :cpp:func:`hipMemPoolSetAccess` at the time of capturing. + + +Capture graphs from a stream +================================================================================ + +The easy way to integrate HIP graphs into already existing code is to use +:cpp:func:`hipStreamBeginCapture` and :cpp:func:`hipStreamEndCapture` to obtain a :cpp:type:`hipGraph_t` +graph template that includes the captured operations. + +When starting to capture operations for a graph using :cpp:func:`hipStreamBeginCapture`, +the operations assigned to the stream are captured into a graph instead of being +executed. The associated graph is returned when calling :cpp:func:`hipStreamEndCapture`, which +also stops capturing operations. +In order to capture to an already existing graph use :cpp:func:`hipStreamBeginCaptureToGraph`. + +The functions assigned to the capturing stream are not executed, but instead are +captured and defined as nodes in the graph, to be run when the instantiated +graph is launched. + +Functions must be associated with a stream in order to be captured. +This means that non-HIP API functions are not captured by default, but are +executed as standard functions when encountered and not added to the graph. +In order to assign host functions to a stream use +:cpp:func:`hipLaunchHostFunc`, as shown in the following code example. +They will then be captured and defined as a host node in the resulting graph, +and won't be executed when encountered. + +Synchronous HIP API calls that are implicitly assigned to the default stream are +not permitted while capturing a stream and will return an error. This is +because they implicitly synchronize and cause a dependency that can not be +captured within the stream. This includes functions like :cpp:func:`hipMalloc`, +:cpp:func:`hipMemcpy` and :cpp:func:`hipFree`. In order to capture these to the stream, replace +them with the corresponding asynchronous calls like :cpp:func:`hipMallocAsync`, :cpp:func:`hipMemcpyAsync` or :cpp:func:`hipFreeAsync`. + +The general flow for using stream capture to create a graph template is: + +#. Create a stream from which to capture the operations + +#. Call :cpp:func:`hipStreamBeginCapture` before the first operation to be captured + +#. Call :cpp:func:`hipStreamEndCapture` after the last operation to be captured + + #. Define a :cpp:type:`hipGraph_t` graph template to which :cpp:func:`hipStreamEndCapture` + passes the captured graph + +The following code is an example of how to use the HIP graph API to capture a +graph from a stream. + +.. code-block:: cpp + #include + #include + #include + + #define HIP_CHECK(expression) \ + { \ + const hipError_t status = expression; \ + if(status != hipSuccess){ \ + std::cerr << "HIP error " \ + << status << ": " \ + << hipGetErrorString(status) \ + << " at " << __FILE__ << ":" \ + << __LINE__ << std::endl; \ + } \ + } + + + __global__ void kernelA(double* arrayA, size_t size){ + const size_t x = threadIdx.x + blockDim.x * blockIdx.x; + if(x < size){arrayA[x] *= 2.0;} + }; + __global__ void kernelB(int* arrayB, size_t size){ + const size_t x = threadIdx.x + blockDim.x * blockIdx.x; + if(x < size){arrayB[x] = 3;} + }; + __global__ void kernelC(double* arrayA, const int* arrayB, size_t size){ + const size_t x = threadIdx.x + blockDim.x * blockIdx.x; + if(x < size){arrayA[x] += arrayB[x];} + }; + + struct set_vector_args{ + std::vector& h_array; + double value; + }; + + void set_vector(void* args){ + set_vector_args h_args{*(reinterpret_cast(args))}; + + std::vector& vec{h_args.h_array}; + vec.assign(vec.size(), h_args.value); + } + + int main(){ + constexpr int numOfBlocks = 1024; + constexpr int threadsPerBlock = 1024; + constexpr size_t arraySize = 1U << 20; + + // This example assumes that kernelA operates on data that needs to be initialized on + // and copied from the host, while kernelB initializes the array that is passed to it. + // Both arrays are then used as input to kernelC, where arrayA is also used as + // output, that is copied back to the host, while arrayB is only read from and not modified. + + double* d_arrayA; + int* d_arrayB; + std::vector h_array(arraySize); + constexpr double initValue = 2.0; + + hipStream_t captureStream; + HIP_CHECK(hipStreamCreate(&captureStream)); + + // Start capturing the operations assigned to the stream + HIP_CHECK(hipStreamBeginCapture(captureStream, hipStreamCaptureModeGlobal)); + + // hipMallocAsync and hipMemcpyAsync are needed, to be able to assign it to a stream + HIP_CHECK(hipMallocAsync(&d_arrayA, arraySize*sizeof(double), captureStream)); + HIP_CHECK(hipMallocAsync(&d_arrayB, arraySize*sizeof(int), captureStream)); + + // Assign host function to the stream + // Needs a custom struct to pass the arguments + set_vector_args args{h_array, initValue}; + HIP_CHECK(hipLaunchHostFunc(captureStream, set_vector, &args)); + + HIP_CHECK(hipMemcpyAsync(d_arrayA, h_array.data(), arraySize*sizeof(double), hipMemcpyHostToDevice, captureStream)); + + kernelA<<>>(d_arrayA, arraySize); + kernelB<<>>(d_arrayB, arraySize); + kernelC<<>>(d_arrayA, d_arrayB, arraySize); + + HIP_CHECK(hipMemcpyAsync(h_array.data(), d_arrayA, arraySize*sizeof(*d_arrayA), hipMemcpyDeviceToHost, captureStream)); + + HIP_CHECK(hipFreeAsync(d_arrayA, captureStream)); + HIP_CHECK(hipFreeAsync(d_arrayB, captureStream)); + + // Stop capturing + hipGraph_t graph; + HIP_CHECK(hipStreamEndCapture(captureStream, &graph)); + + // Create an executable graph from the captured graph + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + // The graph template can be deleted after the instantiation if it's not needed for later use + HIP_CHECK(hipGraphDestroy(graph)); + + // Actually launch the graph. The stream does not have + // to be the same as the one used for capturing. + HIP_CHECK(hipGraphLaunch(graphExec, captureStream)); + + // Verify results + constexpr double expected = initValue * 2.0 + 3; + bool passed = true; + for(size_t i = 0; i < arraySize; ++i){ + if(h_array[i] != expected){ + passed = false; + std::cerr << "Validation failed! Expected " << expected << " got " << h_array[0] << std::endl; + break; + } + } + if(passed){ + std::cerr << "Validation passed." << std::endl; + } + + // Free graph and stream resources after usage + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(captureStream)); + } + +Explicit graph creation +================================================================================ + +Graphs can also be created directly using the HIP graph API, giving more +fine-grained control over the graph. In this case, the graph nodes are created +explicitly, together with their parameters and dependencies, which specify the +edges of the graph, thereby forming the graph structure. + +The nodes are represented by the generic :cpp:type:`hipGraphNode_t` type. The actual +node type is implicitly defined by the specific function used to add the node to +the graph, for example :cpp:func:`hipGraphAddKernelNode` See the +:doc:`HIP graph API documentation<../doxygen/html/group___graph>` for the +available functions, they are of type ``hipGraphAdd{Type}Node``. Each type of +node also has a predefined set of parameters depending on the operation, for +example :cpp:class:`hipKernelNodeParams` for a kernel launch. See the +:doc:`documentation for the general hipGraphNodeParams type<../doxygen/html/structhip_graph_node_params>` +for a list of available parameter types and their members. + +The general flow for explicitly creating a graph is usually: + +#. Create a graph :cpp:type:`hipGraph_t` + +#. Create the nodes and their parameters and add them to the graph + + #. Define a :cpp:type:`hipGraphNode_t` + + #. Define the parameter struct for the desired operation, by explicitly setting the appropriate struct's members. + + #. Use the appropriate ``hipGraphAdd{Type}Node`` function to add the node to the graph. + + #. The dependencies can be defined when adding the node to the graph, or afterwards by using :cpp:func:`hipGraphAddDependencies` + +The following code example demonstrates how to explicitly create nodes in order to create a graph. + +.. code-block:: cpp + + #include + #include + #include + + #define HIP_CHECK(expression) \ + { \ + const hipError_t status = expression; \ + if(status != hipSuccess){ \ + std::cerr << "HIP error " \ + << status << ": " \ + << hipGetErrorString(status) \ + << " at " << __FILE__ << ":" \ + << __LINE__ << std::endl; \ + } \ + } + + __global__ void kernelA(double* arrayA, size_t size){ + const size_t x = threadIdx.x + blockDim.x * blockIdx.x; + if(x < size){arrayA[x] *= 2.0;} + }; + __global__ void kernelB(int* arrayB, size_t size){ + const size_t x = threadIdx.x + blockDim.x * blockIdx.x; + if(x < size){arrayB[x] = 3;} + }; + __global__ void kernelC(double* arrayA, const int* arrayB, size_t size){ + const size_t x = threadIdx.x + blockDim.x * blockIdx.x; + if(x < size){arrayA[x] += arrayB[x];} + }; + + struct set_vector_args{ + std::vector& h_array; + double value; + }; + + void set_vector(void* args){ + set_vector_args h_args{*(reinterpret_cast(args))}; + + std::vector& vec{h_args.h_array}; + vec.assign(vec.size(), h_args.value); + } + + int main(){ + constexpr int numOfBlocks = 1024; + constexpr int threadsPerBlock = 1024; + size_t arraySize = 1U << 20; + + // The pointers to the device memory don't need to be declared here, + // they are contained within the hipMemAllocNodeParams as the dptr member + std::vector h_array(arraySize); + constexpr double initValue = 2.0; + + // Create graph an empty graph + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Parameters to allocate arrays + hipMemAllocNodeParams allocArrayAParams{}; + allocArrayAParams.poolProps.allocType = hipMemAllocationTypePinned; + allocArrayAParams.poolProps.location.type = hipMemLocationTypeDevice; + allocArrayAParams.poolProps.location.id = 0; // GPU on which memory resides + allocArrayAParams.bytesize = arraySize * sizeof(double); + + hipMemAllocNodeParams allocArrayBParams{}; + allocArrayBParams.poolProps.allocType = hipMemAllocationTypePinned; + allocArrayBParams.poolProps.location.type = hipMemLocationTypeDevice; + allocArrayBParams.poolProps.location.id = 0; // GPU on which memory resides + allocArrayBParams.bytesize = arraySize * sizeof(int); + + // Add the allocation nodes to the graph. They don't have any dependencies + hipGraphNode_t allocNodeA, allocNodeB; + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocArrayAParams)); + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeB, graph, nullptr, 0, &allocArrayBParams)); + + // Parameters for the host function + // Needs custom struct to pass the arguments + set_vector_args args{h_array, initValue}; + hipHostNodeParams hostParams{}; + hostParams.fn = set_vector; + hostParams.userData = static_cast(&args); + + // Add the host node that initializes the host array. It also doesn't have any dependencies + hipGraphNode_t hostNode; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); + + // Add memory copy node, that copies the initialized host array to the device. + // It has to wait for the host array to be initialized and the device memory to be allocated + hipGraphNode_t cpyNodeDependencies[] = {allocNodeA, hostNode}; + hipGraphNode_t cpyToDevNode; + HIP_CHECK(hipGraphAddMemcpyNode1D(&cpyToDevNode, graph, cpyNodeDependencies, 1, allocArrayAParams.dptr, h_array.data(), arraySize * sizeof(double), hipMemcpyHostToDevice)); + + // Parameters for kernelA + hipKernelNodeParams kernelAParams; + void* kernelAArgs[] = {&allocArrayAParams.dptr, static_cast(&arraySize)}; + kernelAParams.func = reinterpret_cast(kernelA); + kernelAParams.gridDim = numOfBlocks; + kernelAParams.blockDim = threadsPerBlock; + kernelAParams.sharedMemBytes = 0; + kernelAParams.kernelParams = kernelAArgs; + kernelAParams.extra = nullptr; + + // Add the node for kernelA. It has to wait for the memory copy to finish, as it depends on the values from the host array. + hipGraphNode_t kernelANode; + HIP_CHECK(hipGraphAddKernelNode(&kernelANode, graph, &cpyToDevNode, 1, &kernelAParams)); + + // Parameters for kernelB + hipKernelNodeParams kernelBParams; + void* kernelBArgs[] = {&allocArrayBParams.dptr, static_cast(&arraySize)}; + kernelBParams.func = reinterpret_cast(kernelB); + kernelBParams.gridDim = numOfBlocks; + kernelBParams.blockDim = threadsPerBlock; + kernelBParams.sharedMemBytes = 0; + kernelBParams.kernelParams = kernelBArgs; + kernelBParams.extra = nullptr; + + // Add the node for kernelB. It only has to wait for the memory to be allocated, as it initializes the array. + hipGraphNode_t kernelBNode; + HIP_CHECK(hipGraphAddKernelNode(&kernelBNode, graph, &allocNodeB, 1, &kernelBParams)); + + // Parameters for kernelC + hipKernelNodeParams kernelCParams; + void* kernelCArgs[] = {&allocArrayAParams.dptr, &allocArrayBParams.dptr, static_cast(&arraySize)}; + kernelCParams.func = reinterpret_cast(kernelC); + kernelCParams.gridDim = numOfBlocks; + kernelCParams.blockDim = threadsPerBlock; + kernelCParams.sharedMemBytes = 0; + kernelCParams.kernelParams = kernelCArgs; + kernelCParams.extra = nullptr; + + // Add the node for kernelC. It has to wait on both kernelA and kernelB to finish, as it depends on their results. + hipGraphNode_t kernelCNode; + hipGraphNode_t kernelCDependencies[] = {kernelANode, kernelBNode}; + HIP_CHECK(hipGraphAddKernelNode(&kernelCNode, graph, kernelCDependencies, 1, &kernelCParams)); + + // Copy the results back to the host. Has to wait for kernelC to finish. + hipGraphNode_t cpyToHostNode; + HIP_CHECK(hipGraphAddMemcpyNode1D(&cpyToHostNode, graph, &kernelCNode, 1, h_array.data(), allocArrayAParams.dptr, arraySize * sizeof(double), hipMemcpyDeviceToHost)); + + // Free array of allocNodeA. It needs to wait for the copy to finish, as kernelC stores its results in it. + hipGraphNode_t freeNodeA; + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &cpyToHostNode, 1, allocArrayAParams.dptr)); + // Free array of allocNodeB. It only needs to wait for kernelC to finish, as it is not written back to the host. + hipGraphNode_t freeNodeB; + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeB, graph, &kernelCNode, 1, allocArrayBParams.dptr)); + + // Instantiate the graph in order to execute it + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + // The graph can be freed after the instantiation if it's not needed for other purposes + HIP_CHECK(hipGraphDestroy(graph)); + + // Actually launch the graph + hipStream_t graphStream; + HIP_CHECK(hipStreamCreate(&graphStream)); + HIP_CHECK(hipGraphLaunch(graphExec, graphStream)); + + // Verify results + constexpr double expected = initValue * 2.0 + 3; + bool passed = true; + for(size_t i = 0; i < arraySize; ++i){ + if(h_array[i] != expected){ + passed = false; + std::cerr << "Validation failed! Expected " << expected << " got " << h_array[0] << std::endl; + break; + } + } + if(passed){ + std::cerr << "Validation passed." << std::endl; + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(graphStream)); + } diff --git a/docs/how-to/programming_manual.md b/docs/how-to/programming_manual.md index 33ab58de93..bac20c9996 100644 --- a/docs/how-to/programming_manual.md +++ b/docs/how-to/programming_manual.md @@ -146,7 +146,7 @@ For Linux developers, the link [here](https://github.com/ROCm/hip-tests/blob/dev ## HIP Graph -HIP graph is supported. For more details, refer to the HIP API Guide. +HIP graphs are supported. For more details, refer to the [HIP API Guide](../doxygen/html/group___graph) or the [how-to section for HIP graphs](../how-to/hipgraph). ## Device-Side Malloc diff --git a/docs/index.md b/docs/index.md index b256580a87..54fa9cc411 100644 --- a/docs/index.md +++ b/docs/index.md @@ -49,6 +49,7 @@ On non-AMD platforms, like NVIDIA, HIP provides header files required to support * [Virtual memory](./how-to/virtual_memory) * {doc}`./how-to/stream_ordered_allocator` * [Cooperative groups](./how-to/cooperative_groups) +* [HIP graphs](./how-to/hipgraph) * {doc}`./how-to/faq` ::: diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index 398b0299de..56abaee03f 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -36,6 +36,8 @@ subtrees: - file: how-to/virtual_memory title: Virtual memory - file: how-to/stream_ordered_allocator + - file: how-to/hipgraph + title: HIP graphs - file: how-to/faq - caption: Reference diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 708ab1d8c2..f8855da989 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1399,8 +1399,8 @@ typedef struct hipMemAllocNodeParams { hipMemPoolProps poolProps; ///< Pool properties, which contain where ///< the location should reside const hipMemAccessDesc* accessDescs;///< The number of memory access descriptors. + size_t accessDescCount; ///< The number of access descriptors. ///< Must not be bigger than the number of GPUs - size_t accessDescCount; ///< The number of access descriptors size_t bytesize; ///< The size of the requested allocation in bytes void* dptr; ///< Returned device address of the allocation } hipMemAllocNodeParams; @@ -1517,9 +1517,9 @@ typedef enum hipGraphInstantiateFlags { hipGraphInstantiateFlagAutoFreeOnLaunch = 1, ///< Automatically free memory allocated in a graph before relaunching. hipGraphInstantiateFlagUpload = - 2, ///< Automatically upload the graph after instantiaton. + 2, ///< Automatically upload the graph after instantiation. hipGraphInstantiateFlagDeviceLaunch = - 4, ///< Instantiate the graph to be launchable from the device. + 4, ///< Instantiate the graph to be launched from the device. hipGraphInstantiateFlagUseNodePriority = 8, ///< Run the graph using the per-node priority attributes rather than the priority of the stream it is launched into. } hipGraphInstantiateFlags; @@ -7191,7 +7191,7 @@ are not safe. * * @returns #hipSuccess, #hipErrorInvalidValue * -* @warning param "const hipGraphEdgeData* dependencyData" is currently not supported and has to +* @warning param "const hipGraphEdgeData* dependencyData" is currently not supported and has to be passed as nullptr. This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues. * @@ -7205,7 +7205,7 @@ hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph, * @brief Ends capture on a stream, returning the captured graph. * * @param [in] stream - Stream to end capture. - * @param [out] pGraph - returns the graph captured. + * @param [out] pGraph - Captured graph. * * @returns #hipSuccess, #hipErrorInvalidValue * @@ -7218,9 +7218,9 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph); /** * @brief Get capture status of a stream. * - * @param [in] stream - Stream under capture. - * @param [out] pCaptureStatus - returns current status of the capture. - * @param [out] pId - unique ID of the capture. + * @param [in] stream - Stream of which to get capture status from. + * @param [out] pCaptureStatus - Returns current capture status. + * @param [out] pId - Unique capture ID. * * @returns #hipSuccess, #hipErrorStreamCaptureImplicit * @@ -7234,12 +7234,12 @@ hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* p /** * @brief Get stream's capture state * - * @param [in] stream - Stream under capture. - * @param [out] captureStatus_out - returns current status of the capture. - * @param [out] id_out - unique ID of the capture. - * @param [in] graph_out - returns the graph being captured into. - * @param [out] dependencies_out - returns pointer to an array of nodes. - * @param [out] numDependencies_out - returns size of the array returned in dependencies_out. + * @param [in] stream - Stream of which to get capture status from. + * @param [out] captureStatus_out - Returns current capture status. + * @param [out] id_out - Unique capture ID. + * @param [out] graph_out - Returns the graph being captured into. + * @param [out] dependencies_out - Pointer to an array of nodes representing the graphs dependencies. + * @param [out] numDependencies_out - Returns size of the array returned in dependencies_out. * * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorStreamCaptureImplicit * @@ -7256,8 +7256,8 @@ hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus /** * @brief Get stream's capture state * - * @param [in] stream - Stream under capture. - * @param [out] pCaptureStatus - returns current status of the capture. + * @param [in] stream - Stream of which to get capture status from. + * @param [out] pCaptureStatus - Returns current capture status. * * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorStreamCaptureImplicit * @@ -7270,11 +7270,11 @@ hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCap /** * @brief Update the set of dependencies in a capturing stream * - * @param [in] stream Stream under capture. - * @param [in] dependencies pointer to an array of nodes to Add/Replace. - * @param [in] numDependencies size of the array in dependencies. - * @param [in] flags Flag how to update dependency set. Should be one of value in enum - * #hipStreamUpdateCaptureDependenciesFlags + * @param [in] stream Stream that is being captured. + * @param [in] dependencies Pointer to an array of nodes to add/replace. + * @param [in] numDependencies Size of the dependencies array. + * @param [in] flags Flag to update dependency set. Should be one of the values + * in enum #hipStreamUpdateCaptureDependenciesFlags. * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorIllegalState * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7288,7 +7288,7 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t /** * @brief Swaps the stream capture mode of a thread. * - * @param [in] mode - Pointer to mode value to swap with the current mode + * @param [in] mode - Pointer to mode value to swap with the current mode. * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7327,10 +7327,10 @@ hipError_t hipGraphDestroy(hipGraph_t graph); /** * @brief Adds dependency edges to a graph. * - * @param [in] graph - instance of the graph to add dependencies. - * @param [in] from - pointer to the graph nodes with dependenties to add from. - * @param [in] to - pointer to the graph nodes to add dependenties to. - * @param [in] numDependencies - the number of dependencies to add. + * @param [in] graph - Instance of the graph to add dependencies to. + * @param [in] from - Pointer to the graph nodes with dependencies to add from. + * @param [in] to - Pointer to the graph nodes to add dependencies to. + * @param [in] numDependencies - Number of dependencies to add. * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7343,10 +7343,10 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, /** * @brief Removes dependency edges from a graph. * - * @param [in] graph - instance of the graph to remove dependencies. + * @param [in] graph - Instance of the graph to remove dependencies from. * @param [in] from - Array of nodes that provide the dependencies. * @param [in] to - Array of dependent nodes. - * @param [in] numDependencies - the number of dependencies to remove. + * @param [in] numDependencies - Number of dependencies to remove. * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7359,16 +7359,16 @@ hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* fr /** * @brief Returns a graph's dependency edges. * - * @param [in] graph - instance of the graph to get the edges from. - * @param [out] from - pointer to the graph nodes to return edge endpoints. - * @param [out] to - pointer to the graph nodes to return edge endpoints. - * @param [out] numEdges - returns number of edges. + * @param [in] graph - Instance of the graph to get the edges from. + * @param [out] from - Pointer to the graph nodes to return edge endpoints. + * @param [out] to - Pointer to the graph nodes to return edge endpoints. + * @param [out] numEdges - Returns number of edges. * @returns #hipSuccess, #hipErrorInvalidValue * * from and to may both be NULL, in which case this function only returns the number of edges in * numEdges. Otherwise, numEdges entries will be filled in. If numEdges is higher than the actual * number of edges, the remaining entries in from and to will be set to NULL, and the number of - * edges actually returned will be written to numEdges + * edges actually returned will be written to numEdges. * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. * @@ -7377,11 +7377,11 @@ hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode size_t* numEdges); /** - * @brief Returns graph nodes. + * @brief Returns a graph's nodes. * - * @param [in] graph - instance of graph to get the nodes. - * @param [out] nodes - pointer to return the graph nodes. - * @param [out] numNodes - returns number of graph nodes. + * @param [in] graph - Instance of graph to get the nodes from. + * @param [out] nodes - Pointer to return the graph nodes. + * @param [out] numNodes - Returns the number of graph nodes. * @returns #hipSuccess, #hipErrorInvalidValue * * nodes may be NULL, in which case this function will return the number of nodes in numNodes. @@ -7395,11 +7395,11 @@ hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes); /** - * @brief Returns graph's root nodes. + * @brief Returns a graph's root nodes. * - * @param [in] graph - instance of the graph to get the nodes. - * @param [out] pRootNodes - pointer to return the graph's root nodes. - * @param [out] pNumRootNodes - returns the number of graph's root nodes. + * @param [in] graph - Instance of the graph to get the nodes from. + * @param [out] pRootNodes - Pointer to return the graph's root nodes. + * @param [out] pNumRootNodes - Returns the number of graph's root nodes. * @returns #hipSuccess, #hipErrorInvalidValue * * pRootNodes may be NULL, in which case this function will return the number of root nodes in @@ -7416,9 +7416,9 @@ hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, /** * @brief Returns a node's dependencies. * - * @param [in] node - graph node to get the dependencies from. - * @param [out] pDependencies - pointer to to return the dependencies. - * @param [out] pNumDependencies - returns the number of graph node dependencies. + * @param [in] node - Graph node to get the dependencies from. + * @param [out] pDependencies - Pointer to return the dependencies. + * @param [out] pNumDependencies - Returns the number of graph node dependencies. * @returns #hipSuccess, #hipErrorInvalidValue * * pDependencies may be NULL, in which case this function will return the number of dependencies in @@ -7435,12 +7435,12 @@ hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDep /** * @brief Returns a node's dependent nodes. * - * @param [in] node - graph node to get the Dependent nodes from. - * @param [out] pDependentNodes - pointer to return the graph dependent nodes. - * @param [out] pNumDependentNodes - returns the number of graph node dependent nodes. + * @param [in] node - Graph node to get the dependent nodes from. + * @param [out] pDependentNodes - Pointer to return the graph dependent nodes. + * @param [out] pNumDependentNodes - Returns the number of graph node dependent nodes. * @returns #hipSuccess, #hipErrorInvalidValue * - * DependentNodes may be NULL, in which case this function will return the number of dependent nodes + * pDependentNodes may be NULL, in which case this function will return the number of dependent nodes * in pNumDependentNodes. Otherwise, pNumDependentNodes entries will be filled in. If * pNumDependentNodes is higher than the actual number of dependent nodes, the remaining entries in * pDependentNodes will be set to NULL, and the number of nodes actually obtained will be returned @@ -7455,8 +7455,8 @@ hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pD /** * @brief Returns a node's type. * - * @param [in] node - instance of the graph to add dependencies. - * @param [out] pType - pointer to the return the type + * @param [in] node - Node to get type of. + * @param [out] pType - Returns the node's type. * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7508,12 +7508,12 @@ hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t origina /** * @brief Creates an executable graph from a graph * - * @param [out] pGraphExec - pointer to instantiated executable graph that is created. - * @param [in] graph - instance of graph to instantiate. - * @param [out] pErrorNode - pointer to error node in case error occured in graph instantiation, - * it could modify the correponding node. - * @param [out] pLogBuffer - pointer to log buffer. - * @param [out] bufferSize - the size of log buffer. + * @param [out] pGraphExec - Pointer to instantiated executable graph. + * @param [in] graph - Instance of graph to instantiate. + * @param [out] pErrorNode - Pointer to error node. In case an error occured during + * graph instantiation, it could modify the corresponding node. + * @param [out] pLogBuffer - Pointer to log buffer. + * @param [out] bufferSize - Size of the log buffer. * * @returns #hipSuccess, #hipErrorOutOfMemory * @@ -7527,8 +7527,8 @@ hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, /** * @brief Creates an executable graph from a graph. * - * @param [out] pGraphExec - pointer to instantiated executable graph that is created. - * @param [in] graph - instance of graph to instantiate. + * @param [out] pGraphExec - Pointer to instantiated executable graph. + * @param [in] graph - Instance of graph to instantiate. * @param [in] flags - Flags to control instantiation. * @returns #hipSuccess, #hipErrorInvalidValue * @@ -7542,9 +7542,9 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g /** * @brief Creates an executable graph from a graph. * - * @param [out] pGraphExec - pointer to instantiated executable graph that is created. - * @param [in] graph - instance of graph to instantiate. - * @param [in] instantiateParams - Graph Instantiate Params + * @param [out] pGraphExec - Pointer to instantiated executable graph. + * @param [in] graph - Instance of graph to instantiate. + * @param [in] instantiateParams - Graph instantiation Params * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7553,10 +7553,10 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g hipError_t hipGraphInstantiateWithParams(hipGraphExec_t* pGraphExec, hipGraph_t graph, hipGraphInstantiateParams *instantiateParams); /** - * @brief launches an executable graph in a stream + * @brief Launches an executable graph in the specified stream. * - * @param [in] graphExec - instance of executable graph to launch. - * @param [in] stream - instance of stream in which to launch executable graph. + * @param [in] graphExec - Instance of executable graph to launch. + * @param [in] stream - Instance of stream in which to launch executable graph. * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7565,10 +7565,10 @@ hipError_t hipGraphInstantiateWithParams(hipGraphExec_t* pGraphExec, hipGraph_t hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream); /** - * @brief uploads an executable graph in a stream + * @brief Uploads an executable graph to a stream * - * @param [in] graphExec - instance of executable graph to launch. - * @param [in] stream - instance of stream in which to launch executable graph. + * @param [in] graphExec - Instance of executable graph to be uploaded. + * @param [in] stream - Instance of stream to which the executable graph is uploaded to. * @returns #hipSuccess, #hipErrorInvalidValue * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7579,11 +7579,11 @@ hipError_t hipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream); /** * @brief Creates a kernel execution node and adds it to a graph. * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in] graph - instance of graph to add the created node. - * @param [in] pDependencies - pointer to the dependencies on the kernel execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] nodeParams - pointer to the parameters for the node. + * @param [out] pGraphNode - Pointer to kernel graph node that is created. + * @param [in] graph - Instance of graph to add the created node to. + * @param [in] pDependencies - Pointer to the dependencies on the kernel execution node. + * @param [in] numDependencies - Number of dependencies. + * @param [in] nodeParams - Pointer to the node parameters. * @returns #hipSuccess, #hipErrorInvalidValue. * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7593,9 +7593,9 @@ hipError_t hipGraphAddNode(hipGraphNode_t *pGraphNode, hipGraph_t graph, hipGraphNodeParams *nodeParams); /** - * @brief Return the flags on executable graph. + * @brief Return the flags of an executable graph. * - * @param [in] graphExec - Executable graph to get the flags. + * @param [in] graphExec - Executable graph to get the flags from. * @param [out] flags - Flags used to instantiate this executable graph. * @returns #hipSuccess, #hipErrorInvalidValue. * @warning This API is marked as Beta. While this feature is complete, it can @@ -7604,10 +7604,10 @@ hipError_t hipGraphAddNode(hipGraphNode_t *pGraphNode, hipGraph_t graph, hipError_t hipGraphExecGetFlags(hipGraphExec_t graphExec, unsigned long long* flags); /** - * @brief Updates parameters of a created node. + * @brief Updates parameters of a graph's node. * - * @param [in] node - instance of the node to set parameters to. - * @param [in] nodeParams - pointer to the parameters. + * @param [in] node - Instance of the node to set parameters for. + * @param [in] nodeParams - Pointer to the parameters to be set. * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDeviceFunction, #hipErrorNotSupported. * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7615,11 +7615,11 @@ hipError_t hipGraphExecGetFlags(hipGraphExec_t graphExec, unsigned long long* fl hipError_t hipGraphNodeSetParams(hipGraphNode_t node, hipGraphNodeParams *nodeParams); /** - * @brief Updates parameters of a created node on executable graph. + * @brief Updates parameters of an executable graph's node. * - * @param [in] graphExec - instance of executable graph. - * @param [in] node - instance of the node to set parameters to. - * @param [in] nodeParams - pointer to the parameters. + * @param [in] graphExec - Instance of the executable graph. + * @param [in] node - Instance of the node to set parameters to. + * @param [in] nodeParams - Pointer to the parameters to be set. * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDeviceFunction, #hipErrorNotSupported. * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7629,7 +7629,7 @@ hipError_t hipGraphExecNodeSetParams(hipGraphExec_t graphExec, hipGraphNode_t no /** * @brief Destroys an executable graph * - * @param [in] graphExec - instance of executable graph to destry. + * @param [in] graphExec - Instance of executable graph to destroy. * * @returns #hipSuccess. * @@ -7646,7 +7646,7 @@ hipError_t hipGraphExecDestroy(hipGraphExec_t graphExec); * @param [in] hGraphExec - instance of executable graph to update. * @param [in] hGraph - graph that contains the updated parameters. * @param [in] hErrorNode_out - node which caused the permissibility check to forbid the update. - * @param [in] updateResult_out - Whether the graph update was permitted. + * @param [in] updateResult_out - Return code whether the graph update was performed. * @returns #hipSuccess, #hipErrorGraphExecUpdateFailure * * @warning This API is marked as Beta. While this feature is complete, it can @@ -7659,11 +7659,11 @@ hipError_t hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, /** * @brief Creates a kernel execution node and adds it to a graph. * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in] graph - instance of graph to add the created node. - * @param [in] pDependencies - pointer to the dependencies on the kernel execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] pNodeParams - pointer to the parameters to the kernel execution node on the GPU. + * @param [out] pGraphNode - Pointer to graph node that is created + * @param [in] graph - Instance of graph to add the created node to. + * @param [in] pDependencies - Pointer to the dependencies of the kernel execution node. + * @param [in] numDependencies - The number of the dependencies. + * @param [in] pNodeParams - Pointer to the parameters of the kernel execution node. * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDeviceFunction * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7686,7 +7686,7 @@ hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* /** * @brief Sets a kernel node's parameters. * - * @param [in] node - instance of the node to set parameters to. + * @param [in] node - Instance of the node to set parameters of. * @param [in] pNodeParams - const pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7697,8 +7697,8 @@ hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeP /** * @brief Sets the parameters for a kernel node in the given graphExec. * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. * @param [in] pNodeParams - const pointer to the kernel node parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7710,12 +7710,12 @@ hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo /** * @brief Creates a memcpy node and adds it to a graph. * - * @param [out] phGraphNode - pointer to graph node to create. - * @param [in] hGraph - instance of graph to add the created node. - * @param [in] dependencies - const pointer to the dependencies on the memcpy execution node. - * @param [in] numDependencies - the number of the dependencies. + * @param [out] phGraphNode - Pointer to graph node that is created. + * @param [in] hGraph - Instance of graph to add the created node to. + * @param [in] dependencies - const pointer to the dependencies of the memcpy execution node. + * @param [in] numDependencies - The number of dependencies. * @param [in] copyParams - const pointer to the parameters for the memory copy. - * @param [in] ctx - cotext related to current device. + * @param [in] ctx - context related to current device. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7727,10 +7727,10 @@ hipError_t hipDrvGraphAddMemcpyNode(hipGraphNode_t* phGraphNode, hipGraph_t hGra /** * @brief Creates a memcpy node and adds it to a graph. * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in] graph - instance of graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. - * @param [in] numDependencies - the number of the dependencies. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of graph to add the created node to. + * @param [in] pDependencies - const pointer to the dependencies of the memcpy execution node. + * @param [in] numDependencies - The number of dependencies. * @param [in] pCopyParams - const pointer to the parameters for the memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7762,10 +7762,10 @@ hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pN hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams); /** - * @brief Sets a node attribute. + * @brief Sets a node's attribute. * - * @param [in] hNode - instance of the node to set parameters to. - * @param [in] attr - the attribute node is set to. + * @param [in] hNode - Instance of the node to set parameters of. + * @param [in] attr - The attribute type to be set. * @param [in] value - const pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7774,10 +7774,10 @@ hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DPar hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr, const hipKernelNodeAttrValue* value); /** - * @brief Gets a node attribute. + * @brief Gets a node's attribute. * - * @param [in] hNode - instance of the node to set parameters to. - * @param [in] attr - the attribute node is set to. + * @param [in] hNode - Instance of the node to set parameters of. + * @param [in] attr - The attribute type to be set. * @param [in] value - const pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7786,10 +7786,10 @@ hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr, hipKernelNodeAttrValue* value); /** - * @brief Sets the parameters for a memcpy node in the given graphExec. + * @brief Sets the parameters of a memcpy node in the given graphExec. * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. * @param [in] pNodeParams - const pointer to the kernel node parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7801,14 +7801,14 @@ hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo /** * @brief Creates a 1D memcpy node and adds it to a graph. * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in] graph - instance of graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] dst - pointer to memory address to the destination. - * @param [in] src - pointer to memory address to the source. - * @param [in] count - the size of the memory to copy. - * @param [in] kind - the type of memory copy. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of graph to add the created node to. + * @param [in] pDependencies - const pointer to the dependencies of the memcpy execution node. + * @param [in] numDependencies - The number of dependencies. + * @param [in] dst - Pointer to memory address of the destination. + * @param [in] src - Pointer to memory address of the source. + * @param [in] count - Size of the memory to copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7820,11 +7820,11 @@ hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, /** * @brief Sets a memcpy node's parameters to perform a 1-dimensional copy. * - * @param [in] node - instance of the node to set parameters to. - * @param [in] dst - pointer to memory address to the destination. - * @param [in] src - pointer to memory address to the source. - * @param [in] count - the size of the memory to copy. - * @param [in] kind - the type of memory copy. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] dst - Pointer to memory address of the destination. + * @param [in] src - Pointer to memory address of the source. + * @param [in] count - Size of the memory to copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7836,12 +7836,12 @@ hipError_t hipGraphMemcpyNodeSetParams1D(hipGraphNode_t node, void* dst, const v * @brief Sets the parameters for a memcpy node in the given graphExec to perform a 1-dimensional * copy. * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. - * @param [in] dst - pointer to memory address to the destination. - * @param [in] src - pointer to memory address to the source. - * @param [in] count - the size of the memory to copy. - * @param [in] kind - the type of memory copy. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] dst - Pointer to memory address of the destination. + * @param [in] src - Pointer to memory address of the source. + * @param [in] count - Size of the memory to copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7853,15 +7853,15 @@ hipError_t hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraph /** * @brief Creates a memcpy node to copy from a symbol on the device and adds it to a graph. * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in] graph - instance of graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] dst - pointer to memory address to the destination. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of graph to add the created node to. + * @param [in] pDependencies - const pointer to the dependencies of the memcpy execution node. + * @param [in] numDependencies - Number of the dependencies. + * @param [in] dst - Pointer to memory address of the destination. * @param [in] symbol - Device symbol address. - * @param [in] count - the size of the memory to copy. + * @param [in] count - Size of the memory to copy. * @param [in] offset - Offset from start of symbol in bytes. - * @param [in] kind - the type of memory copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7874,12 +7874,12 @@ hipError_t hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_ /** * @brief Sets a memcpy node's parameters to copy from a symbol on the device. * - * @param [in] node - instance of the node to set parameters to. - * @param [in] dst - pointer to memory address to the destination. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] dst - Pointer to memory address of the destination. * @param [in] symbol - Device symbol address. - * @param [in] count - the size of the memory to copy. + * @param [in] count - Size of the memory to copy. * @param [in] offset - Offset from start of symbol in bytes. - * @param [in] kind - the type of memory copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7891,13 +7891,13 @@ hipError_t hipGraphMemcpyNodeSetParamsFromSymbol(hipGraphNode_t node, void* dst, * @brief Sets the parameters for a memcpy node in the given graphExec to copy from a symbol on the * * device. * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. - * @param [in] dst - pointer to memory address to the destination. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] dst - Pointer to memory address of the destination. * @param [in] symbol - Device symbol address. - * @param [in] count - the size of the memory to copy. + * @param [in] count - Size of the memory to copy. * @param [in] offset - Offset from start of symbol in bytes. - * @param [in] kind - the type of memory copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7909,15 +7909,15 @@ hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, /** * @brief Creates a memcpy node to copy to a symbol on the device and adds it to a graph. * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in] graph - instance of graph to add the created node. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of graph to add the created node to. * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. - * @param [in] numDependencies - the number of the dependencies. + * @param [in] numDependencies - Number of dependencies. * @param [in] symbol - Device symbol address. - * @param [in] src - pointer to memory address of the src. - * @param [in] count - the size of the memory to copy. + * @param [in] src - Pointer to memory address of the src. + * @param [in] count - Size of the memory to copy. * @param [in] offset - Offset from start of symbol in bytes. - * @param [in] kind - the type of memory copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7931,12 +7931,12 @@ hipError_t hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t /** * @brief Sets a memcpy node's parameters to copy to a symbol on the device. * - * @param [in] node - instance of the node to set parameters to. + * @param [in] node - Instance of the node to set parameters of. * @param [in] symbol - Device symbol address. - * @param [in] src - pointer to memory address of the src. - * @param [in] count - the size of the memory to copy. + * @param [in] src - Pointer to memory address of the src. + * @param [in] count - Size of the memory to copy. * @param [in] offset - Offset from start of symbol in bytes. - * @param [in] kind - the type of memory copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7949,13 +7949,13 @@ hipError_t hipGraphMemcpyNodeSetParamsToSymbol(hipGraphNode_t node, const void* /** * @brief Sets the parameters for a memcpy node in the given graphExec to copy to a symbol on the * device. - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. * @param [in] symbol - Device symbol address. - * @param [in] src - pointer to memory address of the src. - * @param [in] count - the size of the memory to copy. + * @param [in] src - Pointer to memory address of the src. + * @param [in] count - Size of the memory to copy. * @param [in] offset - Offset from start of symbol in bytes. - * @param [in] kind - the type of memory copy. + * @param [in] kind - Type of memory copy. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7967,10 +7967,10 @@ hipError_t hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hi /** * @brief Creates a memset node and adds it to a graph. * - * @param [out] pGraphNode - pointer to the graph node to create. - * @param [in] graph - instance of the graph to add the created node. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of the graph to add the created node to. * @param [in] pDependencies - const pointer to the dependencies on the memset execution node. - * @param [in] numDependencies - the number of the dependencies. + * @param [in] numDependencies - Number of dependencies. * @param [in] pMemsetParams - const pointer to the parameters for the memory set. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -7983,8 +7983,8 @@ hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, /** * @brief Gets a memset node's parameters. * - * @param [in] node - instane of the node to get parameters from. - * @param [out] pNodeParams - pointer to the parameters. + * @param [in] node - Instance of the node to get parameters of. + * @param [out] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -7994,8 +7994,8 @@ hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNo /** * @brief Sets a memset node's parameters. * - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - pointer to the parameters. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8005,9 +8005,9 @@ hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParam /** * @brief Sets the parameters for a memset node in the given graphExec. * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - pointer to the parameters. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8018,11 +8018,11 @@ hipError_t hipGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo /** * @brief Creates a host execution node and adds it to a graph. * - * @param [out] pGraphNode - pointer to the graph node to create. - * @param [in] graph - instance of the graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the memset execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] pNodeParams -pointer to the parameters. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of the graph to add the created node to. + * @param [in] pDependencies - const pointer to the dependencies of the memset execution node. + * @param [in] numDependencies - Number of dependencies. + * @param [in] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8034,8 +8034,8 @@ hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, /** * @brief Returns a host node's parameters. * - * @param [in] node - instane of the node to get parameters from. - * @param [out] pNodeParams - pointer to the parameters. + * @param [in] node - Instance of the node to get parameters of. + * @param [out] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8045,8 +8045,8 @@ hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNo /** * @brief Sets a host node's parameters. * - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - pointer to the parameters. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8056,9 +8056,9 @@ hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParam /** * @brief Sets the parameters for a host node in the given graphExec. * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - pointer to the parameters. + * @param [in] hGraphExec - Instance of the executable graph with the node. + * @param [in] node - Instance of the node to set parameters of. + * @param [in] pNodeParams - Pointer to the parameters. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8069,11 +8069,11 @@ hipError_t hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode /** * @brief Creates a child graph node and adds it to a graph. * - * @param [out] pGraphNode - pointer to the graph node to create. - * @param [in] graph - instance of the graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the memset execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] childGraph - the graph to clone into this node + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of the graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies of the memset execution node. + * @param [in] numDependencies - Number of dependencies. + * @param [in] childGraph - Graph to clone into this node * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8085,8 +8085,8 @@ hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t grap /** * @brief Gets a handle to the embedded graph of a child graph node. * - * @param [in] node - instane of the node to get child graph. - * @param [out] pGraph - pointer to get the graph. + * @param [in] node - Instance of the node to get child graph of. + * @param [out] pGraph - Pointer to get the graph. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8109,10 +8109,10 @@ hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGra /** * @brief Creates an empty node and adds it to a graph. * - * @param [out] pGraphNode - pointer to the graph node to create and add to the graph. - * @param [in] graph - instane of the graph the node is add to. - * @param [in] pDependencies - const pointer to the node dependenties. - * @param [in] numDependencies - the number of dependencies. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of the graph the node is added to. + * @param [in] pDependencies - const pointer to the node dependencies. + * @param [in] numDependencies - Number of dependencies. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8124,11 +8124,11 @@ hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, /** * @brief Creates an event record node and adds it to a graph. * - * @param [out] pGraphNode - pointer to the graph node to create and add to the graph. - * @param [in] graph - instane of the graph the node to be added. - * @param [in] pDependencies - const pointer to the node dependenties. - * @param [in] numDependencies - the number of dependencies. - * @param [in] event - Event for the node. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of the graph the node is added to. + * @param [in] pDependencies - const pointer to the node dependencies. + * @param [in] numDependencies - Number of dependencies. + * @param [in] event - Event of the node. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8140,7 +8140,7 @@ hipError_t hipGraphAddEventRecordNode(hipGraphNode_t* pGraphNode, hipGraph_t gra /** * @brief Returns the event associated with an event record node. * - * @param [in] node - instane of the node to get event from. + * @param [in] node - Instance of the node to get event of. * @param [out] event_out - Pointer to return the event. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -8151,8 +8151,8 @@ hipError_t hipGraphEventRecordNodeGetEvent(hipGraphNode_t node, hipEvent_t* even /** * @brief Sets an event record node's event. * - * @param [in] node - instane of the node to set event to. - * @param [in] event - pointer to the event. + * @param [in] node - Instance of the node to set event to. + * @param [in] event - Pointer to the event. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8175,10 +8175,10 @@ hipError_t hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, hipGra /** * @brief Creates an event wait node and adds it to a graph. * - * @param [out] pGraphNode - pointer to the graph node to create and add to the graph. - * @param [in] graph - instane of the graph the node to be added. - * @param [in] pDependencies - const pointer to the node dependenties. - * @param [in] numDependencies - the number of dependencies. + * @param [out] pGraphNode - Pointer to graph node that is created. + * @param [in] graph - Instance of the graph the node to be added. + * @param [in] pDependencies - const pointer to the node dependencies. + * @param [in] numDependencies - Number of dependencies. * @param [in] event - Event for the node. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -8192,7 +8192,7 @@ hipError_t hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph /** * @brief Returns the event associated with an event wait node. * - * @param [in] node - instane of the node to get event from. + * @param [in] node - Instance of the node to get event of. * @param [out] event_out - Pointer to return the event. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -8203,8 +8203,8 @@ hipError_t hipGraphEventWaitNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_ /** * @brief Sets an event wait node's event. * - * @param [in] node - instane of the node to set event to. - * @param [in] event - pointer to the event. + * @param [in] node - Instance of the node to set event of. + * @param [in] event - Pointer to the event. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8228,10 +8228,10 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph * @brief Creates a memory allocation node and adds it to a graph * * @param [out] pGraphNode - Pointer to the graph node to create and add to the graph - * @param [in] graph - Instane of the graph the node to be added - * @param [in] pDependencies - Const pointer to the node dependenties + * @param [in] graph - Instance of the graph node to be added + * @param [in] pDependencies - Const pointer to the node dependencies * @param [in] numDependencies - The number of dependencies - * @param [in] pNodeParams - Node parameters for memory allocation + * @param [in, out] pNodeParams - Node parameters for memory allocation, returns a pointer to the allocated memory. * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8242,7 +8242,7 @@ hipError_t hipGraphAddMemAllocNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, /** * @brief Returns parameters for memory allocation node * - * @param [in] node - Memory allocation node for a query + * @param [in] node - Memory allocation node to query * @param [out] pNodeParams - Parameters for the specified memory allocation node * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can @@ -8254,8 +8254,8 @@ hipError_t hipGraphMemAllocNodeGetParams(hipGraphNode_t node, hipMemAllocNodePar * @brief Creates a memory free node and adds it to a graph * * @param [out] pGraphNode - Pointer to the graph node to create and add to the graph - * @param [in] graph - Instane of the graph the node to be added - * @param [in] pDependencies - Const pointer to the node dependenties + * @param [in] graph - Instance of the graph node to be added + * @param [in] pDependencies - Const pointer to the node dependencies * @param [in] numDependencies - The number of dependencies * @param [in] dev_ptr - Pointer to the memory to be freed * @returns #hipSuccess, #hipErrorInvalidValue @@ -8268,8 +8268,8 @@ hipError_t hipGraphAddMemFreeNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, /** * @brief Returns parameters for memory free node * - * @param [in] node - Memory free node for a query - * @param [out] dev_ptr - Device pointer for the specified memory free node + * @param [in] node - Memory free node to query + * @param [out] dev_ptr - Device pointer of the specified memory free node * @returns #hipSuccess, #hipErrorInvalidValue * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8279,9 +8279,9 @@ hipError_t hipGraphMemFreeNodeGetParams(hipGraphNode_t node, void* dev_ptr); /** * @brief Get the mem attribute for graphs. * - * @param [in] device - device the attr is get for. - * @param [in] attr - attr to get. - * @param [out] value - value for specific attr. + * @param [in] device - Device to get attributes from + * @param [in] attr - Attribute type to be queried + * @param [out] value - Value of the queried attribute * @returns #hipSuccess, #hipErrorInvalidDevice * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8291,9 +8291,9 @@ hipError_t hipDeviceGetGraphMemAttribute(int device, hipGraphMemAttributeType at /** * @brief Set the mem attribute for graphs. * - * @param [in] device - device the attr is set for. - * @param [in] attr - attr to set. - * @param [in] value - value for specific attr. + * @param [in] device - Device to set attribute of. + * @param [in] attr - Attribute type to be set. + * @param [in] value - Value of the attribute. * @returns #hipSuccess, #hipErrorInvalidDevice * @warning This API is marked as Beta. While this feature is complete, it can * change and might have outstanding issues. @@ -8301,9 +8301,9 @@ hipError_t hipDeviceGetGraphMemAttribute(int device, hipGraphMemAttributeType at hipError_t hipDeviceSetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value); /** - * @brief Free unused memory on specific device used for graph back to OS. + * @brief Free unused memory reserved for graphs on a specific device and return it back to the OS. * - * @param [in] device - device the memory is used for graphs + * @param [in] device - Device for which memory should be trimmed * @returns #hipSuccess, #hipErrorInvalidDevice * * @warning This API is marked as Beta. While this feature is complete, it can @@ -8593,8 +8593,8 @@ hipError_t hipDrvGraphAddMemsetNode(hipGraphNode_t* phGraphNode, hipGraph_t hGra * @brief Creates a memory free node and adds it to a graph * * @param [out] phGraphNode - Pointer to the graph node to create and add to the graph - * @param [in] hGraph - Instane of the graph the node to be added - * @param [in] dependencies - Const pointer to the node dependenties + * @param [in] hGraph - Instance of the graph the node to be added + * @param [in] dependencies - Const pointer to the node dependencies * @param [in] numDependencies - The number of dependencies * @param [in] dptr - Pointer to the memory to be freed * @returns #hipSuccess, #hipErrorInvalidValue