diff --git a/INSTALL.md b/INSTALL.md index 3f9305387e..0aecb06384 100644 --- a/INSTALL.md +++ b/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) @@ -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 /hip (can be overridden by setting HIP_PATH environment variable). * Optionally, consider adding /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 /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). diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 6fb56a0325..7d4d58ec24 100644 --- a/include/hip/hip_runtime_api.h +++ b/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); diff --git a/samples/2_Cookbook/8_peer2peer/peer2peer.cpp b/samples/2_Cookbook/8_peer2peer/peer2peer.cpp index c2de71ce40..4ec6372108 100644 --- a/samples/2_Cookbook/8_peer2peer/peer2peer.cpp +++ b/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)); diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index 56de9a6e63..67f5673d30 100644 --- a/tests/catch/CMakeLists.txt +++ b/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") diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 167a984c30..0e8b37c993 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/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) diff --git a/tests/catch/unit/memory/hipMemCoherencyTst.cc b/tests/catch/unit/memory/hipMemCoherencyTst.cc index f03c549d11..f04f9835d0 100644 --- a/tests/catch/unit/memory/hipMemCoherencyTst.cc +++ b/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(&Ptr), SIZE*2, - hipDeviceMallocFinegrained)); - } - SECTION("hipExtMallocWithFlags with hipDeviceMallocSignalMemory flag") { - // for hipMallocSignalMemory flag the size of memory must be 8 - HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&Ptr), SIZE*2, - hipMallocSignalMemory)); - } - SECTION("hipExtMallocWithFlags with hipDeviceMallocDefault flag") { - /* hipExtMallocWithFlags() with flag - hipDeviceMallocDefault allocates CoarseGrain memory */ - FineGrain = false; - HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&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(&Ptr), SIZE*2, + hipDeviceMallocFinegrained)); } + SECTION("hipExtMallocWithFlags with hipDeviceMallocSignalMemory flag") { + // for hipMallocSignalMemory flag the size of memory must be 8 + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&Ptr), SIZE*2, + hipMallocSignalMemory)); + } + SECTION("hipExtMallocWithFlags with hipDeviceMallocDefault flag") { + /* hipExtMallocWithFlags() with flag + hipDeviceMallocDefault allocates CoarseGrain memory */ + FineGrain = false; + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&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 diff --git a/tests/catch/unit/memory/hipMemcpy3D.cc b/tests/catch/unit/memory/hipMemcpy3D.cc index 1fa4ce0576..75fc8efe0c 100644 --- a/tests/catch/unit/memory/hipMemcpy3D.cc +++ b/tests/catch/unit/memory/hipMemcpy3D.cc @@ -272,6 +272,7 @@ void Memcpy3D::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); diff --git a/tests/catch/unit/memory/hipMemcpy3DAsync.cc b/tests/catch/unit/memory/hipMemcpy3DAsync.cc index 017556967c..fb2fa8e1dc 100644 --- a/tests/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy3DAsync.cc @@ -290,6 +290,7 @@ void Memcpy3DAsync::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); diff --git a/tests/catch/unit/streamperthread/CMakeLists.txt b/tests/catch/unit/streamperthread/CMakeLists.txt new file mode 100644 index 0000000000..51a789d426 --- /dev/null +++ b/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) diff --git a/tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc b/tests/catch/unit/streamperthread/hipStreamPerThread_Basic.cc new file mode 100644 index 0000000000..3780631141 --- /dev/null +++ b/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 + +#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 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 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); + } +} \ No newline at end of file diff --git a/tests/catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc b/tests/catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc new file mode 100644 index 0000000000..837446d2a9 --- /dev/null +++ b/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 +#include +#include + +/* + 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 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); +} \ No newline at end of file diff --git a/tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc b/tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc new file mode 100644 index 0000000000..6fef9914eb --- /dev/null +++ b/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 + +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); + } + } +} \ No newline at end of file diff --git a/tests/catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc b/tests/catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc new file mode 100644 index 0000000000..a0ff7ed6be --- /dev/null +++ b/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 +#include +#include + +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 threads(MAX_THREAD_CNT); + + for (auto &th : threads) { + th = std::thread(Copy_to_device); + } + + for (auto& th : threads) { + th.detach(); + } +}