From 6d2e70d8da3f984934a0086ff98d3053ab21a0dc Mon Sep 17 00:00:00 2001
From: "Bhardwaj, Gopesh" <Gopesh.Bhardwaj@amd.com>
Date: Thu, 5 Dec 2024 19:47:57 +0530
Subject: [PATCH] updating roctx documentation for functions (#30)

updating roctx documentation for funcitons
---
 source/docs/how-to/using-rocprofv3.rst | 61 ++++++++++++++++++++++++++
 1 file changed, 61 insertions(+)

diff --git a/source/docs/how-to/using-rocprofv3.rst b/source/docs/how-to/using-rocprofv3.rst
index 09fe980c0b..82bfcb7bec 100644
--- a/source/docs/how-to/using-rocprofv3.rst
+++ b/source/docs/how-to/using-rocprofv3.rst
@@ -276,6 +276,14 @@ Here is a list of useful APIs for code instrumentation.
 - ``roctxRangePush``: Starts a new nested range.
 - ``roctxRangePop``: Stops the current nested range.
 - ``roctxRangeStop``: Stops the given range.
+- ``roctxProfilerPause``: Request any currently running profiling tool that it should stop collecting data.
+- ``roctxProfilerResume``: Request any currently running profiling tool that it should resume collecting data.
+- ``roctxGetThreadId``: Retrieve a id value for the current thread which will be identical to the id value a profiling tool gets via `rocprofiler_get_thread_id(rocprofiler_thread_id_t*)`.
+- ``roctxNameOsThread``: Current CPU OS thread to be labeled by the provided name in the output of the profiling tool.
+- ``roctxNameHsaAgent``: Given HSA agent to be labeled by the provided name in the output of the profiling tool.
+- ``roctxNameHipDevice``: Given HIP device id to be labeled by the provided name in the output of the profiling tool.
+- ``roctxNameHipStream``: Given HIP stream to be labeled by the provided name in the output of the profiling tool.
+
 
 .. note::
   To use ``rocprofv3`` for marker tracing, including and linking to old ROCTx works but it is recommended to switch to new ROCTx because
@@ -328,6 +336,59 @@ Here are the contents of ``marker_api_trace.csv`` file:
 
 For the description of the fields in the output file, see :ref:`output-file-fields`.
 
+``roctxProfilerPause`` and ``roctxProfilerResume`` can be used to hide the calls between them. This is useful when you want to hide the calls that are not relevant to your profiling session.
+
+.. code-block:: bash
+
+    #include <rocprofiler-sdk-roctx/roctx.h>
+
+    // Memory transfer from host to device
+    HIP_API_CALL(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice));
+
+    auto tid = roctx_thread_id_t{};
+    roctxGetThreadId(&tid);
+    roctxProfilerPause(tid);
+    // Memory transfer that should be hidden by profiling tool
+    HIP_API_CALL(
+        hipMemcpy(gpuTransposeMatrix, gpuMatrix, NUM * sizeof(float), hipMemcpyDeviceToDevice));
+    roctxProfilerResume(tid);
+
+    // Lauching kernel from host
+    hipLaunchKernelGGL(matrixTranspose,
+                       dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y),
+                       dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                       0,
+                       0,
+                       gpuTransposeMatrix,
+                       gpuMatrix,
+                       WIDTH);
+
+    // Memory transfer from device to host
+    HIP_API_CALL(
+        hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost));
+
+.. code-block:: shell
+
+    rocprofv3 --marker-trace --hip-trace -- <application_path>
+
+    The above command generates a ``hip_api_trace.csv`` file prefixed with the process ID, which has only 2  `hipMemcpy` calls and the in between ``hipMemcpyDeviceToHost`` is hidden .
+
+.. code-block:: shell
+
+   "Domain","Function","Process_Id","Thread_Id","Correlation_Id","Start_Timestamp","End_Timestamp"
+   "HIP_COMPILER_API","__hipRegisterFatBinary",1643920,1643920,1,320301257609216,320301257636427
+   "HIP_COMPILER_API","__hipRegisterFunction",1643920,1643920,2,320301257650707,320301257678857
+   "HIP_RUNTIME_API","hipGetDevicePropertiesR0600",1643920,1643920,4,320301258114239,320301337764472
+   "HIP_RUNTIME_API","hipMalloc",1643920,1643920,5,320301338073823,320301338247374
+   "HIP_RUNTIME_API","hipMalloc",1643920,1643920,6,320301338248284,320301338399595
+   "HIP_RUNTIME_API","hipMemcpy",1643920,1643920,7,320301338410995,320301631549262
+   "HIP_COMPILER_API","__hipPushCallConfiguration",1643920,1643920,10,320301632131175,320301632134215
+   "HIP_COMPILER_API","__hipPopCallConfiguration",1643920,1643920,11,320301632137745,320301632139735
+   "HIP_RUNTIME_API","hipLaunchKernel",1643920,1643920,12,320301632142615,320301632898289
+   "HIP_RUNTIME_API","hipMemcpy",1643920,1643920,14,320301632901249,320301633934395
+   "HIP_RUNTIME_API","hipFree",1643920,1643920,15,320301643320908,320301643511479
+   "HIP_RUNTIME_API","hipFree",1643920,1643920,16,320301643512629,320301643585639
+
 Kernel Rename
 ++++++++++++++