Handle special cases when stream value is hipStreamLegacy (0x01) or hipStreamPerThread (0x02) (#343)

* Updated stream code to handle special cases when stream value is 0x01 or 0x02

* Removed extra definitions and updated tests to account for special case

* Modified stream.cpp so that each thread assigned a unique stream ID when hipStreamPerThread is used as stream value. Modified tests to check that threads are assigned unique, repeated values when hipStreamPerThread is called

* Updated idx_offset, stream_map, and thread counter to be in one struct.

* Update stream.cpp to only use add_stream() and update tests for seperate unit test for hipStreamPerThread

* Remove unecessary comment

* Removed unecessary line

* Updated tests and stream.cpp to update stream ID correctly

* Updated test structure

projects/rocprofiler-sdk/tests/bin/CMakeLists.txt

@@ -39,3 +39,4 @@ if(rocJPEG_FOUND AND rocJPEG_VERSION VERSION_GREATER 0.6.0)
 endif()
 add_subdirectory(hsa-code-object)
 add_subdirectory(hip-streams)
+add_subdirectory(hip-streams-per-thread)

projects/rocprofiler-sdk/tests/bin/hip-streams-per-thread/CMakeLists.txt

@@ -0,0 +1,42 @@
+#
+#
+#
+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-per-thread 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(hip-streams-per-thread.cpp PROPERTIES LANGUAGE HIP)
+
+add_executable(hip-streams-per-thread)
+target_sources(hip-streams-per-thread PRIVATE hip-streams-per-thread.cpp)
+target_link_libraries(hip-streams-per-thread PRIVATE rocprofiler-sdk::tests-build-flags)
+
+find_package(Threads REQUIRED)
+target_link_libraries(hip-streams-per-thread PRIVATE Threads::Threads)

projects/rocprofiler-sdk/tests/bin/hip-streams-per-thread/hip-streams-per-thread.cpp

@@ -0,0 +1,98 @@
+// MIT License
+//
+// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in
+// all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+// THE SOFTWARE.
+
+#include <array>
+#include <cstdlib>
+#include <thread>
+#include <vector>
+
+#include "hip/hip_runtime.h"
+
+/* 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)
+
+static void
+copy_to_dev(const hipStream_t stream)
+{
+    unsigned int n   = (32 * 1024);  // 32KB
+    double*      A_h = nullptr;
+    double*      A_d = nullptr;
+
+    HIP_ASSERT(hipHostMalloc(&A_h, n * sizeof(double)));
+    HIP_ASSERT(hipMalloc(&A_d, n * sizeof(double)));
+
+    for(unsigned int i = 0; i < n; ++i)
+    {
+        A_h[i] = 123.5;
+    }
+    HIP_ASSERT(hipMemcpyAsync(A_d, A_h, n * sizeof(double), hipMemcpyHostToDevice, stream));
+    // Repeat to make sure streams remain the same
+    HIP_ASSERT(hipMemcpyAsync(A_d, A_h, n * sizeof(double), hipMemcpyHostToDevice, stream));
+
+    // Release device memory
+    HIP_ASSERT(hipFree(A_d));
+    // Release host memory
+    HIP_ASSERT(hipHostFree(A_h));
+}
+
+int
+main(int argc, char** argv)
+{
+    // Test hipStreamPerThread with multiple threads
+    const size_t                         num_streams = 3;
+    const size_t                         thread_cnt  = argc < 2 ? 9 : atoi(argv[1]);
+    std::vector<std::thread>             threads{};
+    std::array<hipStream_t, num_streams> streams{};
+    threads.reserve(thread_cnt);
+    threads.emplace_back(std::thread(copy_to_dev, nullptr));
+    for(size_t i = 1, j = 0; i < thread_cnt; ++i)
+    {
+        if(i % 3 == 0)
+        {
+            threads.emplace_back(std::thread(copy_to_dev, hipStreamLegacy));
+        }
+        else if(i % 3 == 1)
+        {
+            threads.emplace_back(std::thread(copy_to_dev, hipStreamPerThread));
+        }
+        else
+        {
+            HIP_ASSERT(hipStreamCreate(&streams[j]));
+            threads.emplace_back(std::thread(copy_to_dev, streams[j++]));
+        }
+    }
+    for(auto& thread : threads)
+    {
+        thread.join();
+    }
+    return 0;
+}

projects/rocprofiler-sdk/tests/bin/hip-streams/compute_comm_overlap.cpp

@@ -110,6 +110,9 @@ main()
     const int gridSizePerStream = 104;  //(int)ceil((float)elements_per_stream/blockSize);
 
     HIP_ASSERT(hipEventRecord(start));
+    // Extra copy with null stream
+    HIP_ASSERT(hipMemcpyAsync(
+        &d_input1[0], &h_input1[0], bytes_per_stream, hipMemcpyHostToDevice, nullptr));
     // split H2D copies and kernel calls into separate loops
     for(int i = 0; i < num_streams; i++)
     {