SWDEV-291783 - Catch2 tests for some of Stream APIs

Change-Id: I656fc59184863ef142cf89d7541982492f6b8484

catch/hipTestMain/CMakeLists.txt

@@ -11,6 +11,7 @@ endif()
 
 target_link_libraries(UnitTests PRIVATE UnitDeviceTests
                                         MemoryTest
+                                        StreamTest
                                         stdc++fs)
 
 # Add AMD Only Tests

catch/include/hip_test_common.hh

@@ -43,5 +43,9 @@ THE SOFTWARE.
 
 // Utility Functions
 namespace HipTest {
-int getDeviceCount();
+static inline int getDeviceCount() {
+  int dev = 0;
+  HIP_CHECK(hipGetDeviceCount(&dev));
+  return dev;
+}
 }

catch/unit/CMakeLists.txt

@@ -1,4 +1,5 @@
 add_subdirectory(memory)
 add_subdirectory(deviceLib)
+add_subdirectory(stream)
 # Disable Saxpy test temporarily to see if CI Passes
 # add_subdirectory(rtc)

catch/unit/stream/CMakeLists.txt

@@ -0,0 +1,12 @@
+set(TEST_SRC
+    hipStreamCreate.cc
+    hipStreamGetFlags.cc
+    hipStreamGetPriority.cc
+    hipMultiStream.cc
+)
+
+# Create shared lib of all tests
+add_library(StreamTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+
+# Add dependency on build_tests to build it on this custom target
+add_dependencies(build_tests StreamTest)

catch/unit/stream/hipMultiStream.cc

@@ -0,0 +1,97 @@
+
+/*
+Copyright (c) 2021-Present 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 <hip_test_common.hh>
+#include <iostream>
+#include <vector>
+constexpr int NN = 1 << 21;
+__global__ void kernel_do_nothing(__attribute__((unused))int a) {
+  // empty kernel
+}
+__global__ void kernel(float* x, float* y, int n) {
+  size_t tid{threadIdx.x};
+  if (tid < 1) {
+    for (int i = 0; i < n; i++) {
+      x[i] = sqrt(powf(3.14159, i));
+    }
+    y[tid] = y[tid] + 1.0f;
+  }
+}
+__global__ void nKernel(float* y) {
+  size_t tid{threadIdx.x};
+  y[tid] = y[tid] + 1.0f;
+}
+TEST_CASE("Unit_hipMultiStream_sameDevice") {
+  constexpr int num_streams{8};
+  hipStream_t streams[num_streams];
+  float *data[num_streams], *yd, *xd;
+  float y{1.0f}, x{1.0f};
+  HIP_CHECK(hipMalloc((void**)&yd, sizeof(float)));
+  HIP_CHECK(hipMalloc((void**)&xd, sizeof(float)));
+  HIP_CHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice));
+  for (int i = 0; i < num_streams; i++) {
+    HIP_CHECK(hipStreamCreate(&streams[i]));
+    HIP_CHECK(hipMalloc(&data[i], NN * sizeof(float)));
+    hipLaunchKernelGGL(kernel, dim3(1), dim3(1), 0, streams[i], data[i], xd, NN);
+    hipLaunchKernelGGL(HIP_KERNEL_NAME(nKernel), dim3(1), dim3(1), 0, 0, yd);
+  }
+  HIP_CHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
+  REQUIRE(x == Approx(y));
+}
+TEST_CASE("Unit_hipMultiStream_multimeDevice") {
+  constexpr int nLoops = 100000;
+  constexpr int nStreams = 2;
+  std::vector<hipStream_t> streams(nStreams);
+  int nGpu = 0;
+  HIP_CHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+    std::cout << "info: didn't find any GPU! skipping the test!\n";
+    REQUIRE(true);
+  }
+  static int device = 0;
+  HIP_CHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+  std::cout << "info: running on bus "
+            << "0x" << props.pciBusID << " " << props.name << std::endl;
+  for (int i = 0; i < nStreams; i++) {
+    HIP_CHECK(hipStreamCreate(&streams[i]));
+  }
+  for (int k = 0; k <= nLoops; ++k) {
+    HIP_CHECK(hipDeviceSynchronize());
+    // Launch kernel with default stream
+    hipLaunchKernelGGL(kernel_do_nothing, dim3(1), dim3(1), 0, 0, 1);
+    // Launch kernel on all streams
+    for (int i = 0; i < nStreams; i++) {
+      hipLaunchKernelGGL(kernel_do_nothing, dim3(1), dim3(1), 0, streams[i], 1);
+    }
+    // Sync stream 1
+    HIP_CHECK(hipStreamSynchronize(streams[0]));
+    if (k % 10000 == 0 || k == nLoops) {
+      std::cout << "Info: Iteration = " << k << std::endl;
+    }
+  }
+  HIP_CHECK(hipDeviceSynchronize());
+  // Clean up
+  for (int i = 0; i < nStreams; i++) {
+    HIP_CHECK(hipStreamDestroy(streams[i]));
+  }
+}

catch/unit/stream/hipStreamCreate.cc

@@ -0,0 +1,52 @@
+
+/*
+Copyright (c) 2021-Present 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 <hip_test_common.hh>
+TEST_CASE("Unit_hipStreamCreate_default") {
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+}
+TEST_CASE("Unit_hipStreamCreateWithFlags_Negative") {
+  hipStream_t stream;
+  auto status = hipStreamCreateWithFlags(&stream, 0xFF);
+  REQUIRE(status == hipErrorInvalidValue);
+  status = hipStreamCreateWithFlags(nullptr, hipStreamDefault);
+  REQUIRE(status == hipErrorInvalidValue);
+}
+TEST_CASE("Unit_hipStreamCreateWithFlags") {
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+TEST_CASE("Unit_hipStreamCreateWithPriority") {
+  int priority_low = 0;
+  int priority_high = 0;
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream;
+
+  SECTION("Setting high prirority") {
+    HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
+  }
+  SECTION("Setting low priority") {
+    HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
+  }
+  HIP_CHECK(hipStreamDestroy(stream));
+}

catch/unit/stream/hipStreamGetFlags.cc

@@ -0,0 +1,38 @@
+
+/*
+Copyright (c) 2021-Present 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 <hip_test_common.hh>
+TEST_CASE("Unit_hipStreamGetFlags_Negative") {
+  // Get flags for uninitialized stream
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
+  REQUIRE(hipStreamGetFlags(stream, nullptr) == hipErrorInvalidValue);
+}
+TEST_CASE("Unit_hipStreamGetFlags") {
+    hipStream_t stream;
+    unsigned int flags;
+    HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
+    HIP_CHECK(hipStreamGetFlags(stream, &flags));
+    REQUIRE(flags == hipStreamDefault);
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
+    HIP_CHECK(hipStreamGetFlags(stream, &flags));
+    REQUIRE(flags == hipStreamNonBlocking);
+    HIP_CHECK(hipStreamDestroy(stream));
+}

catch/unit/stream/hipStreamGetPriority.cc

@@ -0,0 +1,92 @@
+
+/*
+Copyright (c) 2021-Present 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 <hip_test_common.hh>
+TEST_CASE("Unit_hipStreamGetPriority_Negative") {
+  hipStream_t stream = 0;
+  REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue);
+}
+TEST_CASE("Unit_hipStreamGetPriority_default") {
+  int priority_low = 0;
+  int priority_high = 0;
+  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
+  HIP_CHECK(hipSetDevice(devID));
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  int priority = 0;
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  // valid priority
+  // Lower the value higher the priority, higher the value lower the priority
+  REQUIRE(priority_low >= priority);
+  REQUIRE(priority >= priority_high);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+TEST_CASE("Unit_hipStreamGetPriority_high") {
+  int priority_low = 0;
+  int priority_high = 0;
+  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
+  HIP_CHECK(hipSetDevice(devID));
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
+  int priority = 0;
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  REQUIRE(priority == priority_high);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+TEST_CASE("Unit_hipStreamGetPriority_higher") {
+  int priority_low = 0;
+  int priority_high = 0;
+  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
+  HIP_CHECK(hipSetDevice(devID));
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1));
+  int priority = 0;
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  REQUIRE(priority == priority_high);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+TEST_CASE("Unit_hipStreamGetPriority_low") {
+  int priority_low = 0;
+  int priority_high = 0;
+  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
+  HIP_CHECK(hipSetDevice(devID));
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
+  int priority = 0;
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  REQUIRE(priority_low == priority);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+TEST_CASE("Unit_hipStreamGetPriority_lower") {
+  int priority_low = 0;
+  int priority_high = 0;
+  int devID = GENERATE(range(0, HipTest::getDeviceCount()));
+  HIP_CHECK(hipSetDevice(devID));
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1));
+  int priority = 0;
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  REQUIRE(priority_low == priority);
+  HIP_CHECK(hipStreamDestroy(stream));
+}