Fix HIP Streams Duplication Error (#313)

* Fix stream duplication and fixed tests * Added comments to explain stream.cpp code, change stream nullptr check to occur in update table to prevent readding null stream, simplified hip-streams bin file code, add destroyStreams to hip-streams bin file code * Removed roctx from CMakeLists.txt * Updated documentation * Fix documentation * Removed update_table for HIP compiler table and updated stream.cpp to remove support for HIP compiler table * Added runtime initialization check for HIP * Changed tool name, working on fixing memory management * Added context for counter collection kernel rename combination * Changed name from map to set and changed description * Fix documentation description for group-by-queue * Merged memory copy and kernel operations onto a single track when on the same stream * Updated perfetto output to remove hardware information from track name to merge all memory copy and kernel operations on the same stream to the same track: * Most pr comments addressed * Added filter for counter collection and removed kernel buffer tracing hack * Added PR comment fixes --------- Co-authored-by: Madsen, Jonathan <Jonathan.Madsen@amd.com> [ROCm/rocprofiler-sdk commit: e626df43eb]
2025-05-01 00:56:15 -05:00
parent 49486fee5e
commit 24f054f509
26 changed files with 465 additions and 287 deletions
@@ -38,3 +38,4 @@ if(rocJPEG_FOUND AND rocJPEG_VERSION VERSION_GREATER 0.6.0)
    add_subdirectory(rocjpeg)
 endif()
 add_subdirectory(hsa-code-object)
+add_subdirectory(hip-streams)
@@ -0,0 +1,41 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+if(NOT CMAKE_HIP_COMPILER)
+    find_program(
+        amdclangpp_EXECUTABLE
+        NAMES amdclang++
+        HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATH_SUFFIXES bin llvm/bin NO_CACHE)
+    mark_as_advanced(amdclangpp_EXECUTABLE)
+
+    if(amdclangpp_EXECUTABLE)
+        set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
+    endif()
+endif()
+
+project(rocprofiler-sdk-tests-bin-hip-streams LANGUAGES CXX HIP)
+
+foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
+    if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
+        set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
+    endif()
+endforeach()
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_HIP_STANDARD 17)
+set(CMAKE_HIP_EXTENSIONS OFF)
+set(CMAKE_HIP_STANDARD_REQUIRED ON)
+
+set_source_files_properties(compute_comm_overlap.cpp PROPERTIES LANGUAGE HIP)
+add_executable(hip-streams)
+target_sources(hip-streams PRIVATE compute_comm_overlap.cpp)
+target_link_libraries(hip-streams PRIVATE rocprofiler-sdk::tests-build-flags)
+
+find_package(Threads REQUIRED)
+target_link_libraries(hip-streams PRIVATE Threads::Threads)
@@ -0,0 +1,138 @@
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "hip/hip_runtime.h"
+
+#define BLOCKDIM 64
+
+/* Macro for checking GPU API return values */
+#define HIP_ASSERT(call)                                                                           \
+    do                                                                                             \
+    {                                                                                              \
+        hipError_t gpuErr = call;                                                                  \
+        if(hipSuccess != gpuErr)                                                                   \
+        {                                                                                          \
+            printf(                                                                                \
+                "GPU API Error - %s:%d: '%s'\n", __FILE__, __LINE__, hipGetErrorString(gpuErr));   \
+            exit(1);                                                                               \
+        }                                                                                          \
+    } while(0)
+
+// HIP kernel. Each thread takes care of one element of input
+__global__ void
+cube(double* input, double* output, size_t offset, size_t elements_per_stream)
+{
+    size_t tid     = blockIdx.x * blockDim.x + threadIdx.x;
+    size_t gstride = blockDim.x * gridDim.x;
+
+    // Span all elements assigned to this stream
+    for(size_t id = tid + offset; id < offset + elements_per_stream; id += gstride)
+        for(size_t i = 0; i < 1000; ++i)
+            output[id] = input[id] * input[id] * input[id];
+}
+
+int
+main()
+{
+    // number of streams
+    const int num_streams = 8;
+
+    // Number of threads in each thread block
+    const int blockSize = 512;
+
+    // Size of vectors
+    int n                   = 100000000;
+    int elements_per_stream = n / num_streams;
+    int bytes_per_stream    = elements_per_stream * sizeof(double);
+
+    // Host input vectors
+    double* h_input1{nullptr};
+    // Host output vector
+    double* h_output1{nullptr};
+
+    // Device input vectors
+    double* d_input1{nullptr};
+    // Device output vector
+    double* d_output1{nullptr};
+
+    // Creating events for timers
+    hipEvent_t start{}, stop{};
+    HIP_ASSERT(hipEventCreate(&start));
+    HIP_ASSERT(hipEventCreate(&stop));
+
+    // Creating streams
+    hipStream_t streams[num_streams];
+    for(int i = 0; i < num_streams; ++i)
+    {
+        HIP_ASSERT(hipStreamCreate(&streams[i]));
+    }
+
+    // Size, in bytes, of each vector
+    size_t bytes = n * sizeof(double);
+
+    // Allocate page locked memory for these vectors on host
+    HIP_ASSERT(hipHostMalloc(&h_input1, bytes));
+    HIP_ASSERT(hipHostMalloc(&h_output1, bytes));
+
+    // Allocate memory for each vector on GPU
+    HIP_ASSERT(hipMalloc(&d_input1, bytes));
+    HIP_ASSERT(hipMalloc(&d_output1, bytes));
+
+    // Initialize vectors on host
+    for(int i = 0; i < n; i++)
+    {
+        h_input1[i] = sin(i);
+    }
+
+    // Number of thread blocks in grid
+    const int gridSizePerStream = 104;  //(int)ceil((float)elements_per_stream/blockSize);
+
+    HIP_ASSERT(hipEventRecord(start));
+    // split H2D copies and kernel calls into separate loops
+    for(int i = 0; i < num_streams; i++)
+    {
+        int offset = i * elements_per_stream;
+        HIP_ASSERT(hipMemcpyAsync(&d_input1[offset],
+                                  &h_input1[offset],
+                                  bytes_per_stream,
+                                  hipMemcpyHostToDevice,
+                                  streams[i]));
+    }
+    for(int i = 0; i < num_streams; i++)
+    {
+        int offset = i * elements_per_stream;
+        cube<<<gridSizePerStream, blockSize, 0, streams[i]>>>(
+            d_input1, d_output1, offset, elements_per_stream);
+    }
+    for(int i = 0; i < num_streams; i++)
+    {
+        int offset = i * elements_per_stream;
+        HIP_ASSERT(hipMemcpyAsync(&h_output1[offset],
+                                  &d_output1[offset],
+                                  bytes_per_stream,
+                                  hipMemcpyDeviceToHost,
+                                  streams[i]));
+    }
+
+    HIP_ASSERT(hipEventRecord(stop));
+    HIP_ASSERT(hipEventSynchronize(stop));
+
+    float milliseconds = 0;
+    HIP_ASSERT(hipEventElapsedTime(&milliseconds, start, stop));
+
+    // Release device memory
+    HIP_ASSERT(hipFree(d_input1));
+    HIP_ASSERT(hipFree(d_output1));
+
+    // Release host memory
+    HIP_ASSERT(hipHostFree(h_input1));
+    HIP_ASSERT(hipHostFree(h_output1));
+
+    // Destroy streams
+    for(int i = 0; i < num_streams; ++i)
+    {
+        HIP_ASSERT(hipStreamDestroy(streams[i]));
+    }
+
+    return 0;
+}
@@ -29,7 +29,6 @@ project(
    VERSION 0.0.0)

 find_package(rocprofiler-sdk REQUIRED)
-find_package(rocDecode)

 rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py)

@@ -41,9 +40,9 @@ set(hip-stream-display-env "${PRELOAD_ENV}")
 add_test(
    NAME rocprofv3-test-hip-stream-display-execute
    COMMAND
-        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --kernel-rename -d
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --kernel-rename -s -d
        ${CMAKE_CURRENT_BINARY_DIR}/%tag%-trace -o out --output-format json pftrace
-        --log-level env -- $<TARGET_FILE:transpose>)
+        --log-level env -- $<TARGET_FILE:hip-streams>)

 set_tests_properties(
    rocprofv3-test-hip-stream-display-execute
@@ -56,14 +55,14 @@ set_tests_properties(
               FAIL_REGULAR_EXPRESSION
               "threw an exception"
               DISABLED
-               $<NOT:$<TARGET_EXISTS:transpose>>)
+               $<NOT:$<TARGET_EXISTS:hip-streams>>)

 add_test(
    NAME rocprofv3-test-hip-stream-display-validate
    COMMAND
        ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input
-        ${CMAKE_CURRENT_BINARY_DIR}/hip-stream-display/out_results.json --pftrace-input
-        ${CMAKE_CURRENT_BINARY_DIR}/hip-stream-display/out_results.pftrace)
+        ${CMAKE_CURRENT_BINARY_DIR}/hip-streams-trace/out_results.json --pftrace-input
+        ${CMAKE_CURRENT_BINARY_DIR}/hip-streams-trace/out_results.pftrace)

 set_tests_properties(
    rocprofv3-test-hip-stream-display-validate
@@ -76,4 +75,4 @@ set_tests_properties(
               FAIL_REGULAR_EXPRESSION
               "AssertionError"
               DISABLED
-               $<NOT:$<TARGET_EXISTS:transpose>>)
+               $<NOT:$<TARGET_EXISTS:hip-streams>>)
@@ -54,15 +54,17 @@ def test_stream_trace(json_data):
    buffer_records = data["buffer_records"]

    kernel_dispatch_data = buffer_records["kernel_dispatch"]
-    memory_copies_data = buffer_records["memory_copies"]
+    memory_copies_data = buffer_records["memory_copy"]
    assert len(kernel_dispatch_data) > 0
    assert len(memory_copies_data) > 0

-    # Expect stream ids to be set to 1 or 2 for transpose executable
-    expected_stream_ids = set((1, 2))
+    # Expect stream ids to be set between 1 and 8 inclusive for transpose executable
+    expected_stream_ids = set([i for i in range(1, 9)])
+
    # check buffering data
    for titr in (kernel_dispatch_data, memory_copies_data):
-        for node in rocdecode_data:
+        stream_id_set = set()
+        for node in titr:
            assert "size" in node
            assert "kind" in node
            assert "operation" in node
@@ -70,7 +72,7 @@ def test_stream_trace(json_data):
            assert "end_timestamp" in node
            assert "start_timestamp" in node
            assert "thread_id" in node
-            assert "_stream_id" in node
+            assert "stream_id" in node

            assert node.size > 0
            assert node.thread_id > 0
@@ -78,13 +80,15 @@ def test_stream_trace(json_data):
            assert node.end_timestamp > 0
            assert node.start_timestamp < node.end_timestamp

-            assert node._stream_id.handle in expected_stream_ids
+            stream_id = node.stream_id.handle
+            stream_id_set.add(stream_id)
+        assert stream_id_set == expected_stream_ids


 def test_perfetto_data(pftrace_data, json_data):
    import rocprofiler_sdk.tests.rocprofv3 as rocprofv3

-    assert pftrace_data != None
+    assert pftrace_data.empty == False
    rocprofv3.test_perfetto_data(
        pftrace_data,
        json_data,