SWDEV-292643 - [dtest] Catch2 additional unit tests for stream management apis. (#2349)

APIs covered :
hipStreamGetPriority, hipStreamCreate, hipStreamGetFlags, hipExtStreamGetCUMask apis.

Change-Id: I238b4e631938471eab05c598f91477eeb0856054

[ROCm/hip commit: bef1aa7ee6]

projects/hip/tests/catch/unit/stream/CMakeLists.txt

@@ -9,6 +9,8 @@ set(TEST_SRC
     hipStreamCreateWithFlags.cc
     hipStreamCreateWithPriority.cc
     hipStreamWithCUMask.cc
+    hipStreamGetCUMask.cc
+    hipAPIStreamDisable.cc
 )
 else()
 set(TEST_SRC
@@ -20,6 +22,7 @@ set(TEST_SRC
     hipStreamAddCallback.cc
     hipStreamCreateWithFlags.cc
     hipStreamCreateWithPriority.cc
+    hipAPIStreamDisable.cc
 )
 endif()
 

projects/hip/tests/catch/unit/stream/hipAPIStreamDisable.cc

@@ -0,0 +1,68 @@
+/*
+Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <hip_test_common.hh>
+#include "hip/math_functions.h"
+
+#define NUM_STREAMS 8
+
+namespace hipAPIStreamDisableTest {
+const int NN = 1 << 21;
+
+__global__ void kernel(float* x, float* y, int n) {
+  int 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) {
+  int tid = threadIdx.x;
+  y[tid] = y[tid] + 1.0f;
+}
+}  // namespace hipAPIStreamDisableTest
+
+/**
+ * Validate basic multistream functionalities
+ */
+TEST_CASE("Unit_hipStreamCreate_MultistreamBasicFunctionalities") {
+  hipStream_t streams[NUM_STREAMS];
+  float *data[NUM_STREAMS], *yd, *xd;
+  float y = 1.0f, x = 1.0f;
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&yd), sizeof(float)));
+  HIP_CHECK(hipMalloc(reinterpret_cast<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],
+            (hipAPIStreamDisableTest::NN)*sizeof(float)));
+    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::kernel),
+                       dim3(1), dim3(1), 0, streams[i], data[i], xd,
+                       hipAPIStreamDisableTest::NN);
+    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::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 == y);
+}

projects/hip/tests/catch/unit/stream/hipStreamGetCUMask.cc

@@ -0,0 +1,177 @@
+/*
+Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/**
+Testcase Scenarios :
+1) Test to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask.
+2) Test to verify hipExtStreamGetCUMask api returns custom mask set.
+3) Negative tests for hipExtStreamGetCUMask api.
+*/
+
+#include <hip_test_common.hh>
+#include <vector>
+
+
+/**
+ * Scenario to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask.
+ * Scenario to verify hipExtStreamGetCUMask api returns custom mask set.
+ */
+TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") {
+  constexpr int maxNum = 6;
+  std::vector<uint32_t> cuMask(maxNum);
+  hipDeviceProp_t props;
+  std::stringstream ss;
+  char* gCUMask{nullptr};
+  std::string globalCUMask("");
+  std::vector<uint32_t> defaultCUMask;
+
+  int nGpu = 0;
+  HIP_CHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+    INFO("info: didn't find any GPU! skipping the test!");
+    return;
+  }
+
+  HIP_CHECK(hipSetDevice(0));
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  INFO("info: running on bus " << "0x" << props.pciBusID << " " <<
+         props.name << " with " << props.multiProcessorCount << " CUs");
+
+  // Get global CU Mask if exists
+  gCUMask = getenv("ROC_GLOBAL_CU_MASK");
+  if (gCUMask != nullptr && gCUMask[0] != '\0') {
+    globalCUMask.assign(gCUMask);
+
+    for_each(globalCUMask.begin(), globalCUMask.end(), [](char & c) {
+      c = ::tolower(c);
+    });
+  }
+
+  // Create default CU Mask
+  uint32_t temp = 0;
+  uint32_t bit_index = 0;
+  for (uint32_t i = 0; i < (uint32_t)props.multiProcessorCount; i++) {
+    temp |= 1UL << bit_index;
+    if (bit_index >= 32) {
+      defaultCUMask.push_back(temp);
+      temp = 0;
+      bit_index = 0;
+      temp |= 1UL << bit_index;
+    }
+    bit_index += 1;
+  }
+  if (bit_index != 0) {
+    defaultCUMask.push_back(temp);
+  }
+
+  SECTION("Verify with default CU Mask or global CU Mask") {
+    // make a default CU mask bit-array where all CUs are active
+    // this default mask is expected to be returned when there is no
+    // custom or global CU mask defined
+
+    HIP_CHECK(hipExtStreamGetCUMask(0, cuMask.size(), &cuMask[0]));
+
+    ss << std::hex;
+    for (int i = cuMask.size() - 1; i >= 0; i--) {
+      ss << cuMask[i];
+    }
+
+    // remove extra 0 from ss if any present
+    size_t found = ss.str().find_first_not_of("0");
+    if (found != std::string::npos) {
+      ss.str(ss.str().substr(found, ss.str().length()));
+    }
+
+    INFO("info: CU mask for the default stream is: 0x" << ss.str().c_str());
+    if (globalCUMask.size() > 0) {
+      if (ss.str().compare(globalCUMask) != 0) {
+        INFO("Expected CU mask:" << globalCUMask.c_str() <<
+                                     ", api returned:" << ss.str().c_str());
+        REQUIRE(false);
+      }
+    } else {
+      for (int i = 0 ; i < min(cuMask.size(), defaultCUMask.size()); i++) {
+        if (cuMask[i] != defaultCUMask[i]) {
+          INFO("Expected CU mask " << defaultCUMask[i] <<
+                                      ", api returned:" << cuMask[i]);
+          REQUIRE(false);
+        }
+      }
+    }
+  }
+
+  SECTION("Verify with custom mask set") {
+    std::vector<uint32_t> customMask(defaultCUMask);
+    hipStream_t stream;
+    customMask[0] = 0xe;
+
+    HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, customMask.size(),
+                                                      customMask.data()));
+    ss.str("");
+    for (int i = customMask.size() - 1; i >= 0; i--) {
+      ss << customMask[i];
+    }
+    INFO("info: setting a custom CU mask 0x" << ss.str());
+
+    HIP_CHECK(hipExtStreamGetCUMask(stream, cuMask.size(), &cuMask[0]));
+    ss.str("");
+    for (int i = cuMask.size() - 1; i >= 0; i--) {
+      ss << cuMask[i];
+    }
+
+    size_t found = ss.str().find_first_not_of("0");
+    if (found != std::string::npos) {
+      ss.str(ss.str().substr(found, ss.str().length()));
+    }
+
+    INFO("info: reading back CU mask 0x" << ss.str() <<
+                                                " for stream " << stream);
+
+    if (!gCUMask) {
+      for (size_t i = 0; i < customMask.size(); i++) {
+        if (customMask[i] != cuMask[i]) {
+          INFO("Error! expected CU mask:" << customMask[i]
+                            << ", Received CU mask:" << cuMask[i]);
+          REQUIRE(false);
+        }
+      }
+    }
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+/**
+ * Negative tests for hipExtStreamGetCUMask.
+ */
+TEST_CASE("Unit_hipExtStreamGetCUMask_Negative") {
+  hipError_t ret;
+  constexpr int maxNum = 6;
+  std::vector<uint32_t> cuMask(maxNum);
+
+  SECTION("cuMask is nullptr") {
+    ret = hipExtStreamGetCUMask(0, cuMask.size(), nullptr);
+    REQUIRE(ret == hipErrorInvalidValue);
+  }
+
+  SECTION("cuMaskSize is 0") {
+    ret = hipExtStreamGetCUMask(0, 0, &cuMask[0]);
+    REQUIRE(ret == hipErrorInvalidValue);
+  }
+}

projects/hip/tests/catch/unit/stream/hipStreamGetFlags.cc

@@ -16,16 +16,24 @@ 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.
 */
+/**
+Testcase Scenarios :
+1) Test flag value of stream created with hipStreamCreateWithFlags/
+ /hipStreamCreate/hipStreamCreateWithPriority.
+2) Negative tests for hipStreamGetFlags api.
+3) Test flag value when streams created with CUMask.
+*/
+
 #include <hip_test_common.hh>
-TEST_CASE("Unit_hipStreamGetFlags_Negative") {
-  // Get flags for uninitialized stream
+
+/**
+ * Test flag value of stream created with various types.
+ */
+TEST_CASE("Unit_hipStreamGetFlags_BasicFunctionalities") {
   hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
-  REQUIRE(hipStreamGetFlags(stream, nullptr) == hipErrorInvalidValue);
-}
-TEST_CASE("Unit_hipStreamGetFlags") {
-    hipStream_t stream;
-    unsigned int flags;
+  unsigned int flags;
+  // Check flag value of stream created with hipStreamCreateWithFlags
+  SECTION("Check flag value of streams hipStreamCreateWithFlags") {
     HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
     HIP_CHECK(hipStreamGetFlags(stream, &flags));
     REQUIRE(flags == hipStreamDefault);
@@ -34,4 +42,50 @@ TEST_CASE("Unit_hipStreamGetFlags") {
     HIP_CHECK(hipStreamGetFlags(stream, &flags));
     REQUIRE(flags == hipStreamNonBlocking);
     HIP_CHECK(hipStreamDestroy(stream));
+  }
+  // Check flag value of stream created with hipStreamCreate
+  SECTION("Check flag value of streams hipStreamCreate") {
+    HIP_CHECK(hipStreamCreate(&stream));
+    HIP_CHECK(hipStreamGetFlags(stream, &flags));
+    REQUIRE(flags == hipStreamDefault);
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+  // Check flag value of stream created with hipStreamCreateWithPriority
+  SECTION("Check flag value of streams hipStreamCreateWithPriority") {
+    HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, 0));
+    HIP_CHECK(hipStreamGetFlags(stream, &flags));
+    REQUIRE(flags == hipStreamDefault);
+    HIP_CHECK(hipStreamDestroy(stream));
+    HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, 0));
+    HIP_CHECK(hipStreamGetFlags(stream, &flags));
+    REQUIRE(flags == hipStreamNonBlocking);
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
 }
+
+/**
+ * Negative Scenarios
+ */
+TEST_CASE("Unit_hipStreamGetFlags_Negative") {
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  // nullptr check
+  REQUIRE(hipStreamGetFlags(stream, nullptr) != hipSuccess);
+  REQUIRE(hipStreamGetFlags(nullptr, hipStreamDefault) != hipSuccess);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+#if HT_AMD
+/**
+ * Test flag value when streams created with CUMask.
+ */
+TEST_CASE("Unit_hipStreamGetFlags_StreamsCreatedWithCUMask") {
+  hipStream_t stream;
+  unsigned int flags;
+  const uint32_t cuMask = 0xffffffff;
+  HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask));
+  HIP_CHECK(hipStreamGetFlags(stream, &flags));
+  REQUIRE(flags == hipStreamDefault);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+#endif

projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc

@@ -16,11 +16,28 @@ 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.
 */
+/**
+Testcase Scenarios :
+1) Negative tests for hipStreamGetPriority api.
+2) Create stream and check default priority of stream is within range.
+3) Create stream with high or low priority and check priority is set as expected.
+4) Create stream with higher priority or lower priority for the priority range returned.
+5) Create stream with CUMask and check priority is returned as expected.
+*/
+
 #include <hip_test_common.hh>
+
+/**
+ * Negative tests for hipStreamGetPriority api.
+ */
 TEST_CASE("Unit_hipStreamGetPriority_Negative") {
   hipStream_t stream = 0;
   REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue);
 }
+
+/**
+ * Create stream and check default priority of stream is within range.
+ */
 TEST_CASE("Unit_hipStreamGetPriority_default") {
   int priority_low = 0;
   int priority_high = 0;
@@ -37,6 +54,10 @@ TEST_CASE("Unit_hipStreamGetPriority_default") {
   REQUIRE(priority >= priority_high);
   HIP_CHECK(hipStreamDestroy(stream));
 }
+
+/**
+ * Create stream with high priority and check priority is set as expected.
+ */
 TEST_CASE("Unit_hipStreamGetPriority_high") {
   int priority_low = 0;
   int priority_high = 0;
@@ -44,12 +65,17 @@ TEST_CASE("Unit_hipStreamGetPriority_high") {
   HIP_CHECK(hipSetDevice(devID));
   HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
   hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
+                                                            priority_high));
   int priority = 0;
   HIP_CHECK(hipStreamGetPriority(stream, &priority));
   REQUIRE(priority == priority_high);
   HIP_CHECK(hipStreamDestroy(stream));
 }
+
+/**
+ * Create stream with higher priority for the priority range returned.
+ */
 TEST_CASE("Unit_hipStreamGetPriority_higher") {
   int priority_low = 0;
   int priority_high = 0;
@@ -57,12 +83,17 @@ TEST_CASE("Unit_hipStreamGetPriority_higher") {
   HIP_CHECK(hipSetDevice(devID));
   HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
   hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1));
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
+                                                           priority_high-1));
   int priority = 0;
   HIP_CHECK(hipStreamGetPriority(stream, &priority));
   REQUIRE(priority == priority_high);
   HIP_CHECK(hipStreamDestroy(stream));
 }
+
+/**
+ * Create stream with low priority and check priority is set as expected.
+ */
 TEST_CASE("Unit_hipStreamGetPriority_low") {
   int priority_low = 0;
   int priority_high = 0;
@@ -70,12 +101,17 @@ TEST_CASE("Unit_hipStreamGetPriority_low") {
   HIP_CHECK(hipSetDevice(devID));
   HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
   hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
+                                                              priority_low));
   int priority = 0;
   HIP_CHECK(hipStreamGetPriority(stream, &priority));
   REQUIRE(priority_low == priority);
   HIP_CHECK(hipStreamDestroy(stream));
 }
+
+/**
+ * Create stream with lower priority for the priority range returned.
+ */
 TEST_CASE("Unit_hipStreamGetPriority_lower") {
   int priority_low = 0;
   int priority_high = 0;
@@ -83,9 +119,35 @@ TEST_CASE("Unit_hipStreamGetPriority_lower") {
   HIP_CHECK(hipSetDevice(devID));
   HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
   hipStream_t stream;
-  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1));
+  HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
+                                                           priority_low+1));
   int priority = 0;
   HIP_CHECK(hipStreamGetPriority(stream, &priority));
   REQUIRE(priority_low == priority);
   HIP_CHECK(hipStreamDestroy(stream));
 }
+
+#if HT_AMD
+/**
+ * Create stream with CUMask and check priority is returned as expected.
+ */
+TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") {
+  hipStream_t stream;
+  int priority = 0;
+  int priority_normal = 0;
+  int priority_low = 0;
+  int priority_high = 0;
+  // Test is to get the Stream Priority Range
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  priority_normal = priority_low + priority_high;
+  // Check if priorities are indeed supported
+  REQUIRE_FALSE(priority_low == priority_high);
+  // Creating a stream with hipExtStreamCreateWithCUMask and checking
+  // priority.
+  const uint32_t cuMask = 0xffffffff;
+  HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask));
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  REQUIRE_FALSE(priority_normal != priority);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+#endif