SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: Id77c54027eb0b9960e78887af75cfc3666d8568b

INSTALL.md

@@ -6,10 +6,15 @@
   * [Prerequisites](#prerequisites)
   * [AMD Platform](#amd-platform)
   * [NVIDIA Platform](#nvidia-platform)
-- [Building HIP from source](#building-hip-from-source)
-  * [Build ROCclr](#build-rocclr)
+- [Building HIP from source on AMD platform](#building-hip-from-source-on-amd-platform)
+  * [Get HIP source code](#get-hip-source-code)
+  * [Set the environment variables](#set-the-environment-variables)
   * [Build HIP](#build-hip)
   * [Default paths and environment variables](#default-paths-and-environment-variables)
+- [Building HIP from source on NVIDIA platform](#building-hip-from-source-on-NVIDIA-platform)
+  * [Get HIP source code](#get-hip-source-code)
+  * [Set the environment variables](#set-the-environment-variables)
+  * [Build HIP](#build-hip)
 - [Verify your installation](#verify-your-installation)
 <!-- tocstop -->
 
@@ -57,17 +62,27 @@ sudo make install
 
 HIP-nvcc is the compiler for HIP program compilation on NVIDIA platform.
 
-* Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories).
-* Install the "hip-runtime-nvidia" and "hip-devel" package.  This will install CUDA SDK and the HIP porting layer.
+* Install Nvidia Driver
 ```
-apt-get install hip-runtime-nvidia hip-devel
+sudo apt-get install ubuntu-drivers-common && sudo ubuntu-drivers autoinstall
+sudo reboot
+```
+Or download the latest cuda-toolkit at https://developer.nvidia.com/cuda-downloads
+The driver will be installed automatically.
+
+* Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories).
+* Install the "hip-runtime-nvidia" and "hip-dev" package.  This will install CUDA SDK and the HIP porting layer.
+```
+apt-get install hip-runtime-nvidia hip-dev
 ```
 * Default paths and environment variables:
    * By default HIP looks for CUDA SDK in /usr/local/cuda (can be overriden by setting CUDA_PATH env variable).
    * By default HIP is installed into <ROCM_PATH>/hip (can be overridden by setting HIP_PATH environment variable).
    * Optionally, consider adding <ROCM_PATH>/bin to your path to make it easier to use the tools.
 
-# Building HIP from source
+
+# Building HIP from source on AMD platform
+
 
 ## Get HIP source code
 
@@ -117,12 +132,45 @@ By default, release version of AMDHIP is built.
 
 After installation, make sure HIP_PATH is pointed to /where/to/install/hip
 
+
+# Building HIP from source on NVIDIA platform
+
+
+## Get HIP source code
+
+```
+git clone -b develop https://github.com/ROCm-Developer-Tools/hip.git
+git clone -b develop https://github.com/ROCm-Developer-Tools/hipamd.git
+```
+
+## Set the environment variables
+
+```
+export HIP_DIR="$(readlink -f hip)"
+export HIPAMD_DIR="$(readlink -f hipamd)"
+```
+
+## Build HIP
+
+```
+cd "$HIPAMD_DIR"
+mkdir -p build; cd build
+cmake -DHIP_COMMON_DIR=$HIP_DIR -DHIP_PLATFORM=nvidia -DCMAKE_INSTALL_PREFIX=$PWD/install ..
+make -j$(nproc)
+sudo make install
+```
+
 # Verify your installation
 
 Run hipconfig (instructions below assume default installation path) :
 ```shell
 <ROCM_PATH>/bin/hipconfig --full
 ```
+or
 
-Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/main/samples/0_Intro/square).
+```shell
+$PWD/install/bin/hipconfig --full
+```
+
+Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/rocm-4.5.x/samples/0_Intro/square).
 

include/hip/hip_runtime_api.h

@@ -4276,7 +4276,7 @@ typedef enum hipGraphExecUpdateResult {
                                   ///< in the return value of the function
   hipGraphExecUpdateErrorTopologyChanged = 0x2,  ///< The update failed because the topology changed
   hipGraphExecUpdateErrorNodeTypeChanged = 0x3,  ///< The update failed because a node type changed
-  hipGraphExecUpdateErrorFunctionChanged = 
+  hipGraphExecUpdateErrorFunctionChanged =
       0x4,  ///< The update failed because the function of a kernel node changed
   hipGraphExecUpdateErrorParametersChanged =
       0x5,  ///< The update failed because the parameters changed in a way that is not supported
@@ -5223,6 +5223,7 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph
  *  This section describes Stream Memory Wait and Write functions of HIP runtime API.
  */
 typedef unsigned int GLuint;
+typedef unsigned int GLenum;
 
 // Queries devices associated with GL Context.
 hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices,
@@ -5230,9 +5231,15 @@ hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices,
 // Registers a GL Buffer for interop and returns corresponding graphics resource.
 hipError_t hipGraphicsGLRegisterBuffer(hipGraphicsResource** resource, GLuint buffer,
                                        unsigned int flags);
+// Register a GL Image for interop and returns the corresponding graphic resource
+hipError_t hipGraphicsGLRegisterImage(hipGraphicsResource** resource, GLuint image,
+                                      GLenum target, unsigned int flags);
 // Maps a graphics resource for hip access.
 hipError_t hipGraphicsMapResources(int count, hipGraphicsResource_t* resources,
                                    hipStream_t stream  __dparm(0) );
+// Get an array through which to access a subresource of a mapped graphics resource.
+hipError_t hipGraphicsSubResourceGetMappedArray(hipArray_t* array, hipGraphicsResource_t resource,
+                                                unsigned int arrayIndex, unsigned int mipLevel);
 // Gets device accessible address of a graphics resource.
 hipError_t hipGraphicsResourceGetMappedPointer(void** devPtr, size_t* size,
                                                hipGraphicsResource_t resource);

samples/2_Cookbook/8_peer2peer/peer2peer.cpp

@@ -174,6 +174,7 @@ int main() {
                     dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y),
                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix[0],
                     data[0], width);
+    HIPCHECK(hipDeviceSynchronize());
 
     HIPCHECK(hipSetDevice(peerGpu));
     TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));

tests/catch/CMakeLists.txt

@@ -1,4 +1,5 @@
 cmake_minimum_required(VERSION 3.16.8)
+project(hiptests)
 
 # Check if platform and compiler are set
 if(HIP_PLATFORM STREQUAL "amd")
@@ -33,12 +34,6 @@ else()
     cmake_path(SET CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc.bat")
 endif()
 
-if(NOT UNIX)
-    # In linux this reruns the cmake and fails with incorrect vars.
-    # so the project command is used only for windows
-    project(hiptests)
-endif()
-
 if(NOT DEFINED CATCH2_PATH)
     if(DEFINED ENV{CATCH2_PATH})
         set(CATCH2_PATH $ENV{CATCH2_PATH} CACHE STRING "Catch2 Path")

tests/catch/unit/CMakeLists.txt

@@ -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)

tests/catch/unit/memory/hipMemCoherencyTst.cc

@@ -91,6 +91,31 @@ static void TstCoherency(int *Ptr, bool HmmMem) {
   }
 }
 
+static int HmmAttrPrint() {
+  int managed = 0;
+  INFO("The following are the attribute values related to HMM for"
+         " device 0:\n");
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
+  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributeConcurrentManagedAccess, 0));
+  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
+  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
+         << managed);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  return managed;
+}
+
+
 /* Test case description: The following test validates if fine grain
    behavior is observed or not with memory allocated using hipHostMalloc()*/
 // The following tests are disabled for Nvidia as they are not consistently
@@ -127,16 +152,23 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
   int *Ptr = nullptr, SIZE = sizeof(int);
   bool HmmMem = true;
   YES_COHERENT = false;
-  // Allocating hipMallocManaged() memory
-  SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
-    HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
+
+  int managed =  HmmAttrPrint();
+  if (managed == 1) {
+    // Allocating hipMallocManaged() memory
+    SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
+      HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
+    }
+    SECTION("hipMallocManaged with hipMemAttachHost flag") {
+      HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
+    }
+    TstCoherency(Ptr, HmmMem);
+    HIP_CHECK(hipFree(Ptr));
+    REQUIRE(YES_COHERENT);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
   }
-  SECTION("hipMallocManaged with hipMemAttachHost flag") {
-    HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
-  }
-  TstCoherency(Ptr, HmmMem);
-  HIP_CHECK(hipFree(Ptr));
-  REQUIRE(YES_COHERENT);
 }
 #endif
 
@@ -197,37 +229,44 @@ TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
   int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9;
   bool FineGrain = true;
   YES_COHERENT = false;
-  // Allocating hipExtMallocWithFlags() memory with flags
-  SECTION("hipExtMallocWithFlags with hipDeviceMallocFinegrained flag") {
-    HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
-                                    hipDeviceMallocFinegrained));
-  }
-  SECTION("hipExtMallocWithFlags with hipDeviceMallocSignalMemory flag") {
-    // for hipMallocSignalMemory flag the size of memory must be 8
-    HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
-                                    hipMallocSignalMemory));
-  }
-  SECTION("hipExtMallocWithFlags with hipDeviceMallocDefault flag") {
-    /* hipExtMallocWithFlags() with flag
-    hipDeviceMallocDefault allocates CoarseGrain memory */
-    FineGrain = false;
-    HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
-                                    hipDeviceMallocDefault));
-  }
-  if (FineGrain) {
-    TstCoherency(Ptr, FineGrain);
-  } else {
-    *Ptr = InitVal;
-    hipStream_t strm;
-    HIP_CHECK(hipStreamCreate(&strm));
-    SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
-    HIP_CHECK(hipStreamSynchronize(strm));
-    if (*Ptr == (InitVal * InitVal)) {
-      YES_COHERENT = true;
+
+  int managed =  HmmAttrPrint();
+  if (managed == 1) {
+    // Allocating hipExtMallocWithFlags() memory with flags
+    SECTION("hipExtMallocWithFlags with hipDeviceMallocFinegrained flag") {
+      HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
+                                      hipDeviceMallocFinegrained));
     }
+    SECTION("hipExtMallocWithFlags with hipDeviceMallocSignalMemory flag") {
+      // for hipMallocSignalMemory flag the size of memory must be 8
+      HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
+                                      hipMallocSignalMemory));
+    }
+    SECTION("hipExtMallocWithFlags with hipDeviceMallocDefault flag") {
+      /* hipExtMallocWithFlags() with flag
+      hipDeviceMallocDefault allocates CoarseGrain memory */
+      FineGrain = false;
+      HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
+                                      hipDeviceMallocDefault));
+    }
+    if (FineGrain) {
+      TstCoherency(Ptr, FineGrain);
+    } else {
+      *Ptr = InitVal;
+      hipStream_t strm;
+      HIP_CHECK(hipStreamCreate(&strm));
+      SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
+      HIP_CHECK(hipStreamSynchronize(strm));
+      if (*Ptr == (InitVal * InitVal)) {
+        YES_COHERENT = true;
+      }
+    }
+    HIP_CHECK(hipFree(Ptr));
+    REQUIRE(YES_COHERENT);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
   }
-  HIP_CHECK(hipFree(Ptr));
-  REQUIRE(YES_COHERENT);
 }
 #endif
 

tests/catch/unit/memory/hipMemcpy3D.cc

@@ -272,6 +272,7 @@ void Memcpy3D<T>::NegativeTests() {
   AllocateMemory();
 
   // Initialization of data
+  memset(&myparms, 0, sizeof(myparms));
   myparms.srcPos = make_hipPos(0, 0, 0);
   myparms.dstPos = make_hipPos(0, 0, 0);
   myparms.extent = make_hipExtent(width , height, depth);

tests/catch/unit/memory/hipMemcpy3DAsync.cc

@@ -290,6 +290,7 @@ void Memcpy3DAsync<T>::NegativeTests() {
   HIP_CHECK(hipStreamCreate(&stream));
 
   // Initialization of data
+  memset(&myparms, 0, sizeof(myparms));
   myparms.srcPos = make_hipPos(0, 0, 0);
   myparms.dstPos = make_hipPos(0, 0, 0);
   myparms.extent = make_hipExtent(width , height, depth);

tests/catch/unit/streamperthread/CMakeLists.txt

@@ -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)

tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc

@@ -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);
+  }
+}

tests/catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc

@@ -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);
+}

tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc

@@ -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);
+    }
+  }
+}

tests/catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc

@@ -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();
+  }
+}