SWDEV-292547 - dtests for hipStreamPerThread (#2417)

[ROCm/hip-tests commit: d756a47a41]
2022-01-06 16:18:30 +05:30
parent 385d8ae8ef
commit 0390880b84
6 changed files with 355 additions and 0 deletions
@@ -21,6 +21,7 @@
 add_subdirectory(memory)
 add_subdirectory(deviceLib)
 add_subdirectory(stream)
+add_subdirectory(streamperthread)
 add_subdirectory(event)
 add_subdirectory(occupancy)
 add_subdirectory(device)
@@ -0,0 +1,13 @@
+# Common Tests - Test independent of all platforms
+set(TEST_SRC
+  hipStreamPerThread_Basic.cc
+  hipStreamPerThread_Event.cc
+  hipStreamPerThread_MultiThread.cc
+  hipStreamPerThread_DeviceReset.cc
+)
+
+# Create shared lib of all tests
+add_library(StreamPerThreadTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+
+# Add dependency on build_tests to build it on this custom target
+add_dependencies(build_tests StreamPerThreadTest)
@@ -0,0 +1,133 @@
+/*
+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>
+
+#define MEM_SIZE (1024*1024*32)
+#define SEED 5
+
+constexpr unsigned int MAX_THREAD_CNT = 10;
+
+__global__ void copy_kernl(int* devPtr) {
+  for (int i = 0; i < MEM_SIZE; ++i) {
+    devPtr[i] = (i+1) + SEED;
+  }
+}
+
+TEST_CASE("Unit_hipStreamPerThread_Basic") {
+  constexpr int size = sizeof(int) * MEM_SIZE;
+  int* hostMem = nullptr;
+  int* devMem = nullptr;
+
+  HIP_CHECK(hipHostMalloc(&hostMem, size));
+  HIP_CHECK(hipMalloc(&devMem, size));
+
+  // Init host mem with different values
+  for (int i = 0; i < MEM_SIZE; ++i) {
+    hostMem[i] = i;
+  }
+
+  /*
+   hipStreamPerThread is an implicit stream which works independent of null stream.
+   Null stream synchronize will account hipStreamPerThread into account.
+   test scenario: Launch kernel + Async mem copy on hipStreamPerThread and call synchronize on null stream.
+   Result : Null stream synchronize should sync hipStreamPerThread as well
+   */
+  copy_kernl<<<1, 1, 0, hipStreamPerThread>>>(devMem);
+
+  HIP_CHECK(hipMemcpyAsync(hostMem, devMem, size, hipMemcpyDeviceToHost, hipStreamPerThread));
+
+  HIP_CHECK(hipStreamSynchronize(0));
+
+  // validate result
+  for (int i = MEM_SIZE-1; i >= 0; --i) {
+    CHECK(hostMem[i] == (i+1+SEED));
+  }
+}
+
+TEST_CASE("Unit_hipStreamPerThread_StreamQuery") {
+  std::vector<std::thread> threads(MAX_THREAD_CNT);
+
+  for (auto &th : threads) {
+    th = std::thread([](){HIP_CHECK(hipStreamQuery(hipStreamPerThread));});
+  }
+
+  for (auto& th : threads) {
+    th.join();
+  }
+  REQUIRE(true);
+}
+
+TEST_CASE("Unit_hipStreamPerThread_StreamSynchronize") {
+  constexpr unsigned int MAX_THREAD_CNT = 10;
+  std::vector<std::thread> threads(MAX_THREAD_CNT);
+
+  for (auto &th : threads) {
+    th = std::thread([](){HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));});
+  }
+
+  for (auto& th : threads) {
+    th.join();
+  }
+  REQUIRE(true);
+}
+
+TEST_CASE("Unit_hipStreamPerThread_StreamGetPriority") {
+  int priority = 0;
+  HIP_CHECK(hipStreamGetPriority(hipStreamPerThread, &priority));
+}
+
+TEST_CASE("Unit_hipStreamPerThread_StreamGetFlags") {
+  unsigned int flags = 0;
+  HIP_CHECK(hipStreamGetFlags(hipStreamPerThread, &flags));
+}
+
+TEST_CASE("Unit_hipStreamPerThread_StreamDestroy") {
+  hipError_t status = hipStreamDestroy(hipStreamPerThread);
+  REQUIRE(status != hipSuccess);
+}
+
+TEST_CASE("Unit_hipStreamPerThread_MemcpyAsync") {
+  unsigned int ele_size = (16 * 1024);  // 16KB
+  int* A_h = nullptr;
+  int* A_d = nullptr;
+
+  HIP_CHECK(hipHostMalloc(&A_h, ele_size*sizeof(int)));
+  HIP_CHECK(hipMalloc(&A_d, ele_size * sizeof(int)));
+
+  for (unsigned int i = 0; i < ele_size; ++i) {
+    A_h[i] = 123;
+  }
+
+  HIP_CHECK(hipMemcpy(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice));
+
+  // Rest host memory
+  for (unsigned int i = 0; i < ele_size; ++i) {
+    A_h[i] = 0;
+  }
+
+  HIP_CHECK(hipMemcpyAsync(A_h, A_d, ele_size * sizeof(int), hipMemcpyDeviceToHost,
+                           hipStreamPerThread));
+  HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+
+  // Verify result
+  for (unsigned int i = 0; i < ele_size; ++i) {
+    REQUIRE(A_h[i] == 123);
+  }
+}
@@ -0,0 +1,89 @@
+/*
+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 <vector>
+#include <thread>
+
+/*
+ hipDeviceReset deletes all active streams including hipStreamPerThread.
+ Scenario: App calls hipDeviceReset while in other thread some Async operation is in
+          progress on hipStreamPerThread.
+ Watch out: hipDeviceRest should be successfull without any crash
+ */
+static void Copy_to_device() {
+  unsigned int ele_size = (32 * 1024);  // 32KB
+  int* A_h = nullptr;
+  int* A_d = nullptr;
+
+  hipError_t status =  hipHostMalloc(&A_h, ele_size*sizeof(int));
+  if (status != hipSuccess) return;
+
+  status = hipMalloc(&A_d, ele_size * sizeof(int));
+  if (status != hipSuccess) return;
+
+  for(unsigned int i = 0; i < ele_size; ++i) {
+    A_h[i] = 123;
+  }
+  hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
+                 hipStreamPerThread);
+}
+
+TEST_CASE("Unit_hipStreamPerThread_DeviceReset_1") {
+  constexpr unsigned int MAX_THREAD_CNT = 10;
+  std::vector<std::thread> threads(MAX_THREAD_CNT);
+
+  for (auto &th : threads) {
+    th = std::thread(Copy_to_device);
+    th.detach();
+  }
+  HIP_CHECK(hipDeviceReset());
+}
+
+/*
+ hipDeviceReset deletes all active streams including hipStreamPerThread.
+ Scenario: i) Launch Async task on hipStreamPerThread and waits for it to complete.
+          ii) Call hipDeviceReset to delete all active stream
+         iii) Again try to launch Async task on hipStreamPerThread
+ Watch out: Since hipStreamPerThread is an implicit stream hence even after device reset
+            it should available to use.
+ */
+TEST_CASE("Unit_hipStreamPerThread_DeviceReset_2") {
+  unsigned int ele_size = (32 * 1024);  // 32KB
+  int* A_h = nullptr;
+  int* A_d = nullptr;
+
+  hipError_t status =  hipHostMalloc(&A_h, ele_size*sizeof(int));
+  if (status != hipSuccess) return;
+  status = hipMalloc(&A_d, ele_size * sizeof(int));
+  if (status != hipSuccess) return;
+
+  for (unsigned int i = 0; i < ele_size; ++i) {
+    A_h[i] = 123;
+  }
+  hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
+                 hipStreamPerThread);
+  hipStreamSynchronize(hipStreamPerThread);
+
+  hipDeviceReset();
+
+  hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
+                 hipStreamPerThread);
+  hipStreamSynchronize(hipStreamPerThread);
+}
@@ -0,0 +1,64 @@
+/*
+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>
+
+TEST_CASE("Unit_hipStreamPerThread_EventRecord") {
+  hipEvent_t event;
+  HIP_CHECK(hipEventCreate(&event));
+  HIP_CHECK(hipEventRecord(event, hipStreamPerThread));
+}
+
+__global__ void update_even_odd(unsigned int N, int* out) {
+  for (unsigned int i = 0; i < N; ++i) {
+    if (i%2 == 0) {
+      out[i] = 2;
+    } else {
+      out[i] = 3;
+    }
+  }
+}
+TEST_CASE("Unit_hipStreamPerThread_EventSynchronize") {
+  int* A_h = nullptr;
+  int* A_d = nullptr;
+  unsigned int size = 1000;
+
+  HIP_CHECK(hipHostMalloc(&A_h, size*sizeof(int)));
+  HIP_CHECK(hipMalloc(&A_d, size * sizeof(int)));
+
+  hipEvent_t start, end;
+  HIP_CHECK(hipEventCreate(&start));
+  HIP_CHECK(hipEventCreate(&end));
+
+  HIP_CHECK(hipEventRecord(start, hipStreamPerThread));
+  update_even_odd<<<1, 1>>>(size, A_d);
+  HIP_CHECK(hipEventRecord(end, hipStreamPerThread));
+
+  HIP_CHECK(hipEventSynchronize(end));
+  HIP_CHECK(hipMemcpy(A_h, A_d, size*sizeof(int), hipMemcpyDeviceToHost));
+
+  // Verify result
+  for (unsigned int i = 0; i < size; ++i) {
+    if (i%2 == 0 && A_h[i] != 2)
+      REQUIRE(false);
+    else if (i%2 != 0 && A_h[i] != 3) {
+      REQUIRE(false);
+    }
+  }
+}
@@ -0,0 +1,55 @@
+/*
+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 <vector>
+#include <thread>
+
+static void Copy_to_device() {
+  unsigned int ele_size = (32 * 1024);  // 32KB
+  int* A_h = nullptr;
+  int* A_d = nullptr;
+
+  hipHostMalloc(&A_h, ele_size*sizeof(int));
+  hipMalloc(&A_d, ele_size * sizeof(int));
+
+  for (unsigned int i = 0; i < ele_size; ++i) {
+    A_h[i] = 123;
+  }
+  hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
+                 hipStreamPerThread);
+}
+
+/*
+hipStreamPerThread is an implicit stream which gets destroyed once thread is completed.
+Scenario : App pushes Async task(s) into hipStreamPerThread and did not wait for it to complete.
+Watch out : Incomplete task in hipStreamPerThread should not cause any crash due to thread exit.
+ */
+TEST_CASE("Unit_hipStreamPerThread_MultiThread") {
+  constexpr unsigned int MAX_THREAD_CNT = 10;
+  std::vector<std::thread> threads(MAX_THREAD_CNT);
+
+  for (auto &th : threads) {
+    th = std::thread(Copy_to_device);
+  }
+
+  for (auto& th : threads) {
+    th.detach();
+  }
+}