diff --git a/projects/hip/bin/hipcc.pl b/projects/hip/bin/hipcc.pl index 606727f4cd..48ed8e1871 100755 --- a/projects/hip/bin/hipcc.pl +++ b/projects/hip/bin/hipcc.pl @@ -488,7 +488,7 @@ foreach $arg (@ARGV) my $fileType = `file $obj`; my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/); if ($fileType =~ m/ELF/) { - my $sections = `readelf -e -W $obj`; + my $sections = `$HIP_CLANG_PATH/llvm-readelf -e -W $obj`; $isObj = !($sections =~ m/__CLANG_OFFLOAD_BUNDLE__/); } $allIsObj = ($allIsObj and $isObj); diff --git a/projects/hip/include/hip/library_types.h b/projects/hip/include/hip/library_types.h index 3e844e6bad..1088cd4c90 100644 --- a/projects/hip/include/hip/library_types.h +++ b/projects/hip/include/hip/library_types.h @@ -28,12 +28,34 @@ THE SOFTWARE. #if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__)) typedef enum hipDataType { - HIP_R_16F = 2, - HIP_R_32F = 0, - HIP_R_64F = 1, - HIP_C_16F = 6, - HIP_C_32F = 4, - HIP_C_64F = 5 + HIP_R_32F = 0, + HIP_R_64F = 1, + HIP_R_16F = 2, + HIP_R_8I = 3, + HIP_C_32F = 4, + HIP_C_64F = 5, + HIP_C_16F = 6, + HIP_C_8I = 7, + HIP_R_8U = 8, + HIP_C_8U = 9, + HIP_R_32I = 10, + HIP_C_32I = 11, + HIP_R_32U = 12, + HIP_C_32U = 13, + HIP_R_16BF = 14, + HIP_C_16BF = 15, + HIP_R_4I = 16, + HIP_C_4I = 17, + HIP_R_4U = 18, + HIP_C_4U = 19, + HIP_R_16I = 20, + HIP_C_16I = 21, + HIP_R_16U = 22, + HIP_C_16U = 23, + HIP_R_64I = 24, + HIP_C_64I = 25, + HIP_R_64U = 26, + HIP_C_64U = 27 } hipDataType; typedef enum hipLibraryPropertyType { diff --git a/projects/hip/tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip/tests/catch/hipTestMain/config/config_amd_linux_common.json index 60ef935548..8b8c42fcd4 100644 --- a/projects/hip/tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip/tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -2,6 +2,8 @@ "DisabledTests": [ "Unit_hipStreamPerThread_DeviceReset_1", - "Unit_hipMallocManaged_OverSubscription" + "Unit_hipMallocManaged_OverSubscription", + "Unit_hipDeviceGetCacheConfig_Positive_Basic", + "Unit_hipDeviceGetCacheConfig_Positive_Threaded" ] } diff --git a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json index 07a4714117..2094761dc2 100644 --- a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -79,6 +79,8 @@ "Unit_hipStreamQuery_WithPendingWork", "Unit_hipStreamWaitEvent_DifferentStreams", "Unit_hipStreamQuery_WithFinishedWork", + "Unit_hipDeviceGetCacheConfig_Positive_Basic", + "Unit_hipDeviceGetCacheConfig_Positive_Threaded", "Unit_hipStreamValue_Wait32_Blocking_Mask_Gte", "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1", "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_2", diff --git a/projects/hip/tests/catch/unit/CMakeLists.txt b/projects/hip/tests/catch/unit/CMakeLists.txt index 0ce9be5ed9..6e97f24663 100644 --- a/projects/hip/tests/catch/unit/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/CMakeLists.txt @@ -31,6 +31,8 @@ add_subdirectory(texture) add_subdirectory(streamperthread) add_subdirectory(kernel) add_subdirectory(multiThread) +add_subdirectory(compiler) +add_subdirectory(errorHandling) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(clock) endif() diff --git a/projects/hip/tests/catch/unit/compiler/CMakeLists.txt b/projects/hip/tests/catch/unit/compiler/CMakeLists.txt new file mode 100644 index 0000000000..67eb7afcf5 --- /dev/null +++ b/projects/hip/tests/catch/unit/compiler/CMakeLists.txt @@ -0,0 +1,8 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipClassKernel.cc +) + +hip_add_exe_to_target(NAME CompilerTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) diff --git a/projects/hip/tests/catch/unit/compiler/hipClassKernel.cc b/projects/hip/tests/catch/unit/compiler/hipClassKernel.cc new file mode 100644 index 0000000000..b205c6bffb --- /dev/null +++ b/projects/hip/tests/catch/unit/compiler/hipClassKernel.cc @@ -0,0 +1,220 @@ +/* +Copyright (c) 2015 - 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 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. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "hipClassKernel.h" + +__global__ void +ovrdClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testOvrD tobj1; + result_ecd[tid] = (tobj1.ovrdFunc1() == 30); +} + +__global__ void +ovldClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFuncOvld tfo1; + result_ecd[tid] = (tfo1.func1(10) == 20) + && (tfo1.func1(10,10) == 30); +} + +TEST_CASE("Unit_hipClassKernel_Overload_Override") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovrdClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); + + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovldClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); +} + +// check for friend +__global__ void +friendClassKernel(bool* result_ecd){ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFrndB tfb1; + result_ecd[tid] = (tfb1.showA() == 10); +} + +TEST_CASE("Unit_hipClassKernel_Friend") { + bool *result_ecd; + result_ecd = AllocateDeviceMemory(); + hipLaunchKernelGGL(friendClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); +} + +// check sizeof empty class is 1 +__global__ void +emptyClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testClassEmpty ob1,ob2; + result_ecd[tid] = (sizeof(testClassEmpty) == 1) + && (&ob1 != &ob2); +} + +TEST_CASE("Unit_hipClassKernel_Empty") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(emptyClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); +} + +// tests for classes >8 bytes +__global__ void + sizeClassBKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeB) == 12) + && (sizeof(testSizeC) == 16) + && (sizeof(testSizeP1) == 6) + && (sizeof(testSizeP2) == 13) + && (sizeof(testSizeP3) == 8); +} + +TEST_CASE("Unit_hipClassKernel_BSize") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassBKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); +} + +__global__ void +sizeClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeA) == 16) + && (sizeof(testSizeDerived) == 24) + && (sizeof(testSizeDerived2) == 20); +} + +TEST_CASE("Unit_hipClassKernel_Size") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); +} + +__global__ void + sizeVirtualClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeDV) == 16) + && (sizeof(testSizeDerivedDV) == 16) + && (sizeof(testSizeVirtDerPack) == 24) + && (sizeof(testSizeVirtDer) == 24) + && (sizeof(testSizeDerMulti) == 48) ; + } + +TEST_CASE("Unit_hipClassKernel_Virtual") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeVirtualClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); +} + +// check pass by value +__global__ void +passByValueKernel(testPassByValue obj, bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (obj.exI == 10) + && (obj.exC == 'C'); +} + +TEST_CASE("Unit_hipClassKernel_Value") { + bool *result_ecd,*result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + testPassByValue exObj; + exObj.exI = 10; + exObj.exC = 'C'; + hipLaunchKernelGGL(passByValueKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + exObj, + result_ecd); + + VerifyResult(result_ech,result_ecd); + FreeMem(result_ech,result_ecd); +} \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/compiler/hipClassKernel.h b/projects/hip/tests/catch/unit/compiler/hipClassKernel.h new file mode 100644 index 0000000000..6da07f61d6 --- /dev/null +++ b/projects/hip/tests/catch/unit/compiler/hipClassKernel.h @@ -0,0 +1,234 @@ +/* +Copyright (c) 2015 - 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 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. +*/ +#ifndef _COMPILER_HIPCLASSKERNEL_H_ +#define _COMPILER_HIPCLASSKERNEL_H_ + +#include + +static const int BLOCKS = 512; +static const int THREADS_PER_BLOCK = 1; +size_t NBOOL = BLOCKS * sizeof(bool); + +class testFuncOvld{ + public: + int __host__ __device__ func1(int a){ + return a + 10; + } + + int __host__ __device__ func1(int a , int b){ + return a + b + 10; + } + +}; + +class testOvrB{ + public: + int __host__ __device__ ovrdFunc1(){ + return 10; + } + + +}; + +class testOvrD: public testOvrB{ + public: + int __host__ __device__ ovrdFunc1(){ + int x = testOvrB::ovrdFunc1(); + return x + 20; + } + +}; + +class testFrndA{ + private: + int fa = 10; + public: + friend class testFrndB; +}; + +class testFrndB{ + public: + __host__ __device__ int showA(){ + testFrndA x; + return x.fa; + } +}; + +class testClassEmpty {}; + +class testPassByValue{ + public: + int exI; + char exC; +}; + +class testSizeA { + public: + float xa; + int ia; + double da; + static char ca; +}; + +class testSizeDerived : testSizeA { + public: + float fd; +}; + +#pragma pack(push,4) +class testSizeDerived2 : testSizeA { + public: + float fd; +}; +#pragma pack(pop) + +class testSizeB { + public: + char ab; + int ib; + char cb; +}; + +class testSizeVirtDer : public virtual testSizeB { + public: + int ivd; +}; + +class testSizeVirtDer1 : public virtual testSizeB { + public: + int ivd1; +}; + +class testSizeDerMulti : public testSizeVirtDer, public testSizeVirtDer1 { + public: + int ivd2; +}; + +#pragma pack(push,4) +class testSizeVirtDerPack : public virtual testSizeB { + public: + int ivd; +}; +#pragma pack(pop) + +class testSizeC { + public: + char ac; + int ic; + int bc[2]; +}; + +class testSizeDV { + public: + virtual void __host__ __device__ func1(); + private: + int iDV; + +}; + +class testSizeDerivedDV : testSizeDV { + public: + virtual void __host__ __device__ funcD1(); + private: + int iDDV; +}; + +#pragma pack(push, 1) +class testSizeP1 { + public: + char ap; + int ip; + char cp; +}; +#pragma pack(pop) + +#pragma pack(push, 1) +class testSizeP2 { + public: + char ap1; + int ip1; + int bp1[2]; +}; +#pragma pack(pop) + +#pragma pack(push, 2) +class testSizeP3 { + public: + char ap2; + int ip2; + char cp2; +}; +#pragma pack(pop) + +class testDeviceClass { + public: + int iVar; + __host__ __device__ testDeviceClass(); + __host__ __device__ testDeviceClass(int a); + __host__ __device__ ~testDeviceClass(); +}; + +__host__ __device__ +testDeviceClass::testDeviceClass() { + iVar = 5; +} + +__host__ __device__ +testDeviceClass::testDeviceClass(int a) { + iVar = a; +} + +bool* AllocateHostMemory(void){ + bool *result_ech; + HIPCHECK(hipHostMalloc(&result_ech, + NBOOL, + hipHostMallocDefault)); + return result_ech; +} + +bool* AllocateDeviceMemory(void){ + bool* result_ecd; + HIPCHECK(hipMalloc(&result_ecd, + NBOOL)); + HIPCHECK(hipMemset(result_ecd, + false, + NBOOL)); + return result_ecd; +} + +void VerifyResult(bool* result_ech, bool* result_ecd){ + HIPCHECK(hipMemcpy(result_ech, + result_ecd, + BLOCKS*sizeof(bool), + hipMemcpyDeviceToHost)); + // validation on host side + for (int i = 0; i < BLOCKS; i++) { + HIPASSERT(result_ech[i] == true); + } +} + +void FreeMem(bool* result_ech, bool* result_ecd){ + HIPCHECK(hipHostFree(result_ech)); + HIPCHECK(hipFree(result_ecd)); +} + +#endif // _HIPCLASSKERNEL_H_ \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/device/CMakeLists.txt b/projects/hip/tests/catch/unit/device/CMakeLists.txt index 4dc6c79021..2bc7ea2498 100644 --- a/projects/hip/tests/catch/unit/device/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/device/CMakeLists.txt @@ -17,8 +17,10 @@ set(TEST_SRC hipSetGetDevice.cc hipDeviceGetUuid.cc hipDeviceGetP2PAttribute.cc + hipDeviceGetDefaultMemPool.cc hipExtGetLinkTypeAndHopCount.cc hipDeviceSetLimit.cc + hipDeviceSetGetMemPool.cc ) set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17) diff --git a/projects/hip/tests/catch/unit/device/hipDeviceGetDefaultMemPool.cc b/projects/hip/tests/catch/unit/device/hipDeviceGetDefaultMemPool.cc new file mode 100644 index 0000000000..822c2ab81c --- /dev/null +++ b/projects/hip/tests/catch/unit/device/hipDeviceGetDefaultMemPool.cc @@ -0,0 +1,57 @@ +/* +Copyright (c) 2022 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 +#include + +TEST_CASE("Unit_hipDeviceGetDefaultMemPool_Positive_Basic") { + const int device = GENERATE(range(0, HipTest::getDeviceCount())); + + int mem_pool_support = 0; + HIP_CHECK( + hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Test only runs on devices with memory pool support"); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, device)); + REQUIRE(mem_pool != nullptr); +} + +TEST_CASE("Unit_hipDeviceGetDefaultMemPool_Negative_Parameters") { + hipMemPool_t mem_pool; + + SECTION("mem_pool == nullptr") { + HIP_CHECK_ERROR(hipDeviceGetDefaultMemPool(nullptr, 0), hipErrorInvalidValue); + } + + SECTION("device < 0") { + HIP_CHECK_ERROR(hipDeviceGetDefaultMemPool(&mem_pool, -1), hipErrorInvalidDevice); + } + + SECTION("device ordinance too large") { + HIP_CHECK_ERROR(hipDeviceGetDefaultMemPool(&mem_pool, HipTest::getDeviceCount()), + hipErrorInvalidDevice); + } +} \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc b/projects/hip/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc index 01dfdc3d63..b093336d6a 100644 --- a/projects/hip/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc +++ b/projects/hip/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc @@ -1,45 +1,108 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 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 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 -/** - * hipDeviceGetCacheConfig tests - * Scenario1: Validates if pConfig = nullptr returns hip error code. - * Scenario2: Validates if the value returned by hipDeviceGetCacheConfig is valid. - */ -TEST_CASE("Unit_hipDeviceGetCacheConfig_NegTst") { - // Scenario1 - REQUIRE_FALSE(hipSuccess == hipDeviceGetCacheConfig(nullptr)); -} -TEST_CASE("Unit_hipDeviceGetCacheConfig_FuncTst") { - hipFuncCache_t cacheConfig; - // Scenario2 - HIP_CHECK(hipDeviceGetCacheConfig(&cacheConfig)); - REQUIRE_FALSE(((cacheConfig != hipFuncCachePreferNone) && - (cacheConfig != hipFuncCachePreferShared) && - (cacheConfig != hipFuncCachePreferL1) && - (cacheConfig != hipFuncCachePreferEqual))); - // This code exists to test the dummy implementation of - // hipDeviceSetCacheConfig. -#if HT_NVIDIA - HIP_CHECK(hipDeviceSetCacheConfig(cacheConfig)); -#else - REQUIRE_FALSE(hipSuccess == hipDeviceSetCacheConfig(cacheConfig)); +#include + +#include +#include + +namespace { +constexpr std::array kCacheConfigs{ + hipFuncCachePreferNone, hipFuncCachePreferShared, hipFuncCachePreferL1, + hipFuncCachePreferEqual}; +} // anonymous namespace + + +TEST_CASE("Unit_hipDeviceSetCacheConfig_Positive_Basic") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + INFO("Current device is: " << device); + + const auto cache_config = + GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs))); +#if HT_AMD + HIP_CHECK_ERROR(hipDeviceSetCacheConfig(cache_config), hipErrorNotSupported); +#elif HT_NVIDIA + HIP_CHECK(hipDeviceSetCacheConfig(cache_config)); #endif } + +TEST_CASE("Unit_hipDeviceSetCacheConfig_Negative_Parameters") { +#if HT_AMD + HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast(-1)), hipErrorNotSupported); +#elif HT_NVIDIA + HIP_CHECK_ERROR(hipDeviceSetCacheConfig(static_cast(-1)), hipErrorInvalidValue); +#endif +} + +TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Default") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + INFO("Current device is: " << device); + + hipFuncCache_t cache_config; + HIP_CHECK(hipDeviceGetCacheConfig(&cache_config)); + REQUIRE(cache_config == hipFuncCachePreferNone); +} + +TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Basic") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + INFO("Current device is: " << device); + + const auto cache_config = + GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs))); + + HIP_CHECK(hipDeviceSetCacheConfig(cache_config)); + hipFuncCache_t returned_cache_config; + HIP_CHECK(hipDeviceGetCacheConfig(&returned_cache_config)); + + REQUIRE(returned_cache_config == cache_config); +} + +TEST_CASE("Unit_hipDeviceGetCacheConfig_Positive_Threaded") { + class HipDeviceSetGetCacheConfigTest : public ThreadedZigZagTest { + public: + HipDeviceSetGetCacheConfigTest(const hipFuncCache_t cache_config) + : cache_config_{cache_config} {} + + void TestPart2() { HIP_CHECK_THREAD(hipDeviceSetCacheConfig(cache_config_)); } + + void TestPart3() { + hipFuncCache_t returned_cache_config; + HIP_CHECK(hipDeviceGetCacheConfig(&returned_cache_config)); + REQUIRE(returned_cache_config == cache_config_); + } + + private: + const hipFuncCache_t cache_config_; + }; + + const auto cache_config = + GENERATE(from_range(std::begin(kCacheConfigs), std::end(kCacheConfigs))); + + HipDeviceSetGetCacheConfigTest test(cache_config); + test.run(); +} + +TEST_CASE("Unit_HipDeviceGetCacheConfig_Negative_Parameters") { + HIP_CHECK_ERROR(hipDeviceGetCacheConfig(nullptr), hipErrorInvalidValue); +} \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/device/hipDeviceSetGetMemPool.cc b/projects/hip/tests/catch/unit/device/hipDeviceSetGetMemPool.cc new file mode 100644 index 0000000000..089c8003ac --- /dev/null +++ b/projects/hip/tests/catch/unit/device/hipDeviceSetGetMemPool.cc @@ -0,0 +1,159 @@ +/* +Copyright (c) 2022 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 +#include +#include + +static inline bool CheckMemPoolSupport(const int device) { + int mem_pool_support = 0; + HIP_CHECK( + hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Test only runs on devices with memory pool support"); + return false; + } + return true; +} + +static inline hipMemPool_t CreateMemPool(const int device) { + hipMemPoolProps kPoolProps; + kPoolProps.allocType = hipMemAllocationTypePinned; + kPoolProps.handleTypes = hipMemHandleTypeNone; + kPoolProps.location.type = hipMemLocationTypeDevice; + kPoolProps.location.id = device; + kPoolProps.win32SecurityAttributes = nullptr; + memset(kPoolProps.reserved, 0, sizeof(kPoolProps.reserved)); + + hipMemPool_t mem_pool; + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + + return mem_pool; +} + +TEST_CASE("Unit_hipDeviceSetMemPool_Positive_Basic") { + const int device = GENERATE(range(0, HipTest::getDeviceCount())); + + if (!CheckMemPoolSupport(device)) { + return; + } + + hipMemPool_t mem_pool = CreateMemPool(device); + HIP_CHECK(hipDeviceSetMemPool(device, mem_pool)); + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); +} + +TEST_CASE("Unit_hipDeviceSetMemPool_Negative_Parameters") { + hipMemPool_t mem_pool; + HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0)); + + SECTION("mem_pool == nullptr") { + HIP_CHECK_ERROR(hipDeviceSetMemPool(0, nullptr), hipErrorInvalidValue); + } + + SECTION("device < 0") { + HIP_CHECK_ERROR(hipDeviceSetMemPool(-1, mem_pool), hipErrorInvalidValue); + } + + SECTION("device ordinance too large") { + HIP_CHECK_ERROR(hipDeviceSetMemPool(HipTest::getDeviceCount(), mem_pool), hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipDeviceGetMemPool_Positive_Default") { + const int device = GENERATE(range(0, HipTest::getDeviceCount())); + + if (!CheckMemPoolSupport(device)) { + return; + } + + hipMemPool_t default_mem_pool; + HIP_CHECK(hipDeviceGetDefaultMemPool(&default_mem_pool, device)); + + hipMemPool_t mem_pool; + HIP_CHECK(hipDeviceGetMemPool(&mem_pool, device)); + + REQUIRE(mem_pool == default_mem_pool); +} + +TEST_CASE("Unit_hipDeviceGetMemPool_Positive_Basic") { + const int device = GENERATE(range(0, HipTest::getDeviceCount())); + + if (!CheckMemPoolSupport(device)) { + return; + } + + hipMemPool_t mem_pool = CreateMemPool(device); + HIP_CHECK(hipDeviceSetMemPool(device, mem_pool)); + + hipMemPool_t returned_mem_pool; + HIP_CHECK(hipDeviceGetMemPool(&returned_mem_pool, device)); + + REQUIRE(returned_mem_pool == mem_pool); + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); +} + +TEST_CASE("Unit_hipDeviceGetMemPool_Positive_Threaded") { + class HipDeviceGetMemPoolTest : public ThreadedZigZagTest { + public: + void TestPart2() { + mem_pool_ = CreateMemPool(0); + HIP_CHECK_THREAD(hipDeviceSetMemPool(0, mem_pool_)); + } + void TestPart3() { + hipMemPool_t returned_mem_pool; + HIP_CHECK(hipDeviceGetMemPool(&returned_mem_pool, 0)); + + REQUIRE(returned_mem_pool == mem_pool_); + } + void TestPart4() { HIP_CHECK_THREAD(hipMemPoolDestroy(mem_pool_)); } + + private: + hipMemPool_t mem_pool_; + }; + + if (!CheckMemPoolSupport(0)) { + return; + } + + HipDeviceGetMemPoolTest test; + test.run(); +} + +TEST_CASE("Unit_hipDeviceGetMemPool_Negative_Parameters") { + hipMemPool_t mem_pool; + + SECTION("mem_pool == nullptr") { + HIP_CHECK_ERROR(hipDeviceGetMemPool(nullptr, 0), hipErrorInvalidValue); + } + + SECTION("device < 0") { + HIP_CHECK_ERROR(hipDeviceGetMemPool(&mem_pool, -1), hipErrorInvalidValue); + } + + SECTION("device ordinance too large") { + HIP_CHECK_ERROR(hipDeviceGetMemPool(&mem_pool, HipTest::getDeviceCount()), + hipErrorInvalidValue); + } +} diff --git a/projects/hip/tests/catch/unit/device/hipDeviceSynchronize.cc b/projects/hip/tests/catch/unit/device/hipDeviceSynchronize.cc index 0b54463159..bfb02ab784 100644 --- a/projects/hip/tests/catch/unit/device/hipDeviceSynchronize.cc +++ b/projects/hip/tests/catch/unit/device/hipDeviceSynchronize.cc @@ -1,19 +1,22 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 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 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. */ @@ -41,6 +44,36 @@ static __global__ void Iter(int* Ad, int num) { } } +TEST_CASE("Unit_hipDeviceSynchronize_Positive_Empty_Streams") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + INFO("Current device: " << device); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + INFO("Current device: " << device); + + int *A_h = nullptr, *A_d = nullptr; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), _SIZE, hipHostMallocDefault)); + A_h[0] = 1; + HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), _SIZE)); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL)); + hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30); + HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL)); + + CHECK(1 << 30 != A_h[0] - 1); + HIP_CHECK(hipDeviceSynchronize()); + CHECK(1 << 30 == A_h[0] - 1); +} + TEST_CASE("Unit_hipDeviceSynchronize_Functional") { int* A[NUM_STREAMS]; int* Ad[NUM_STREAMS]; diff --git a/projects/hip/tests/catch/unit/errorHandling/CMakeLists.txt b/projects/hip/tests/catch/unit/errorHandling/CMakeLists.txt new file mode 100644 index 0000000000..78933fd3f8 --- /dev/null +++ b/projects/hip/tests/catch/unit/errorHandling/CMakeLists.txt @@ -0,0 +1,12 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipGetErrorName.cc + hipGetErrorString.cc + hipGetLastError.cc + hipPeekAtLastError.cc +) + +hip_add_exe_to_target(NAME ErrorHandlingTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++14) \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/errorHandling/errorEnumerators.h b/projects/hip/tests/catch/unit/errorHandling/errorEnumerators.h new file mode 100644 index 0000000000..a379a47a06 --- /dev/null +++ b/projects/hip/tests/catch/unit/errorHandling/errorEnumerators.h @@ -0,0 +1,99 @@ +/* +Copyright (c) 2022 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. +*/ + +#pragma once + +#include + +constexpr hipError_t kErrorEnumerators[] = {hipSuccess, + hipErrorInvalidValue, + hipErrorOutOfMemory, + hipErrorNotInitialized, + hipErrorDeinitialized, + hipErrorProfilerDisabled, + hipErrorProfilerNotInitialized, + hipErrorProfilerAlreadyStarted, + hipErrorProfilerAlreadyStopped, + hipErrorInvalidConfiguration, + hipErrorInvalidPitchValue, + hipErrorInvalidSymbol, + hipErrorInvalidDevicePointer, + hipErrorInvalidMemcpyDirection, + hipErrorInsufficientDriver, + hipErrorMissingConfiguration, + hipErrorPriorLaunchFailure, + hipErrorInvalidDeviceFunction, + hipErrorNoDevice, + hipErrorInvalidDevice, + hipErrorInvalidImage, + hipErrorInvalidContext, + hipErrorContextAlreadyCurrent, + hipErrorMapFailed, + hipErrorUnmapFailed, + hipErrorArrayIsMapped, + hipErrorAlreadyMapped, + hipErrorNoBinaryForGpu, + hipErrorAlreadyAcquired, + hipErrorNotMapped, + hipErrorNotMappedAsArray, + hipErrorNotMappedAsPointer, + hipErrorECCNotCorrectable, + hipErrorUnsupportedLimit, + hipErrorContextAlreadyInUse, + hipErrorPeerAccessUnsupported, + hipErrorInvalidKernelFile, + hipErrorInvalidGraphicsContext, + hipErrorInvalidSource, + hipErrorFileNotFound, + hipErrorSharedObjectSymbolNotFound, + hipErrorSharedObjectInitFailed, + hipErrorOperatingSystem, + hipErrorInvalidHandle, + hipErrorIllegalState, + hipErrorNotFound, + hipErrorNotReady, + hipErrorIllegalAddress, + hipErrorLaunchOutOfResources, + hipErrorLaunchTimeOut, + hipErrorPeerAccessAlreadyEnabled, + hipErrorPeerAccessNotEnabled, + hipErrorSetOnActiveProcess, + hipErrorContextIsDestroyed, + hipErrorAssert, + hipErrorHostMemoryAlreadyRegistered, + hipErrorHostMemoryNotRegistered, + hipErrorLaunchFailure, + hipErrorCooperativeLaunchTooLarge, + hipErrorNotSupported, + hipErrorStreamCaptureUnsupported, + hipErrorStreamCaptureInvalidated, + hipErrorStreamCaptureMerge, + hipErrorStreamCaptureUnmatched, + hipErrorStreamCaptureUnjoined, + hipErrorStreamCaptureIsolation, + hipErrorStreamCaptureImplicit, + hipErrorCapturedEvent, + hipErrorStreamCaptureWrongThread, + hipErrorGraphExecUpdateFailure, + hipErrorUnknown, + hipErrorRuntimeMemory, + hipErrorRuntimeOther}; \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/errorHandling/hipGetErrorName.cc b/projects/hip/tests/catch/unit/errorHandling/hipGetErrorName.cc new file mode 100644 index 0000000000..f8c7f624be --- /dev/null +++ b/projects/hip/tests/catch/unit/errorHandling/hipGetErrorName.cc @@ -0,0 +1,47 @@ +/* +Copyright (c) 2022 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 "errorEnumerators.h" +#include +#include + +TEST_CASE("Unit_hipGetErrorName_Positive_Basic") { + const char* error_string = nullptr; + const auto enumerator = + GENERATE(from_range(std::begin(kErrorEnumerators), std::end(kErrorEnumerators))); + + error_string = hipGetErrorName(enumerator); + + REQUIRE(error_string != nullptr); + REQUIRE(strlen(error_string) > 0); +} + +TEST_CASE("Unit_hipGetErrorName_Negative_Parameters") { + const char* error_string = hipGetErrorName(static_cast(-1)); + REQUIRE(error_string != nullptr); +#if HT_NVIDIA + REQUIRE_THAT(error_string, Catch::Equals("cudaErrorUnknown")); +#elif HT_AMD + REQUIRE_THAT(error_string, Catch::Equals("hipErrorUnknown")); +#endif +} \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/errorHandling/hipGetErrorString.cc b/projects/hip/tests/catch/unit/errorHandling/hipGetErrorString.cc new file mode 100644 index 0000000000..fe3f5523c2 --- /dev/null +++ b/projects/hip/tests/catch/unit/errorHandling/hipGetErrorString.cc @@ -0,0 +1,41 @@ +/* +Copyright (c) 2022 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 "errorEnumerators.h" +#include +#include + +TEST_CASE("Unit_hipGetErrorString_Positive_Basic") { + const char* error_string = nullptr; + const auto enumerator = + GENERATE(from_range(std::begin(kErrorEnumerators), std::end(kErrorEnumerators))); + + error_string = hipGetErrorString(enumerator); + + REQUIRE(error_string != nullptr); + REQUIRE(strlen(error_string) > 0); +} + +TEST_CASE("Unit_hipGetErrorString_Negative_Parameters") { + const char* error_string = hipGetErrorString(static_cast(-1)); + REQUIRE(error_string != nullptr); +} \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/errorHandling/hipGetLastError.cc b/projects/hip/tests/catch/unit/errorHandling/hipGetLastError.cc new file mode 100644 index 0000000000..f2b8f1cfd3 --- /dev/null +++ b/projects/hip/tests/catch/unit/errorHandling/hipGetLastError.cc @@ -0,0 +1,44 @@ +/* +Copyright (c) 2022 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 +#include +#include + +TEST_CASE("Unit_hipGetLastError_Positive_Basic") { + HIP_CHECK(hipGetLastError()); + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipGetLastError()); +} + +TEST_CASE("Unit_hipGetLastError_Positive_Threaded") { + class HipGetLastErrorThreadedTest : public ThreadedZigZagTest { + public: + void TestPart2() { REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); } + void TestPart3() { HIP_CHECK(hipGetLastError()); } + void TestPart4() { REQUIRE_THREAD(hipGetLastError() == hipErrorInvalidValue); } + }; + + HipGetLastErrorThreadedTest test; + test.run(); +} diff --git a/projects/hip/tests/catch/unit/errorHandling/hipPeekAtLastError.cc b/projects/hip/tests/catch/unit/errorHandling/hipPeekAtLastError.cc new file mode 100644 index 0000000000..ce7c605719 --- /dev/null +++ b/projects/hip/tests/catch/unit/errorHandling/hipPeekAtLastError.cc @@ -0,0 +1,48 @@ +/* +Copyright (c) 2022 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 +#include +#include + +TEST_CASE("Unit_hipPeekAtLastError_Positive_Basic") { + HIP_CHECK(hipPeekAtLastError()); + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipPeekAtLastError()); +} + +TEST_CASE("Unit_hipPeekAtLastError_Positive_Threaded") { + class HipPeekAtLastErrorTest : public ThreadedZigZagTest { + public: + void TestPart2() { REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); } + void TestPart3() { + HIP_CHECK(hipPeekAtLastError()); + HIP_CHECK(hipGetLastError()); + } + void TestPart4() { REQUIRE_THREAD(hipPeekAtLastError() == hipErrorInvalidValue); } + }; + + HipPeekAtLastErrorTest test; + test.run(); +} diff --git a/projects/hip/tests/catch/unit/event/CMakeLists.txt b/projects/hip/tests/catch/unit/event/CMakeLists.txt index b9f57c1451..3054e85670 100644 --- a/projects/hip/tests/catch/unit/event/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/event/CMakeLists.txt @@ -6,6 +6,7 @@ set(TEST_SRC Unit_hipEventRecord.cc Unit_hipEventIpc.cc hipEventDestroy.cc + hipEventCreate.cc hipEventCreateWithFlags.cc hipEventSynchronize.cc ) diff --git a/projects/hip/tests/catch/unit/event/hipEventCreate.cc b/projects/hip/tests/catch/unit/event/hipEventCreate.cc new file mode 100644 index 0000000000..30cad62c77 --- /dev/null +++ b/projects/hip/tests/catch/unit/event/hipEventCreate.cc @@ -0,0 +1,38 @@ +/* +Copyright (c) 2022 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. +*/ +/* +Testcase Scenarios : +Unit_hipEventCreate_Positive - Test simple event creation with hipEventCreate api +*/ + +#include + +TEST_CASE("Unit_hipEventCreate_Positive") { + int id = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(id)); + + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + REQUIRE(event != nullptr); + + HIP_CHECK(hipEventDestroy(event)); +} diff --git a/projects/hip/tests/catch/unit/event/hipEventCreateWithFlags.cc b/projects/hip/tests/catch/unit/event/hipEventCreateWithFlags.cc index 2368fcad13..7f31c8e347 100644 --- a/projects/hip/tests/catch/unit/event/hipEventCreateWithFlags.cc +++ b/projects/hip/tests/catch/unit/event/hipEventCreateWithFlags.cc @@ -29,8 +29,7 @@ Unit_hipEventCreateWithFlags_Positive - Test simple event creation with hipEvent TEST_CASE("Unit_hipEventCreateWithFlags_Positive") { #if HT_AMD - // On AMD platform, hipEventInterprocess support is under development. Use of this flag will return an error. Omitted - const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventReleaseToDevice, hipEventReleaseToSystem); + const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming, hipEventReleaseToDevice, hipEventReleaseToSystem); #else // On Non-AMD platforms hipEventReleaseToDevice / hipEventReleaseToSystem are not defined const unsigned int flagUnderTest = GENERATE(hipEventDefault, hipEventBlockingSync, hipEventDisableTiming, hipEventInterprocess | hipEventDisableTiming); diff --git a/projects/hip/tests/catch/unit/graph/CMakeLists.txt b/projects/hip/tests/catch/unit/graph/CMakeLists.txt index 5af1423bdc..7f15b7e6ff 100644 --- a/projects/hip/tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/graph/CMakeLists.txt @@ -23,6 +23,7 @@ set(TEST_SRC hipGraphAddEmptyNode.cc hipGraphAddDependencies.cc hipGraphAddEventRecordNode.cc + hipGraphAddEventWaitNode.cc hipGraph.cc hipSimpleGraphWithKernel.cc hipGraphAddMemcpyNode.cc @@ -78,6 +79,7 @@ set(TEST_SRC hipGraphNodeGetDependencies.cc hipGraphHostNodeGetParams.cc hipGraphExecChildGraphNodeSetParams.cc + hipStreamGetCaptureInfo_v2.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/projects/hip/tests/catch/unit/graph/hipGraphAddEventWaitNode.cc b/projects/hip/tests/catch/unit/graph/hipGraphAddEventWaitNode.cc new file mode 100644 index 0000000000..a196ac7428 --- /dev/null +++ b/projects/hip/tests/catch/unit/graph/hipGraphAddEventWaitNode.cc @@ -0,0 +1,303 @@ +/* +Copyright (c) 2022 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, INCLUDING 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 ANY 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) Simple Scenario: Create an event record node and then create an event +wait node using the same event and add it to graph. Instantiate and Launch +the Graph. Wait for the graph to complete. The operation must succeed without +any failures. + 2) Create a graph 1 with memcpyh2d, event record node (event A), kernel1 +and memcpyd2h nodes. Create a graph 2 with Event Wait (event A) , kernel2 +and memcpyd2h nodes. Instantiate and launch graph1 on stream1 and graph2 on +stream2. Wait for both graph1 and graph2 to complete. Validate the result of +both graphs. + 3) Execute graph1 and graph2 in scenario 2 multiple times in a loop +(100 times). + 4) Execute scenario 2 with stream1 = stream2. + 5) Repeat scenario 2 for different event flags. + 6) Negative Scenarios + - Pass input node parameter as nullptr. + - Pass input graph parameter as nullptr. + - Pass input dependency parameter as nullptr. + - Pass input event parameter as nullptr. + - Pass uninitialized input graph parameter. + - Pass uninitialized input event parameter. +*/ + +#include +#include +#include + +#define LEN 512 + +/** + * Scenario 1 + */ +TEST_CASE("Unit_hipGraphAddEventWaitNode_Functional_Simple") { + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + hipGraphNode_t event_rec_node, event_wait_node; + // Create a event record node in graph + HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph, nullptr, 0, + event)); + // Create a event wait node in graph + HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0, + event)); + HIP_CHECK(hipGraphAddDependencies(graph, &event_rec_node, + &event_wait_node, 1)); + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Local Function + */ +static void validate_hipGraphAddEventWaitNode_internodedep(int test, + int nstep, unsigned flag = hipEventDefault) { + size_t memsize = LEN * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, LEN); + size_t NElem{LEN}; + hipGraph_t graph1, graph2; + hipStream_t streamForGraph1, streamForGraph2; + hipGraphExec_t graphExec1, graphExec2; + HIP_CHECK(hipStreamCreate(&streamForGraph1)); + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphCreate(&graph2, 0)); + if (0 == test) { + HIP_CHECK(hipStreamCreate(&streamForGraph2)); + } else if (1 == test) { + streamForGraph2 = streamForGraph1; + } + hipEvent_t event1; + HIP_CHECK(hipEventCreateWithFlags(&event1, flag)); + hipGraphNode_t event_rec_node, event_wait_node; + int *inp_h, *inp_d, *out_h_g1, *out_d_g1, *out_h_g2, *out_d_g2; + // Allocate host buffers + inp_h = reinterpret_cast(malloc(memsize)); + REQUIRE(inp_h != nullptr); + out_h_g1 = reinterpret_cast(malloc(memsize)); + REQUIRE(out_h_g1 != nullptr); + out_h_g2 = reinterpret_cast(malloc(memsize)); + REQUIRE(out_h_g2 != nullptr); + // Allocate device buffers + HIP_CHECK(hipMalloc(&inp_d, memsize)); + HIP_CHECK(hipMalloc(&out_d_g1, memsize)); + HIP_CHECK(hipMalloc(&out_d_g2, memsize)); + // Initialize host buffer + for (uint32_t i = 0; i < LEN; i++) { + inp_h[i] = i; + out_h_g1[i] = 0; + out_h_g2[i] = 0; + } + // Graph1 creation ........... + // Create event1 record node in graph1 + HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph1, nullptr, 0, + event1)); + + // Create memcpy and kernel nodes for graph1 + hipGraphNode_t memcpyH2D, memcpyD2H_1, kernelnode_1; + hipKernelNodeParams kernelNodeParams1{}; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d, + inp_h, memsize, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph1, nullptr, 0, + out_h_g1, out_d_g1, memsize, hipMemcpyDeviceToHost)); + + void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast(&NElem)}; + kernelNodeParams1.func = + reinterpret_cast(HipTest::vector_square); + kernelNodeParams1.gridDim = dim3(blocks); + kernelNodeParams1.blockDim = dim3(threadsPerBlock); + kernelNodeParams1.sharedMemBytes = 0; + kernelNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph1, nullptr, 0, + &kernelNodeParams1)); + // Create dependencies for graph1 + HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D, + &event_rec_node, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &event_rec_node, + &kernelnode_1, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode_1, + &memcpyD2H_1, 1)); + + // Graph2 creation ........... + // Create event1 record node in graph2 + HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0, + event1)); + + // Create memcpy and kernel nodes for graph2 + hipGraphNode_t memcpyD2H_2, kernelnode_2; + hipKernelNodeParams kernelNodeParams2{}; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph2, nullptr, 0, + out_h_g2, out_d_g2, memsize, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast(&NElem)}; + kernelNodeParams2.func = + reinterpret_cast(HipTest::vector_cubic); + kernelNodeParams2.gridDim = dim3(blocks); + kernelNodeParams2.blockDim = dim3(threadsPerBlock); + kernelNodeParams2.sharedMemBytes = 0; + kernelNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams2.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph2, nullptr, 0, + &kernelNodeParams2)); + // Create dependencies for graph2 + HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node, + &kernelnode_2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph2, &kernelnode_2, + &memcpyD2H_2, 1)); + + // Instantiate and launch the graphs + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + for (int istep = 0; istep < nstep; istep++) { + HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1)); + HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2)); + HIP_CHECK(hipStreamSynchronize(streamForGraph1)); + HIP_CHECK(hipStreamSynchronize(streamForGraph2)); + // Validate output + bool btestPassed1 = true; + for (uint32_t i = 0; i < LEN; i++) { + if (out_h_g1[i] != (inp_h[i]*inp_h[i])) { + btestPassed1 = false; + break; + } + } + REQUIRE(btestPassed1 == true); + bool btestPassed2 = true; + for (uint32_t i = 0; i < LEN; i++) { + if (out_h_g2[i] != (inp_h[i]*inp_h[i]*inp_h[i])) { + btestPassed2 = false; + break; + } + } + REQUIRE(btestPassed2 == true); + } + // Destroy all resources + HIP_CHECK(hipFree(inp_d)); + HIP_CHECK(hipFree(out_d_g1)); + HIP_CHECK(hipFree(out_d_g2)); + free(inp_h); + free(out_h_g1); + free(out_h_g2); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipStreamDestroy(streamForGraph1)); + if (0 == test) { + HIP_CHECK(hipStreamDestroy(streamForGraph2)); + } +} + +/** + * Scenario 2 + */ +TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphMultStrmDependency") { + validate_hipGraphAddEventWaitNode_internodedep(0, 1); +} + +/** + * Scenario 3 + */ +TEST_CASE("Unit_hipGraphAddEventWaitNode_MultipleRun") { + validate_hipGraphAddEventWaitNode_internodedep(0, 100); +} + +/** + * Scenario 4 + */ +TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphOneStrmDependency") { + validate_hipGraphAddEventWaitNode_internodedep(1, 1); +} + +/** + * Scenario 5 + */ +TEST_CASE("Unit_hipGraphAddEventWaitNode_differentFlags") { + SECTION("flag = hipEventBlockingSync") { + validate_hipGraphAddEventWaitNode_internodedep(0, 1, + hipEventBlockingSync); + } + SECTION("graph = hipEventDisableTiming") { + validate_hipGraphAddEventWaitNode_internodedep(0, 1, + hipEventDisableTiming); + } +} + +/** + * Scenario 6 + */ +TEST_CASE("Unit_hipGraphAddEventWaitNode_Negative") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + hipGraphNode_t eventwait; + + SECTION("pGraphNode = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(nullptr, + graph, nullptr, 0, event)); + } + + SECTION("graph = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait, + nullptr, nullptr, 0, event)); + } + + SECTION("pDependencies = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait, + graph, nullptr, 1, event)); + } + + SECTION("event = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait, + graph, nullptr, 0, nullptr)); + } + + SECTION("graph is uninitialized") { + hipGraph_t graph_uninit{}; + REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait, + graph_uninit, nullptr, 0, event)); + } + + SECTION("event is uninitialized") { + hipEvent_t event_uninit{}; + REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait, + graph, nullptr, 0, event_uninit)); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); +} diff --git a/projects/hip/tests/catch/unit/graph/hipGraphGetNodes.cc b/projects/hip/tests/catch/unit/graph/hipGraphGetNodes.cc index 18b3df1722..93d86ac87f 100644 --- a/projects/hip/tests/catch/unit/graph/hipGraphGetNodes.cc +++ b/projects/hip/tests/catch/unit/graph/hipGraphGetNodes.cc @@ -25,6 +25,8 @@ Functional :: 2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph. 3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes will be set to NULL, and the number of nodes actually obtained will be returned in numNodes. +4) Begin stream capture and push operations to stream. Verify nodes of created graph are matching the +operations pushed. Argument Validation :: 1) Pass graph as nullptr and verify api returns error code. @@ -139,6 +141,110 @@ TEST_CASE("Unit_hipGraphGetNodes_Functional") { free(nodes); } +/** + * Begin stream capture and push operations to stream. + * Verify nodes of created graph are matching the operations pushed. + */ +TEST_CASE("Unit_hipGraphGetNodes_CapturedStream") { + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + size_t Nbytes = N * sizeof(float); + constexpr int numMemcpy{2}, numKernel{1}, numMemset{1}; + int cntMemcpy{}, cntKernel{}, cntMemset{}; + hipStream_t stream, streamForGraph; + hipGraphNodeType nodeType; + float *A_d, *C_d; + float *A_h, *C_h; + + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + REQUIRE(graph != nullptr); + + size_t numNodes{}; + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + INFO("Num of nodes returned by GetNodes : " << numNodes); + REQUIRE(numNodes == numMemcpy + numKernel + numMemset); + + int numBytes = sizeof(hipGraphNode_t) * numNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes)); + for (size_t i = 0; i < numNodes; i++) { + HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType)); + + switch (nodeType) { + case hipGraphNodeTypeMemcpy: + cntMemcpy++; + break; + + case hipGraphNodeTypeKernel: + cntKernel++; + break; + + case hipGraphNodeTypeMemset: + cntMemset++; + break; + + default: + INFO("Unexpected nodetype returned : " << nodeType); + REQUIRE(false); + } + } + + REQUIRE(cntMemcpy == numMemcpy); + REQUIRE(cntKernel == numKernel); + REQUIRE(cntMemset == numMemset); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + INFO("A and C not matching at " << i << " C_h[i] " << C_h[i] + << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} + /** * Test performs api parameter validation by passing various values * as input and output parameters and validates the behavior. diff --git a/projects/hip/tests/catch/unit/graph/hipGraphGetRootNodes.cc b/projects/hip/tests/catch/unit/graph/hipGraphGetRootNodes.cc index c2484fc8b5..2be8ea5110 100644 --- a/projects/hip/tests/catch/unit/graph/hipGraphGetRootNodes.cc +++ b/projects/hip/tests/catch/unit/graph/hipGraphGetRootNodes.cc @@ -26,6 +26,8 @@ Functional :: 2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph. 3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes. + 4) Create a graph with stream capture done on multiple dependent streams. + Verify root nodes of created graph are matching the operations pushed which doesn't have dependencies. Argument Validation :: 1) Pass graph as nullptr and verify api returns error code. @@ -50,6 +52,7 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { constexpr auto addlEntries = 5; hipGraph_t graph; + hipGraphNode_t memcpyNode, kernelNode; hipKernelNodeParams kernelNodeParams{}; hipStream_t streamForGraph; @@ -141,6 +144,112 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { free(rootnodes); } + +/** + * Create a graph with stream capture done on multiple dependent streams. Verify root nodes + * of created graph are matching the operations pushed which doesn't have dependencies. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_CapturedStream") { + hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr}; + hipStream_t streamForGraph{nullptr}; + hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + constexpr int numMemsetNodes = 2; + size_t Nbytes = N * sizeof(float), numRootNodes{}; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&mstream)); + HIP_CHECK(hipEventCreate(&memsetEvent1)); + HIP_CHECK(hipEventCreate(&memsetEvent2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, mstream)); + HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1)); + HIP_CHECK(hipEventRecord(memsetEvent1, stream1)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream2)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, mstream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + + // Verify numof root nodes + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == numMemsetNodes); + INFO("Num of nodes returned by GetRootNodes : " << numRootNodes); + + int numBytes = sizeof(hipGraphNode_t) * numRootNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numRootNodes)); + REQUIRE(numRootNodes == numMemsetNodes); + + // Verify root nodes returned are memset nodes. + HIP_CHECK(hipGraphNodeGetType(nodes[0], &nodeType)); + REQUIRE(nodeType == hipGraphNodeTypeMemset); + HIP_CHECK(hipGraphNodeGetType(nodes[1], &nodeType)); + REQUIRE(nodeType == hipGraphNodeTypeMemset); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + INFO("A and C not matching at " << i << " C_h[i] " << C_h[i] + << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipStreamDestroy(mstream)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memsetEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} + + /** * Test performs api parameter validation by passing various values * as input and output parameters and validates the behavior. diff --git a/projects/hip/tests/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc b/projects/hip/tests/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc new file mode 100644 index 0000000000..d19e059083 --- /dev/null +++ b/projects/hip/tests/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc @@ -0,0 +1,280 @@ +/* +Copyright (c) 2022 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, INCLUDING 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 ANY 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 +------------------ +Functional: +1) Start stream capture and get capture info v2. Verify api is success, capture status is hipStreamCaptureStatusActive, + identifier returned is valid/non-zero, graph object is returned. +2) When stream capture is in progress, create dependent nodes by creating multistream dependencies and verify the api returns + valid dependent nodes. +3) End stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone and + identifier/graph/nodes are not returned by api. +4) When optional parameters are not passed, make sure api still returns capture status of stream. +5) Begin capture on hipStreamPerThread, get capture info v2 and validate results. +6) Perform multiple captures and verify the identifier returned is unique. + +Parameter Validation/Negative: +1) Capture Status location as nullptr and verify api returns error code. +2) Stream as nullptr and verify api returns error code. + +*/ + +#include +#include +#include + +constexpr size_t N = 1000000; +constexpr int LAUNCH_ITERS = 1; + +/** + * Validates stream capture infov2, launches graph and verifies results + */ +void validateStreamCaptureInfoV2(hipStream_t mstream) { + hipStream_t stream1{nullptr}, stream2{nullptr}, streamForLaunch{nullptr}; + hipEvent_t memcpyEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}, capInfoGraph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + const hipGraphNode_t *nodelist{}; + size_t Nbytes = N * sizeof(float), numDependencies; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + HIP_CHECK(hipStreamCreate(&streamForLaunch)); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + // Create cross stream dependencies. + // memset/memcpy operations are done on stream1 and stream2 + // and they are joined back to mainstream + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipEventCreate(&memcpyEvent1)); + HIP_CHECK(hipEventCreate(&memsetEvent2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + + HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, mstream)); + HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1)); + HIP_CHECK(hipEventRecord(memcpyEvent1, stream1)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream2)); + HIP_CHECK(hipStreamWaitEvent(mstream, memcpyEvent1, 0)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0)); + + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + unsigned long long capSequenceID = 0; // NOLINT + constexpr int numDepsCreated = 2; // Num of dependencies created + + HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies)); + + // verify capture status is active, sequence id is valid, graph is returned, + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + REQUIRE(capSequenceID > 0); + REQUIRE(capInfoGraph != nullptr); + REQUIRE(numDependencies == numDepsCreated); + + // verify dependency nodes list returned is the one we created. + hipGraphNodeType type(hipGraphNodeTypeEmpty); + + HIP_CHECK(hipGraphNodeGetType(nodelist[0], &type)); + if ((type != hipGraphNodeTypeMemset) && (type != hipGraphNodeTypeMemcpy)) { + INFO("Type0 returned as " << type); + REQUIRE(false); + } + + HIP_CHECK(hipGraphNodeGetType(nodelist[1], &type)); + if ((type != hipGraphNodeTypeMemset) && (type != hipGraphNodeTypeMemcpy)) { + INFO("Type1 returned as " << type); + REQUIRE(false); + } + + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, mstream, A_d, C_d, N); + + // End capture and verify graph is returned + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + REQUIRE(graph != nullptr); + + // verify capture status is inactive and other params are not updated + capSequenceID = 0; + capInfoGraph = nullptr; + numDependencies = 0; + nodelist = nullptr; + HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies)); + REQUIRE(captureStatus == hipStreamCaptureStatusNone); + REQUIRE(capSequenceID == 0); + REQUIRE(capInfoGraph == nullptr); + REQUIRE(nodelist == nullptr); + REQUIRE(numDependencies == 0); + + // Verify api still returns capture status when optional args are not passed + HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus)); + REQUIRE(captureStatus == hipStreamCaptureStatusNone); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(graphExec != nullptr); + + // Replay the recorded sequence multiple times + for (int i = 0; i < LAUNCH_ITERS; i++) { + HIP_CHECK(hipGraphLaunch(graphExec, streamForLaunch)); + } + + HIP_CHECK(hipStreamSynchronize(streamForLaunch)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForLaunch)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memcpyEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + INFO("A and C not matching at " << i << " C_h[i] " << C_h[i] + << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + free(A_h); + free(C_h); +} + +/** + * Basic Functional Test for stream capture and getting capture info V2. + * Regular/custom stream is used for stream capture. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_BasicFunctional") { + hipStream_t streamForCapture; + + HIP_CHECK(hipStreamCreate(&streamForCapture)); + validateStreamCaptureInfoV2(streamForCapture); + HIP_CHECK(hipStreamDestroy(streamForCapture)); +} + +/** + * Test performs stream capture on hipStreamPerThread and validates + * capture info V2. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_hipStreamPerThread") { + validateStreamCaptureInfoV2(hipStreamPerThread); +} + +/** + * Test starts stream capture on multiple streams and verifies uniqueness of + * identifiers returned from capture Info V2. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_UniqueID") { + constexpr int numStreams = 100; + hipStream_t streams[numStreams]{}; + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + std::vector idlist; + unsigned long long capSequenceID{}; // NOLINT + hipGraph_t graph{nullptr}; + + for (int i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + HIP_CHECK(hipStreamBeginCapture(streams[i], hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamGetCaptureInfo_v2(streams[i], &captureStatus, + &capSequenceID, nullptr, nullptr, nullptr)); + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + REQUIRE(capSequenceID > 0); + idlist.push_back(capSequenceID); + } + + for (int i = 0; i < numStreams; i++) { + for (int j = i+1; j < numStreams; j++) { + if (idlist[i] == idlist[j]) { + INFO("Same identifier returned for stream " + << i << " and stream " << j); + REQUIRE(false); + } + } + } + + for (int i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamEndCapture(streams[i], &graph)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streams[i])); + } +} + +/** + * Parameter validation/Negative tests for api + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_ParamValidation") { + hipError_t ret; + hipStream_t stream; + float *A_d; + hipGraph_t graph{}, capInfoGraph{}; + hipStreamCaptureStatus captureStatus; + unsigned long long capSequenceID; // NOLINT + size_t numDependencies; + const hipGraphNode_t *nodelist{}; + constexpr int numBytes{100}; + + HIP_CHECK(hipMalloc(&A_d, numBytes)); + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemsetAsync(A_d, 0, numBytes, stream)); + + SECTION("Capture Status location as nullptr") { + ret = hipStreamGetCaptureInfo_v2(stream, nullptr, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("Stream as nullptr") { + ret = hipStreamGetCaptureInfo_v2(nullptr, &captureStatus, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies); + if ((ret != hipErrorUnknown) && (ret != hipErrorStreamCaptureImplicit)) { + INFO("Ret : " << ret); + REQUIRE(false); + } + } + + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc b/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc index 09ea32bda8..cef164e2ff 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc @@ -676,8 +676,14 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { int stat = 0; if (fork() == 0) { // The below part should be inside fork - int managed = HmmAttrPrint(); - if (managed == 1) { + int managedMem = 0, pageMemAccess = 0; + HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess); + + HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0)); + WARN("hipDeviceAttributeManagedMemory: " << managedMem); + if ((managedMem == 1) && (pageMemAccess == 1)) { int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123; // Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz)); Mllc = reinterpret_cast(aligned_alloc(4096, 4096*4)); @@ -706,19 +712,27 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { exit(10); // 10 for Pass result } } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + SUCCEED("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + exit(Catch::ResultDisposition::ContinueOnFailure); } } else { wait(&stat); int Result = WEXITSTATUS(stat); - if (Result != 10) { - REQUIRE(false); + if (Result == Catch::ResultDisposition::ContinueOnFailure) { + WARN("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } else { + if (Result != 10) { + REQUIRE(false); + } } } } else { - SUCCEED("Memory model feature is only supported for gfx90a, Hence" - "skipping the testcase for this GPU " << device); + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + WARN("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); } } diff --git a/projects/hip/tests/catch/unit/stream/hipStreamWaitEvent.cc b/projects/hip/tests/catch/unit/stream/hipStreamWaitEvent.cc index cea39fcd54..d050efea2a 100644 --- a/projects/hip/tests/catch/unit/stream/hipStreamWaitEvent.cc +++ b/projects/hip/tests/catch/unit/stream/hipStreamWaitEvent.cc @@ -16,7 +16,13 @@ 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 : +Unit_hipStreamWaitEvent_Negative - Test unsuccessful hipStreamWaitEvent when either event or flags are invalid +Unit_hipStreamWaitEvent_UninitializedStream_Negative - Test unsuccessful hipStreamWaitEvent when stream is uninitialized +Unit_hipStreamWaitEvent_Default - Test simple waiting for an event with hipStreamWaitEvent api +Unit_hipStreamWaitEvent_DifferentStreams - Test waiting for an event on a different stream with hipStreamWaitEvent api +*/ #include @@ -59,6 +65,20 @@ TEST_CASE("Unit_hipStreamWaitEvent_Negative") { } } + /* Test removed for Nvidia devices because it returns unexpected error */ +#if !HT_NVIDIA +TEST_CASE("Unit_hipStreamWaitEvent_UninitializedStream_Negative") { + hipStream_t stream{reinterpret_cast(0xFFFF)}; + hipEvent_t event{nullptr}; + + HIP_CHECK(hipEventCreate(&event)); + + HIP_CHECK_ERROR(hipStreamWaitEvent(stream, event, 0), hipErrorContextIsDestroyed); + + HIP_CHECK(hipEventDestroy(event)); +} +#endif + // Since we can not use atomic*_system on every gpu, we will use wait based on clock rate. // This wont be accurate since clock rate of a GPU varies depending on many variables including // thermals, load, utilization @@ -74,6 +94,38 @@ __global__ void waitKernel(int clockRate, int seconds) { } } +TEST_CASE("Unit_hipStreamWaitEvent_Default") { + hipStream_t stream{nullptr}; + hipEvent_t waitEvent{nullptr}; + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipEventCreate(&waitEvent)); + + REQUIRE(stream != nullptr); + REQUIRE(waitEvent != nullptr); + + int deviceId {}; + HIP_CHECK(hipGetDevice(&deviceId)); + + hipDeviceProp_t prop{}; + HIP_CHECK(hipGetDeviceProperties(&prop, deviceId)); + auto clockRate = prop.clockRate; + + waitKernel<<<1, 1, 0, stream>>>(clockRate, 2); // Wait for 2 seconds + + HIP_CHECK(hipEventRecord(waitEvent, stream)); + + // Make sure stream is waiting for data to be set + HIP_CHECK_ERROR(hipEventQuery(waitEvent), hipErrorNotReady); + + HIP_CHECK(hipStreamWaitEvent(stream, waitEvent, 0)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipEventDestroy(waitEvent)); +} + TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { hipStream_t blockedStreamA{nullptr}, streamBlockedOnStreamA{nullptr}, unblockingStream{nullptr}; hipEvent_t waitEvent{nullptr}; @@ -87,12 +139,15 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { REQUIRE(streamBlockedOnStreamA != nullptr); REQUIRE(waitEvent != nullptr); + int deviceId {}; + HIP_CHECK(hipGetDevice(&deviceId)); + hipDeviceProp_t prop{}; - HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + HIP_CHECK(hipGetDeviceProperties(&prop, deviceId)); auto clockRate = prop.clockRate; waitKernel<<<1, 1, 0, blockedStreamA>>>(clockRate, - 3); // wait for 5 seconds + 3); // wait for 3 seconds HIP_CHECK(hipEventRecord(waitEvent, blockedStreamA)); // Make sure stream is waiting for data to be set @@ -102,12 +157,12 @@ TEST_CASE("Unit_hipStreamWaitEvent_DifferentStreams") { waitKernel<<<1, 1, 0, streamBlockedOnStreamA>>>(clockRate, 2); // Wait for 2 seconds - // Make sure stream is waiting for event on blockedStreamA - HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady); - HIP_CHECK(hipStreamSynchronize(unblockingStream)); HIP_CHECK(hipStreamSynchronize(blockedStreamA)); + + // Make sure streamBlockedOnStreamA waited for event on blockedStreamA + HIP_CHECK_ERROR(hipStreamQuery(streamBlockedOnStreamA), hipErrorNotReady); HIP_CHECK(hipStreamSynchronize(streamBlockedOnStreamA)); // Check that both streams have finished