diff --git a/.jenkins/jenkinsfile b/.jenkins/jenkinsfile index ae4984e5a6..bed5be315e 100644 --- a/.jenkins/jenkinsfile +++ b/.jenkins/jenkinsfile @@ -1,6 +1,6 @@ def hipBuildTest(String backendLabel) { node(backendLabel) { - stage("Source sync ${backendLabel}") { + stage("SYNC - ${backendLabel}") { // Checkout hip repository with the PR patch dir("${WORKSPACE}/hip-tests") { @@ -37,7 +37,7 @@ def hipBuildTest(String backendLabel) { } } - stage("Build - Catch2 framework") { + stage("BUILD HIP - ${backendLabel}") { // Running the build on hipamd workspace dir("${WORKSPACE}/hipamd") { sh """#!/usr/bin/env bash @@ -56,7 +56,7 @@ def hipBuildTest(String backendLabel) { } } - stage("Build - HIP TESTS") { + stage("BUILD HIP TESTS - ${backendLabel}") { // Running the build on HIP TESTS workspace dir("${WORKSPACE}/hip-tests") { env.HIP_PATH = "${HIPAMD_DIR}" + "/build/install" @@ -78,7 +78,7 @@ def hipBuildTest(String backendLabel) { } } - stage('HIP Unit Tests - Catch2 framework') { + stage("TEST - ${backendLabel}") { dir("${WORKSPACE}/hip-tests") { sh """#!/usr/bin/env bash set -x diff --git a/catch/external/Catch2/cmake/Catch2/catch_include.cmake b/catch/external/Catch2/cmake/Catch2/catch_include.cmake index b698e8b96c..9d37d62260 100644 --- a/catch/external/Catch2/cmake/Catch2/catch_include.cmake +++ b/catch/external/Catch2/cmake/Catch2/catch_include.cmake @@ -3,39 +3,41 @@ get_filename_component(_cmake_path cmake ABSOLUTE) -foreach(EXEC_NAME ${exc_names}) - if(WIN32) - set(EXEC_NAME ${EXEC_NAME}.exe) - endif() - if(EXISTS "${EXEC_NAME}") - execute_process( - COMMAND "${_cmake_path}" - -D "TEST_TARGET=${TARGET}" - -D "TEST_EXECUTABLE=${EXEC_NAME}" - -D "TEST_EXECUTOR=${crosscompiling_emulator}" - -D "TEST_WORKING_DIR=${_workdir}" - -D "TEST_SPEC=${_TEST_SPEC}" - -D "TEST_EXTRA_ARGS=${_EXTRA_ARGS}" - -D "TEST_PROPERTIES=${_PROPERTIES}" - -D "TEST_PREFIX=${_TEST_PREFIX}" - -D "TEST_SUFFIX=${_TEST_SUFFIX}" - -D "TEST_LIST=${_TEST_LIST}" - -D "TEST_REPORTER=${_REPORTER}" - -D "TEST_OUTPUT_DIR=${_OUTPUT_DIR}" - -D "TEST_OUTPUT_PREFIX=${_OUTPUT_PREFIX}" - -D "TEST_OUTPUT_SUFFIX=${_OUTPUT_SUFFIX}" - -D "CTEST_FILE=${ctestfilepath}" - -P "${_CATCH_ADD_TEST_SCRIPT}" - OUTPUT_VARIABLE output - RESULT_VARIABLE result - WORKING_DIRECTORY "${TEST_WORKING_DIR}" - ) - else() - message("executable not found : ${EXEC_NAME}" ) - endif() -endforeach() +if(NOT EXISTS "${ctestfilepath}") + foreach(EXEC_NAME ${exc_names}) + if(WIN32) + set(EXEC_NAME ${EXEC_NAME}.exe) + endif() + if(EXISTS "${EXEC_NAME}") + execute_process( + COMMAND "${_cmake_path}" + -D "TEST_TARGET=${TARGET}" + -D "TEST_EXECUTABLE=${EXEC_NAME}" + -D "TEST_EXECUTOR=${crosscompiling_emulator}" + -D "TEST_WORKING_DIR=${_workdir}" + -D "TEST_SPEC=${_TEST_SPEC}" + -D "TEST_EXTRA_ARGS=${_EXTRA_ARGS}" + -D "TEST_PROPERTIES=${_PROPERTIES}" + -D "TEST_PREFIX=${_TEST_PREFIX}" + -D "TEST_SUFFIX=${_TEST_SUFFIX}" + -D "TEST_LIST=${_TEST_LIST}" + -D "TEST_REPORTER=${_REPORTER}" + -D "TEST_OUTPUT_DIR=${_OUTPUT_DIR}" + -D "TEST_OUTPUT_PREFIX=${_OUTPUT_PREFIX}" + -D "TEST_OUTPUT_SUFFIX=${_OUTPUT_SUFFIX}" + -D "CTEST_FILE=${ctestfilepath}" + -P "${_CATCH_ADD_TEST_SCRIPT}" + OUTPUT_VARIABLE output + RESULT_VARIABLE result + WORKING_DIRECTORY "${TEST_WORKING_DIR}" + ) + else() + message("executable not found : ${EXEC_NAME}" ) + endif() + endforeach() +endif() if(EXISTS "${ctestfilepath}") -# include the generated ctest file for execution -include(${ctestfilepath}) + # include the generated ctest file for execution + include(${ctestfilepath}) endif() diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index b5eb2dfca6..a4e1fd0cb1 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -15,6 +15,8 @@ "Unit_hipMemset_Negative_OutOfBoundsPtr", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", + "Unit_hipMemAdvise_AccessedBy_All_Devices", + "Unit_hipMemAdvise_No_Flag_Interference", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", @@ -27,6 +29,8 @@ "Unit_hipStreamAttachMemAsync_Negative_Parameters", "Unit_hipMemGetAddressRange_Positive", "Unit_hipGraphAddMemcpyNode1D_Negative_Basic", - "Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo" + "Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo", + "intermittent issue: corrupted double-linked list", + "Unit_hipGraphRetainUserObject_Functional_2" ] } diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index 6bb11e842c..0e87e64b1f 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -94,6 +94,8 @@ "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", + "Unit_hipMemAdvise_AccessedBy_All_Devices", + "Unit_hipMemAdvise_No_Flag_Interference", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", diff --git a/catch/unit/deviceLib/Atomic_func.cc b/catch/unit/deviceLib/Atomic_func.cc new file mode 100644 index 0000000000..4a28839c6f --- /dev/null +++ b/catch/unit/deviceLib/Atomic_func.cc @@ -0,0 +1,119 @@ +/* +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 to validate atomicInc and atomicDec functions. +// if TestToRun=1, then atomicInc function will be tested and validated +// if TestToRun=2, then atomicDec function will be tested and validated. + + +// kernel function for atomicInc +static __global__ void AtomicCheckInc(int* g_ptr) { + atomicInc(reinterpret_cast(&g_ptr[0]), 17); +} + +// kernel function for atomicDec +static __global__ void AtomicCheckDec(int* g_ptr) { + atomicDec(reinterpret_cast(&g_ptr[0]), 25); +} + +// verify results for atomicInc +static int verifyResultInc(int value) { + int limit = 17; + value = (value >= limit) ? 0 : value + 1; + return value; +} + +// verify results for atomicDec +static int verifyResultDec(int value) { + int limit = 25; + value = ((value == 0) || (value > limit)) ? limit : value - 1; + return value; +} + +// common fuction to launch atomic functions kernel. +static void launchAtomicFunction(int *Hptr, int val, int TestToRun) { + unsigned int memSize = sizeof(int) * 1; + int *dptr{nullptr}; + // allocate device memory + HIP_CHECK(hipMalloc(reinterpret_cast(&dptr), memSize)); + // copy host memory to device + HIP_CHECK(hipMemcpy(dptr, Hptr, memSize, hipMemcpyHostToDevice)); + // launch kernel function + if (TestToRun == 1) { + AtomicCheckInc<<<1, 1>>>(dptr); + } else if (TestToRun == 2) { + AtomicCheckDec<<<1, 1>>>(dptr); + } + // copy back from device to host + HIP_CHECK(hipMemcpy(Hptr, dptr, memSize, hipMemcpyDeviceToHost)); + // verify the results. + if (TestToRun == 1) { + int result = verifyResultInc(val); + REQUIRE(result == Hptr[0]); + } else if (TestToRun == 2) { + int result = verifyResultDec(val); + REQUIRE(result == Hptr[0]); + } + // Cleanup memory + HIP_CHECK(hipFree(dptr)); +} + +TEST_CASE("Unit_AtomicFunctions_Inc") { + int *Hptr{nullptr}; + int val; + // Allocate Host memory + Hptr = reinterpret_cast(malloc(sizeof(int))); + SECTION("Test case when value is lesser than limit") { + val = Hptr[0] = 10; + launchAtomicFunction(Hptr, val, 1); + } + SECTION("Test case when value is greater than limit") { + val = Hptr[0] = 20; + launchAtomicFunction(Hptr, val, 1); + } + SECTION("Test case when value is equal to the limit") { + val = Hptr[0] = 17; + launchAtomicFunction(Hptr, val, 1); + } + free(Hptr); +} + +TEST_CASE("Unit_AtomicFunctions_Dec") { + int *Hptr{nullptr}; + int val; + // Allocate Host memory + Hptr = reinterpret_cast(malloc(sizeof(int))); + SECTION("Test case when value is less than limit") { + val = Hptr[0] = 4; + launchAtomicFunction(Hptr, val, 2); + } + SECTION("Test case when value is greater than limit") { + val = Hptr[0] = 31; + launchAtomicFunction(Hptr, val, 2); + } + SECTION("Test case when value is equal to the limit") { + val = Hptr[0] = 25; + launchAtomicFunction(Hptr, val, 2); + } + free(Hptr); +} diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index 02b542e71e..3ba4bddd39 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -13,6 +13,7 @@ set(TEST_SRC syncthreadsand.cc syncthreadscount.cc syncthreadsor.cc + Atomic_func.cc ) if(UNIX) diff --git a/catch/unit/graph/hipGraphGetRootNodes.cc b/catch/unit/graph/hipGraphGetRootNodes.cc index 2be8ea5110..43b900772f 100644 --- a/catch/unit/graph/hipGraphGetRootNodes.cc +++ b/catch/unit/graph/hipGraphGetRootNodes.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 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 @@ -28,6 +28,11 @@ Functional :: 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. + 5) Functional Test to validate number of root nodes when dependencies in the graph are dynamically varied. + 6) Functional Test to validate number of root nodes when dependencies in the graph are dynamically varied + in a cloned graph. + 7) Functional Test to validate number of root nodes when a graph with N independent nodes is added as a + child node to another graph. Argument Validation :: 1) Pass graph as nullptr and verify api returns error code. @@ -41,6 +46,12 @@ Argument Validation :: #include #include +#define NUM_OF_DUMMY_NODES 8 + +static __global__ void dummyKernel() { + return; +} + /** * Functional Test for API fetching root node list */ @@ -346,3 +357,110 @@ TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(C_d)); } + +/** + * Functional Test to validate number of root nodes when dependencies + * in the graph are dynamically varied. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes") { + hipGraph_t graph; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + size_t numRootNodes{}; + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == NUM_OF_DUMMY_NODES); + // Start creating dependencies in a chain + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + numRootNodes = 0; + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i], + &kernelnode[i+1], 1)); + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == (NUM_OF_DUMMY_NODES - i - 1)); + } + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test to validate number of root nodes when dependencies + * in the graph are dynamically varied in a cloned graph. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes_ClonedGrph") { + hipGraph_t graph, clonedgraph; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&clonedgraph, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + size_t numRootNodes{}; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + HIP_CHECK(hipGraphGetRootNodes(clonedgraph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == NUM_OF_DUMMY_NODES); + // Start creating dependencies in a chain + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + numRootNodes = 0; + hipGraphNode_t node1, node2; + HIP_CHECK(hipGraphNodeFindInClone(&node1, kernelnode[i], clonedgraph)); + HIP_CHECK(hipGraphNodeFindInClone(&node2, kernelnode[i+1], clonedgraph)); + HIP_CHECK(hipGraphAddDependencies(clonedgraph, &node1, &node2, 1)); + HIP_CHECK(hipGraphGetRootNodes(clonedgraph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == (NUM_OF_DUMMY_NODES - i - 1)); + } + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test to validate number of root nodes when a graph with N + * independent nodes is added as a child node to another graph. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NRootNodesAsChildGraph") { + hipGraph_t graph, graph1; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + hipGraphNode_t child_node; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&graph1, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + HIP_CHECK(hipGraphAddChildGraphNode(&child_node, graph1, + nullptr, 0, graph)); + size_t numRootNodes{}; + HIP_CHECK(hipGraphGetRootNodes(graph1, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == 1); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 876dae382b..2fa38dd802 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -108,6 +108,7 @@ set(TEST_SRC hipMemcpySync.cc hipMemsetSync.cc hipMemsetAsync.cc + hipMemAdvise_old.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc hipStreamAttachMemAsync.cc @@ -194,6 +195,7 @@ set(TEST_SRC hipMemcpySync.cc hipMemsetSync.cc hipMemsetAsync.cc + hipMemAdvise_old.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc hipMemRangeGetAttributes_old.cc diff --git a/catch/unit/memory/hipMemAdvise.cc b/catch/unit/memory/hipMemAdvise.cc index e5412ddfc7..87059a594e 100644 --- a/catch/unit/memory/hipMemAdvise.cc +++ b/catch/unit/memory/hipMemAdvise.cc @@ -1,943 +1,322 @@ /* -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. */ -/* Test Case Description: - Scenario-1: The following Function Tests the working of flags which can be - assigned to HMM memory using hipMemAdvise() api - Scenario-2: Negative tests on hipMemAdvise() api - Scenario-3: The following function tests various scenarios around the flag - 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api - Scenario-4: The following function tests various scenarios around the flag - 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api - Scenario-5: The following function verifies if assigning of a flag - invalidates the earlier flag which was assigned to the same memory region - using hipMemAdvise() - Scenario-6: The following function tests if peers can set - hipMemAdviseSetAccessedBy flag - on HMM memory prefetched on each of the other gpus - Scenario-7: Set AccessedBy flag and check value returned by - hipMemRangeGetAttribute() It should be -2(same is observed on cuda) - Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the - memory to device 1, then probe for AccessedBy flag using - hipMemRangeGetAttribute() we should still see the said flag is set for - device 0 - Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to - device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should - return 0 - 2) Unset AccessedBy to 0 and set it to device 1 followed by - PreferredLocation to device 1, check for AccessedBy flag using - hipMemRangeGetAttribute() it should return 1 - Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset - AccessedBy, launch kernel. We should not have any access issues - Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation - flag to the allocated memory and launch a kernel. Kernel should get executed - successfully without hang or segfault - Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then - get attribute using the api hipMemRangeGetAttribute() for - hipMemRangeAttributeLastPrefetchLocation the value returned should be -2 - Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch - the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using - hipMemRangeGetAttribute(), we should get 1 - Scenario-14: Allocate HMM memory, set ReadMostly followed by - PreferredLocation, probe for hipMemRangeAttributeReadMostly and - hipMemRangeAttributePreferredLocation - using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. - In other words setting of hipMemRangeAttributePreferredLocation should not - impact hipMemRangeAttributeReadMostly advise to the memory - Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and - launch kernel on all other gpus except 0. This test case may discover any - effect or access denial case arising due to setting ReadMostly only to a - particular gpu -*/ - #include -#if __linux__ -#include -#include -#include -#endif +#include +#include +#include -// Kernel function -__global__ void MemAdvseKernel(int n, int *x) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - if (index < n) - x[index] = x[index] * x[index]; -} - -// Kernel -__global__ void MemAdvise2(int *Hmm, int n) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) { - Hmm[i] = Hmm[i] + 10; +static inline hipMemoryAdvise GetUnsetMemAdvice(const hipMemoryAdvise advice) { + switch (advice) { + case hipMemAdviseSetAccessedBy: + return hipMemAdviseUnsetAccessedBy; + case hipMemAdviseSetReadMostly: + return hipMemAdviseUnsetReadMostly; + case hipMemAdviseSetPreferredLocation: + return hipMemAdviseUnsetPreferredLocation; + default: + assert("Invalid hipMemoryAdvise enumerator"); + return advice; } } -// Kernel -__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) { - Hmm1[i] = Hmm[i] + 10; +static inline hipMemRangeAttribute GetMemAdviceAttr(const hipMemoryAdvise advice) { + switch (advice) { + case hipMemAdviseSetAccessedBy: + return hipMemRangeAttributeAccessedBy; + case hipMemAdviseSetReadMostly: + return hipMemRangeAttributeReadMostly; + case hipMemAdviseSetPreferredLocation: + return hipMemRangeAttributePreferredLocation; + default: + assert("Invalid hipMemoryAdvise enumerator"); + return static_cast(-1); } } - -static bool CheckError(hipError_t err, int LineNo) { - if (err == hipSuccess) { - WARN("Error expected but received hipSuccess at line no.:" << LineNo); - return false; - } else { - return true; - } -} - -static int HmmAttrPrint() { - int managed = 0; - WARN("The following are the attribute values related to HMM for" - " device 0:\n"); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); - WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeConcurrentManagedAccess, 0)); - WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccess, 0)); - WARN("hipDeviceAttributePageableMemoryAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); - WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" << managed); - - HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, - 0)); - WARN("hipDeviceAttributeManagedMemory: " << managed); - return managed; -} - - -// The following Function Tests the working of flags which can be assigned -// to HMM memory using hipMemAdvise() api -TEST_CASE("Unit_hipMemAdvise_TstFlags") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int NumDevs = 0, *Outpt = nullptr; - int MEM_SIZE = 4*1024, A_CONST = 9999; - float *Hmm = nullptr; - int AttrVal = 0; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - Outpt = new int[NumDevs]; - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); - // With the following for loop we iterate through each of the Gpus in the - // system set and unset the flags and check the behavior. - for (int i = 0; i < NumDevs; ++i) { - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i)); - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributeReadMostly, Hmm, - MEM_SIZE * 2)); - if (AttrVal != 1) { - WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n"); - IfTestPassed = false; - } - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly, - i)); - - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributeReadMostly, Hmm, - (MEM_SIZE * 2))); - if (AttrVal != 0) { - WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n"); - IfTestPassed = false; - } - AttrVal = A_CONST; - // Currently hipMemAdviseSetPreferredLocation and - // hipMemAdviseSetAccessedBy - // flags are resulting in issues: SWDEV-267357 - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, - hipMemAdviseSetPreferredLocation, i)); - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributePreferredLocation, - Hmm, (MEM_SIZE * 2))); - if (AttrVal != i) { - WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n"); - IfTestPassed = false; - } - AttrVal = A_CONST; - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, - hipMemAdviseUnsetPreferredLocation, i)); - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributePreferredLocation, - Hmm, (MEM_SIZE * 2))); - if (AttrVal == i) { - WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation "); - WARN("flag failed!\n"); - IfTestPassed = false; - } - for (int m = 0; m < NumDevs; ++m) { - Outpt[m] = A_CONST; - } - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i)); - HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), - hipMemRangeAttributeAccessedBy, Hmm, - (MEM_SIZE * 2))); - if ((Outpt[0]) != i) { - WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n"); - IfTestPassed = false; - } - for (int m = 0; m < NumDevs; ++m) { - Outpt[m] = A_CONST; - } - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy, - i)); - HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), - hipMemRangeAttributeAccessedBy, Hmm, - (MEM_SIZE * 2))); - if ((Outpt[0]) >= 0) { - WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n"); - IfTestPassed = false; - } +std::vector GetDevicesWithAdviseSupport() { + const auto device_count = HipTest::getDeviceCount(); + std::vector supported_devices; + supported_devices.reserve(device_count + 1); + for (int i = 0; i < device_count; ++i) { + if (DeviceAttributesSupport(i, hipDeviceAttributeManagedMemory, + hipDeviceAttributeConcurrentManagedAccess)) { + supported_devices.push_back(i); + } + } + return supported_devices; +} + +TEST_CASE("Unit_hipMemAdvise_Set_Unset_Basic") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + const auto device = GENERATE_COPY(from_range(supported_devices)); + + const auto SetUnset = [=](const hipMemoryAdvise advice) { + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + int32_t attribute = 0; + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 1 : device) == attribute); + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, GetUnsetMemAdvice(advice), device)); + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 0 : hipInvalidDeviceId) == attribute); + }; + +// Disabled due to defect - EXSWHTEC-132 +#if HT_NVIDIA + SECTION("hipMemAdviseSetAccessedBy") { SetUnset(hipMemAdviseSetAccessedBy); } +#endif + SECTION("hipMemAdviseSetReadMostly") { SetUnset(hipMemAdviseSetReadMostly); } + SECTION("hipMemAdviseSetPreferredLocation") { SetUnset(hipMemAdviseSetPreferredLocation); } +} + +TEST_CASE("Unit_hipMemAdvise_No_Flag_Interference") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + const auto device = GENERATE_COPY(from_range(supported_devices)); + + std::array advice{hipMemAdviseSetReadMostly, hipMemAdviseSetPreferredLocation, + hipMemAdviseSetAccessedBy}; + for (int i = 0; i < 6; ++i) { + std::next_permutation(std::begin(advice), std::end(advice)); + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + + for (const auto a : advice) { + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, a, device)); + } + + for (const auto a : advice) { + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(a), + alloc.ptr(), kPageSize)); + REQUIRE((a == hipMemAdviseSetReadMostly ? 1 : device) == attribute); } - delete [] Outpt; - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); } } -TEST_CASE("Unit_hipMemAdvise_NegtveTsts") { - HipTest::HIP_SKIP_TEST("Fixed few issues to match with Nvidia, Skip now to avoid CI failures"); - return; - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int NumDevs = 0, MEM_SIZE = 4*1024; - float *Hmm = nullptr; - std::string str; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); +TEST_CASE("Unit_hipMemAdvise_Rounding") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + const auto device = supported_devices.front(); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, 3 * kPageSize); + REQUIRE_FALSE(reinterpret_cast(alloc.ptr()) % kPageSize); + const auto [offset, width] = + GENERATE_COPY(std::make_pair(kPageSize / 4, kPageSize / 2), // Withing page + std::make_pair(kPageSize / 2, kPageSize), // Across page border + std::make_pair(kPageSize / 2, kPageSize * 2)); // Across two page borders + HIP_CHECK(hipMemAdvise(alloc.ptr() + offset, width, hipMemAdviseSetAccessedBy, device)); + constexpr auto RoundDown = [](const intptr_t a, const intptr_t n) { return a - a % n; }; + constexpr auto RoundUp = [RoundDown](const intptr_t a, const intptr_t n) { + return RoundDown(a + n - 1, n); + }; + const auto base = alloc.ptr(); + const auto rounded_up = RoundUp(offset + width, kPageSize); + unsigned int attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), hipMemRangeAttributeAccessedBy, + reinterpret_cast(base), rounded_up)); + REQUIRE(device == attribute); + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), hipMemRangeAttributeAccessedBy, + alloc.ptr(), 3 * kPageSize)); + REQUIRE((rounded_up == 3 * kPageSize ? device : hipInvalidDeviceId) == attribute); +} + +TEST_CASE("Unit_hipMemAdvise_Flags_Do_Not_Cause_Prefetch") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + supported_devices.push_back(hipCpuDeviceId); + + const auto Test = [](const int device, const hipMemoryAdvise advice) { + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), + hipMemRangeAttributeLastPrefetchLocation, alloc.ptr(), + kPageSize)); + REQUIRE(attribute == hipInvalidDeviceId); + }; + const auto device = + GENERATE_COPY(from_range(std::begin(supported_devices), std::end(supported_devices))); + + SECTION("hipMemAdviseSetPreferredLocation") { Test(device, hipMemAdviseSetPreferredLocation); } +// Disabled on AMD due to defect - EXSWHTEC-132 +#if HT_NVIDIA + SECTION("hipMemAdviseSetAccessedBy") { Test(device, hipMemAdviseSetAccessedBy); } +#endif +} + +TEST_CASE("Unit_hipMemAdvise_Read_Write_After_Advise") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + constexpr size_t count = kPageSize / sizeof(*alloc.ptr()); + + const auto ReadWriteManagedMemory = [&](const int device, const hipMemoryAdvise advice) { + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + + std::fill_n(alloc.ptr(), count, -1); + ArrayFindIfNot(alloc.ptr(), -1, count); + for (int i = 0; i < supported_devices.size(); ++i) { + HIP_CHECK(hipSetDevice(supported_devices[i])); + VectorIncrement<<>>(alloc.ptr(), 1, count); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + ArrayFindIfNot(alloc.ptr(), i, count); + } + + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 1 : device) == attribute); + }; + +// Disabled on AMD due to defect - EXSWHTEC-133 +#if HT_NVIDIA + SECTION("ReadMostly") { ReadWriteManagedMemory(hipInvalidDeviceId, hipMemAdviseSetReadMostly); } +#endif + supported_devices.push_back(hipCpuDeviceId); + + const auto device = + GENERATE_COPY(from_range(std::begin(supported_devices), std::end(supported_devices))); + supported_devices.pop_back(); + SECTION("PreferredLocation") { ReadWriteManagedMemory(device, hipMemAdviseSetPreferredLocation); } + +// Disabled on AMD due to defect - EXSWHTEC-132 +#if HT_NVIDIA + SECTION("AccessedBy") { ReadWriteManagedMemory(device, hipMemAdviseSetAccessedBy); } +#endif +} + +TEST_CASE("Unit_hipMemAdvise_Prefetch_After_Advise") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + supported_devices.push_back(hipCpuDeviceId); + const auto advice = GENERATE(hipMemAdviseSetReadMostly, hipMemAdviseSetPreferredLocation + // Skipped due to defect - EXSWHTEC - 132 + // hipMemAdviseSetAccessedBy + ); + const auto device = GENERATE_COPY(from_range(supported_devices)); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + + for (const auto d : supported_devices) { + HIP_CHECK(hipMemPrefetchAsync(alloc.ptr(), kPageSize, d)); + HIP_CHECK(hipStreamSynchronize(nullptr)); + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), + hipMemRangeAttributeLastPrefetchLocation, alloc.ptr(), + kPageSize)); + REQUIRE(d == attribute); + } + + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 1 : device) == attribute); +} + +TEST_CASE("Unit_hipMemAdvise_AccessedBy_All_Devices") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + for (const auto device : supported_devices) { + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, hipMemAdviseSetAccessedBy, device)); + } + std::vector accessed_by(supported_devices.size(), hipInvalidDeviceId); + HIP_CHECK(hipMemRangeGetAttribute(accessed_by.data(), sizeof(accessed_by.data()), + hipMemRangeAttributeAccessedBy, alloc.ptr(), kPageSize)); + REQUIRE_THAT(accessed_by, Catch::Matchers::Equals(supported_devices)); +} + +TEST_CASE("Unit_hipMemAdvise_Negative_Parameters") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + const auto device = supported_devices.front(); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + +// Disabled on NVIDIA due to defect - EXSWHTEC-122 #if HT_AMD - // Passing invalid value(99) device param - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, 99), __LINE__); - - // Passing invalid value(-12) device param - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, -12), __LINE__); -#endif - // Passing NULL as first parameter instead of valid pointer to a memory - IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, 0), __LINE__); - - // Passing 0 for count(2nd param) parameter - IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly, - 0), __LINE__); - - // Passing count much more than actually allocated value - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6, - hipMemAdviseSetReadMostly, 0), __LINE__); - - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + SECTION("Invalid advice") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize, static_cast(-1), device), + hipErrorInvalidValue); } -} - -// The following function tests various scenarios around the flag -// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api -TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - // Check that when a page fault occurs for the memory region set to devPtr, - // the data is migrated to the destn processor - int MEM_SIZE = 4096, A_CONST = 9999; - int *Hmm = nullptr, NumDevs = 0, dev = A_CONST; - bool IfTestPassed = true; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal)); - for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) { - Hmm[i] = 4; - } - for (int devId = 0; devId < NumDevs; ++devId) { - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, - hipMemAdviseSetPreferredLocation, devId)); - int NumElms = ((MEM_SIZE * 3)/4); - MemAdvseKernel<<>>(NumElms, Hmm); - int dev = A_CONST; - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), - hipMemRangeAttributePreferredLocation, - Hmm, MEM_SIZE * 3)); - if (dev != devId) { - WARN("Memory observed to be not available on expected location\n"); - WARN("line no: " << __LINE__); - WARN("dev: " << dev); - IfTestPassed = false; - } - } - - // Check that when preferred location is set for a memory region, - // data can still be prefetched using hipMemPrefetchAsync - hipStream_t strm; - dev = A_CONST; - for (int devId = 0; devId < NumDevs; ++devId) { - HIP_CHECK(hipSetDevice(devId)); - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, - hipMemAdviseSetPreferredLocation, devId)); - HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, MEM_SIZE * 3)); - if (dev != devId) { - WARN("Memory reported to be not available at the Prefetched "); - WARN("location with device id: " << devId); - WARN("line no: " << __LINE__); - WARN("dev: " << dev); - IfTestPassed = false; - } - HIP_CHECK(hipStreamDestroy(strm)); - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -// The following function tests various scenarios around the flag -// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api - -TEST_CASE("Unit_hipMemAdvise_ReadMostly") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int MEM_SIZE = 4096, A_CONST = 9999; - float *Hmm = nullptr; - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE)); - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - Hmm[i] = A_CONST; - } - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); - // Checking if the data can be read after setting hipMemAdviseSetReadMostly - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - if (Hmm[i] != A_CONST) { - WARN("Didn't find expected value in Hmm memory after setting"); - WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__); - IfTestPassed = false; - } - } - - // Checking if the memory region can be modified - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - Hmm[i] = A_CONST; - } - - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - if (Hmm[i] != A_CONST) { - WARN("Didn't find expected value in Hmm memory after Modification\n"); - WARN("line no.: " << __LINE__); - IfTestPassed = false; - } - } - - int out = A_CONST; - HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, - Hmm, MEM_SIZE)); - if (out != 1) { - WARN("out value: " << out); - IfTestPassed = false; - } - // Checking the advise attribute after prefetch - HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0)); - HIP_CHECK(hipDeviceSynchronize()); - HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), - hipMemRangeAttributeReadMostly, Hmm, - MEM_SIZE)); - if (out != 1) { - WARN("Attribute assigned to memory changed after calling "); - WARN("hipMemPrefetchAsync(). line no.: " << __LINE__); - WARN("out value: " << out); - IfTestPassed = false; - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -// The following function verifies if assigning of a flag invalidates the -// earlier flag which was assigned to the same memory region using -// hipMemAdvise() -TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int MEM_SIZE = 4*4096, A_CONST = 9999; - float *Hmm = nullptr; - int NumDevs = 0, dev = A_CONST; - - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); - for (int i = 0; i < NumDevs; ++i) { - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i)); - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), - hipMemRangeAttributeReadMostly, Hmm, - MEM_SIZE)); - if (dev != 1) { - WARN("hipMemAdviseSetReadMostly flag did not take affect despite "); - WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); - IfTestPassed = false; - break; - } - dev = A_CONST; - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation, - i)); - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), - hipMemRangeAttributePreferredLocation, - Hmm, MEM_SIZE)); - if (dev != i) { - WARN("hipMemAdviseSetPreferredLocation flag did not take affect "); - WARN("despite setting it using hipMemAdvise()\n"); - WARN("line no.: " << __LINE__); - IfTestPassed = false; - break; - } - - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i)); - dev = A_CONST; - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, - MEM_SIZE)); - if (dev != i) { - WARN("hipMemAdviseSetAccessedBy flag did not take affect despite "); - WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); - IfTestPassed = false; - break; - } - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i)); - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -// The following function tests if peers can set hipMemAdviseSetAccessedBy flag -// on HMM memory prefetched on each of the other gpus -#if HT_AMD -TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;; - int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0; - - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - if (NumDevs < 2) { - SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test" - " the scenario. This system has GPUs less than 2"); - } - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); - for (int i = 0; i < NumDevs; ++i) { - HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0)); - for (int j = 0; j < NumDevs; ++j) { - if (i == j) - continue; - HIP_CHECK(hipSetDevice(j)); - HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i)); - if (CanAccessPeer) { - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j)); - for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { - Hmm[m] = 4; - } - HIP_CHECK(hipDeviceEnablePeerAccess(i, 0)); - MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>( - (MEM_SIZE/sizeof(int)), Hmm); - HIP_CHECK(hipDeviceSynchronize()); - // Verifying the result - for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { - if (Hmm[m] != 16) { - flag = 1; - } - } - if (flag) { - WARN("Didnt get Expected results with device: " << j); - WARN("line no.: " << __LINE__); - IfTestPassed = false; - flag = 0; - } - } - } - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} #endif - -/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute() - It should be -2(same is observed on cuda)*/ -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, 2*4096)); - if (data != -2) { - WARN("Didnt get expected value!!\n"); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to - device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute() - we should still see the said flag is set for device 0*/ -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999, Ngpus = 0; - HIP_CHECK(hipGetDeviceCount(&Ngpus)); - if (Ngpus >= 2) { - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); - HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); - if (data != 0) { - WARN("Didnt get expected behavior at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); - HIP_CHECK(hipStreamDestroy(strm)); - HIP_CHECK(hipFree(Hmm)); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - - -/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1 - check for AccessedBy flag using hipMemRangeGetAttribute() it should - return 0 - 2) Unset AccessedBy to 0 and set it to device 1 followed by - PreferredLocation to device 1, check for AccessedBy flag using - hipMemRangeGetAttribute() it should return 1*/ - -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999, Ngpus = 0; - HIP_CHECK(hipGetDeviceCount(&Ngpus)); - if (Ngpus >= 2) { - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); - if (data != 0) { - WARN("Didnt get expected behavior at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); - if (data != 1) { - WARN("Didnt get expected behavior at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm)); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -/* Set AccessedBy flag to HMM memory launch a kernel and then unset - AccessedBy, launch kernel. We should not have any access issues*/ - -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; - int DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseSetAccessedBy, 0)); - // Initializing memory - for (int i = 0; i < NumElms; ++i) { - Hmm[i] = InitVal; - } - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); - // launching kernel from each one of the gpus - MemAdvise2<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - - // verifying the final result - for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + 10)) { - DataMismatch++; - } - } - - if (DataMismatch != 0) { - WARN("DataMismatch is observed at line: " << __LINE__); - REQUIRE(false); - } - - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseUnsetAccessedBy, 0)); - MemAdvise2<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // verifying the final result - for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + (2*10))) { - DataMismatch++; - } - } - - if (DataMismatch != 0) { - WARN("DataMismatch is observed at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipStreamDestroy(strm)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to - the allocated memory and launch a kernel. Kernel should get executed - successfully without hang or segfault*/ -#if __linux__ && HT_AMD -TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { - if ((setenv("HSA_XNACK", "1", 1)) != 0) { - WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); - REQUIRE(false); - } - // The following code block checks for gfx90a so as to skip if the device is not MI200 - - hipDeviceProp_t prop; - int device; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - std::string gfxName(prop.gcnArchName); - - if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { - int stat = 0; - if (fork() == 0) { - // The below part should be inside fork - 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)); - for (int i = 0; i < NumElms; ++i) { - Mllc[i] = InitVal; - } - hipStream_t strm; - int DataMismatch = 0; - HIP_CHECK(hipStreamCreate(&strm)); - // The following hipMemAdvise() call is made to know if advise on - // aligned_alloc() is causing any issue - HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - for (int i = 0; i < NumElms; ++i) { - if (Mllc[i] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!!"); - exit(9); // 9 for failure - } else { - exit(10); // 10 for Pass result - } - } else { - 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 == 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); - WARN("Memory model feature is only supported for gfx90a, Hence" - "skipping the testcase for this GPU " << device); + const auto advice = GENERATE(hipMemAdviseSetAccessedBy, hipMemAdviseSetReadMostly, + hipMemAdviseSetPreferredLocation); + SECTION("count == 0") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), 0, advice, device), hipErrorInvalidValue); } -} +// Disabled due to defect - EXSWHTEC-131 +#if HT_NVIDIA + SECTION("count larger than allocation size") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize + 1, advice, device), hipErrorInvalidValue); + } #endif -/* Allocate Hmm memory, set advise to PreferredLocation and then get - attribute using the api hipMemRangeGetAttribute() for - hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/ - -TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 4096)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, 4096)); - if (data != -2) { - WARN("Didnt receive expected value."); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + SECTION("dev_ptr == nullptr") { + HIP_CHECK_ERROR(hipMemAdvise(nullptr, kPageSize, advice, device), hipErrorInvalidValue); } -} - -/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem - to device1, probe for hipMemRangeAttributeLastPrefetchLocation using - hipMemRangeGetAttribute(), we should get 1*/ - -TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") { - int NumDevs = 0; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - if (NumDevs >= 2) { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - hipStream_t strm; - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, 4096)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, 4096)); - if (data != 1) { - WARN("Didnt receive expected value!!"); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } - } else { - SUCCEED("This system has less than 2 gpus hence skipping the test.\n"); + SECTION("dev_ptr pointing to non-managed memory") { + LinearAllocGuard alloc(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize, advice, device), hipErrorInvalidValue); } -} - -/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe - for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation - using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. - In other words setting of hipMemRangeAttributePreferredLocation should not - impact hipMemRangeAttributeReadMostly advise to the memory*/ - -TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 4096)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeReadMostly, Hmm, - 4096)); - if (data != 1) { - WARN("Didnt receive expected value at line: " << data); - REQUIRE(false); - } - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributePreferredLocation, Hmm, - 4096)); - if (data != 0) { - WARN("Didnt receive expected value at line: " << data); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); +// Disabled on AMD due to defect - EXSWHTEC-130 +#if HT_NVIDIA + SECTION("Invalid device") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize, advice, hipInvalidDeviceId), + (advice == hipMemAdviseSetReadMostly ? hipSuccess : hipErrorInvalidDevice)); } -} - - - -/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel - on all other gpus except 0. This test case may discover any effect or - access denial case arising due to setting ReadMostly only to a particular - gpu*/ - -TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int Ngpus = 0; - HIP_CHECK(hipGetDeviceCount(&Ngpus)); - if (Ngpus < 2) { - SUCCEED("This test needs atleast two gpus to run." - "Hence skipping the test.\n"); - } - int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; - int *Hmm1 = NULL, DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); - // Initializing memory - for (int i = 0; i < NumElms; ++i) { - Hmm[i] = InitVal; - } - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseSetReadMostly, 0)); - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); -#if HT_AMD - SECTION("Launch Kernel on all other gpus") { - // launching kernel from each one of the gpus - for (int i = 1; i < Ngpus; ++i) { - DataMismatch = 0; - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); - MemAdvise3<<>>(Hmm, Hmm1, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // verifying the results - for (int j = 0; j < NumElms; ++j) { - if (Hmm1[j] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch is observed with the gpu: " << i); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm1)); - } - } - - SECTION("Launch Kernel on all other gpus and manipulate the content") { - for (int i = 0; i < Ngpus; ++i) { - DataMismatch = 0; - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseSetReadMostly, i)); - MemAdvise2<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - } - // verifying the final result - for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + Ngpus * 10)) { - DataMismatch++; - } - } - - if (DataMismatch != 0) { - WARN("DataMismatch is observed at line: " << __LINE__); - REQUIRE(false); - } - } #endif - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipStreamDestroy(strm)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } } - - -TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); - if (data != 0) { - WARN("Didnt receive expected value!!"); - REQUIRE(false); - } - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); - if (data != -2) { - WARN("Didnt receive expected value!!"); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - diff --git a/catch/unit/memory/hipMemAdvise_old.cc b/catch/unit/memory/hipMemAdvise_old.cc new file mode 100644 index 0000000000..b951998053 --- /dev/null +++ b/catch/unit/memory/hipMemAdvise_old.cc @@ -0,0 +1,938 @@ +/* +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. +*/ + +/* Test Case Description: + Scenario-1: The following Function Tests the working of flags which can be + assigned to HMM memory using hipMemAdvise() api + Scenario-2: Negative tests on hipMemAdvise() api + Scenario-3: The following function tests various scenarios around the flag + 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api + Scenario-4: The following function tests various scenarios around the flag + 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api + Scenario-5: The following function verifies if assigning of a flag + invalidates the earlier flag which was assigned to the same memory region + using hipMemAdvise() + Scenario-6: The following function tests if peers can set + hipMemAdviseSetAccessedBy flag + on HMM memory prefetched on each of the other gpus + Scenario-7: Set AccessedBy flag and check value returned by + hipMemRangeGetAttribute() It should be -2(same is observed on cuda) + Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the + memory to device 1, then probe for AccessedBy flag using + hipMemRangeGetAttribute() we should still see the said flag is set for + device 0 + Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to + device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should + return 0 + 2) Unset AccessedBy to 0 and set it to device 1 followed by + PreferredLocation to device 1, check for AccessedBy flag using + hipMemRangeGetAttribute() it should return 1 + Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset + AccessedBy, launch kernel. We should not have any access issues + Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation + flag to the allocated memory and launch a kernel. Kernel should get executed + successfully without hang or segfault + Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then + get attribute using the api hipMemRangeGetAttribute() for + hipMemRangeAttributeLastPrefetchLocation the value returned should be -2 + Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch + the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using + hipMemRangeGetAttribute(), we should get 1 + Scenario-14: Allocate HMM memory, set ReadMostly followed by + PreferredLocation, probe for hipMemRangeAttributeReadMostly and + hipMemRangeAttributePreferredLocation + using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. + In other words setting of hipMemRangeAttributePreferredLocation should not + impact hipMemRangeAttributeReadMostly advise to the memory + Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and + launch kernel on all other gpus except 0. This test case may discover any + effect or access denial case arising due to setting ReadMostly only to a + particular gpu +*/ + +#include +#if __linux__ +#include +#include +#include +#endif + +// Kernel function +__global__ void MemAdvseKernel(int n, int *x) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < n) + x[index] = x[index] * x[index]; +} + +// Kernel +__global__ void MemAdvise2(int *Hmm, int n) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + Hmm[i] = Hmm[i] + 10; + } +} + +// Kernel +__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + Hmm1[i] = Hmm[i] + 10; + } +} + + +static bool CheckError(hipError_t err, int LineNo) { + if (err == hipSuccess) { + WARN("Error expected but received hipSuccess at line no.:" << LineNo); + return false; + } else { + return true; + } +} + +static int HmmAttrPrint() { + int managed = 0; + WARN("The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeConcurrentManagedAccess, 0)); + WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" << managed); + + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + WARN("hipDeviceAttributeManagedMemory: " << managed); + return managed; +} + + +// The following Function Tests the working of flags which can be assigned +// to HMM memory using hipMemAdvise() api +TEST_CASE("Unit_hipMemAdvise_TstFlags") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, *Outpt = nullptr; + int MEM_SIZE = 4*1024, A_CONST = 9999; + float *Hmm = nullptr; + int AttrVal = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + Outpt = new int[NumDevs]; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); + // With the following for loop we iterate through each of the Gpus in the + // system set and unset the flags and check the behavior. + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE * 2)); + if (AttrVal != 1) { + WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly, + i)); + + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributeReadMostly, Hmm, + (MEM_SIZE * 2))); + if (AttrVal != 0) { + WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + AttrVal = A_CONST; + // Currently hipMemAdviseSetPreferredLocation and + // hipMemAdviseSetAccessedBy + // flags are resulting in issues: SWDEV-267357 + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, + hipMemAdviseSetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributePreferredLocation, + Hmm, (MEM_SIZE * 2))); + if (AttrVal != i) { + WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n"); + IfTestPassed = false; + } + AttrVal = A_CONST; + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, + hipMemAdviseUnsetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributePreferredLocation, + Hmm, (MEM_SIZE * 2))); + if (AttrVal == i) { + WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation "); + WARN("flag failed!\n"); + IfTestPassed = false; + } + for (int m = 0; m < NumDevs; ++m) { + Outpt[m] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i)); + HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), + hipMemRangeAttributeAccessedBy, Hmm, + (MEM_SIZE * 2))); + if ((Outpt[0]) != i) { + WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + for (int m = 0; m < NumDevs; ++m) { + Outpt[m] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy, + i)); + HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), + hipMemRangeAttributeAccessedBy, Hmm, + (MEM_SIZE * 2))); + if ((Outpt[0]) >= 0) { + WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + } + delete [] Outpt; + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +TEST_CASE("Unit_hipMemAdvise_NegtveTsts") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, MEM_SIZE = 4*1024; + float *Hmm = nullptr; + std::string str; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); + // Passing NULL as first parameter instead of valid pointer to a memory + IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, 0), __LINE__); + + // Passing 0 for count(2nd param) parameter + IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly, + 0), __LINE__); + + // Passing count much more than actually allocated value + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6, + hipMemAdviseSetReadMostly, 0), __LINE__); + + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function tests various scenarios around the flag +// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api +TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + // Check that when a page fault occurs for the memory region set to devPtr, + // the data is migrated to the destn processor + int MEM_SIZE = 4096, A_CONST = 9999; + int *Hmm = nullptr, NumDevs = 0, dev = A_CONST; + bool IfTestPassed = true; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal)); + for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) { + Hmm[i] = 4; + } + for (int devId = 0; devId < NumDevs; ++devId) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, + hipMemAdviseSetPreferredLocation, devId)); + int NumElms = ((MEM_SIZE * 3)/4); + MemAdvseKernel<<>>(NumElms, Hmm); + int dev = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), + hipMemRangeAttributePreferredLocation, + Hmm, MEM_SIZE * 3)); + if (dev != devId) { + WARN("Memory observed to be not available on expected location\n"); + WARN("line no: " << __LINE__); + WARN("dev: " << dev); + IfTestPassed = false; + } + } + + // Check that when preferred location is set for a memory region, + // data can still be prefetched using hipMemPrefetchAsync + hipStream_t strm; + dev = A_CONST; + for (int devId = 0; devId < NumDevs; ++devId) { + HIP_CHECK(hipSetDevice(devId)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, + hipMemAdviseSetPreferredLocation, devId)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, MEM_SIZE * 3)); + if (dev != devId) { + WARN("Memory reported to be not available at the Prefetched "); + WARN("location with device id: " << devId); + WARN("line no: " << __LINE__); + WARN("dev: " << dev); + IfTestPassed = false; + } + HIP_CHECK(hipStreamDestroy(strm)); + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function tests various scenarios around the flag +// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api + +TEST_CASE("Unit_hipMemAdvise_ReadMostly") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int MEM_SIZE = 4096, A_CONST = 9999; + float *Hmm = nullptr; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE)); + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + Hmm[i] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + // Checking if the data can be read after setting hipMemAdviseSetReadMostly + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + if (Hmm[i] != A_CONST) { + WARN("Didn't find expected value in Hmm memory after setting"); + WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__); + IfTestPassed = false; + } + } + + // Checking if the memory region can be modified + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + Hmm[i] = A_CONST; + } + + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + if (Hmm[i] != A_CONST) { + WARN("Didn't find expected value in Hmm memory after Modification\n"); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + } + } + + int out = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, + Hmm, MEM_SIZE)); + if (out != 1) { + WARN("out value: " << out); + IfTestPassed = false; + } + // Checking the advise attribute after prefetch + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE)); + if (out != 1) { + WARN("Attribute assigned to memory changed after calling "); + WARN("hipMemPrefetchAsync(). line no.: " << __LINE__); + WARN("out value: " << out); + IfTestPassed = false; + } + // hipMemAdvise should succeed for SetReadMostly and UnsetReadMostly + // irrespective of the device + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 99)); + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetReadMostly, -12)); + + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function verifies if assigning of a flag invalidates the +// earlier flag which was assigned to the same memory region using +// hipMemAdvise() +TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int MEM_SIZE = 4*4096, A_CONST = 9999; + float *Hmm = nullptr; + int NumDevs = 0, dev = A_CONST; + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE)); + if (dev != 1) { + WARN("hipMemAdviseSetReadMostly flag did not take affect despite "); + WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); + IfTestPassed = false; + break; + } + dev = A_CONST; + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation, + i)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributePreferredLocation, + Hmm, MEM_SIZE)); + if (dev != i) { + WARN("hipMemAdviseSetPreferredLocation flag did not take affect "); + WARN("despite setting it using hipMemAdvise()\n"); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + break; + } + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i)); + dev = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, + MEM_SIZE)); + if (dev != i) { + WARN("hipMemAdviseSetAccessedBy flag did not take affect despite "); + WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); + IfTestPassed = false; + break; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i)); + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +// The following function tests if peers can set hipMemAdviseSetAccessedBy flag +// on HMM memory prefetched on each of the other gpus +#if HT_AMD +TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;; + int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0; + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + if (NumDevs < 2) { + SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test" + " the scenario. This system has GPUs less than 2"); + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0)); + for (int j = 0; j < NumDevs; ++j) { + if (i == j) + continue; + HIP_CHECK(hipSetDevice(j)); + HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i)); + if (CanAccessPeer) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j)); + for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + Hmm[m] = 4; + } + HIP_CHECK(hipDeviceEnablePeerAccess(i, 0)); + MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>( + (MEM_SIZE/sizeof(int)), Hmm); + HIP_CHECK(hipDeviceSynchronize()); + // Verifying the result + for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + if (Hmm[m] != 16) { + flag = 1; + } + } + if (flag) { + WARN("Didnt get Expected results with device: " << j); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + flag = 0; + } + } + } + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} +#endif + + +/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute() + It should be -2(same is observed on cuda)*/ +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 2*4096)); + if (data != -2) { + WARN("Didnt get expected value!!\n"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to + device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute() + we should still see the said flag is set for device 0*/ +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999, Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus >= 2) { + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); + HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipFree(Hmm)); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + +/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1 + check for AccessedBy flag using hipMemRangeGetAttribute() it should + return 0 + 2) Unset AccessedBy to 0 and set it to device 1 followed by + PreferredLocation to device 1, check for AccessedBy flag using + hipMemRangeGetAttribute() it should return 1*/ + +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999, Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus >= 2) { + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 1) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Set AccessedBy flag to HMM memory launch a kernel and then unset + AccessedBy, launch kernel. We should not have any access issues*/ + +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; + int DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetAccessedBy, 0)); + // Initializing memory + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = InitVal; + } + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + // launching kernel from each one of the gpus + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + 10)) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseUnsetAccessedBy, 0)); + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + (2*10))) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to + the allocated memory and launch a kernel. Kernel should get executed + successfully without hang or segfault*/ +#if __linux__ && HT_AMD +TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { + if ((setenv("HSA_XNACK", "1", 1)) != 0) { + WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); + REQUIRE(false); + } + // The following code block checks for gfx90a so as to skip if the device is not MI200 + + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + int stat = 0; + if (fork() == 0) { + // The below part should be inside fork + 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)); + for (int i = 0; i < NumElms; ++i) { + Mllc[i] = InitVal; + } + hipStream_t strm; + int DataMismatch = 0; + HIP_CHECK(hipStreamCreate(&strm)); + // The following hipMemAdvise() call is made to know if advise on + // aligned_alloc() is causing any issue + HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + for (int i = 0; i < NumElms; ++i) { + if (Mllc[i] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch observed!!"); + exit(9); // 9 for failure + } else { + exit(10); // 10 for Pass result + } + } else { + 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 == 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); + WARN("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } + +} +#endif + +/* Allocate Hmm memory, set advise to PreferredLocation and then get + attribute using the api hipMemRangeGetAttribute() for + hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 4096)); + if (data != -2) { + WARN("Didnt receive expected value."); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem + to device1, probe for hipMemRangeAttributeLastPrefetchLocation using + hipMemRangeGetAttribute(), we should get 1*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") { + int NumDevs = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + if (NumDevs >= 2) { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + hipStream_t strm; + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 4096)); + if (data != 1) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } + } else { + SUCCEED("This system has less than 2 gpus hence skipping the test.\n"); + } +} + + +/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe + for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation + using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. + In other words setting of hipMemRangeAttributePreferredLocation should not + impact hipMemRangeAttributeReadMostly advise to the memory*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + 4096)); + if (data != 1) { + WARN("Didnt receive expected value at line: " << data); + REQUIRE(false); + } + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, + 4096)); + if (data != 0) { + WARN("Didnt receive expected value at line: " << data); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + +/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel + on all other gpus except 0. This test case may discover any effect or + access denial case arising due to setting ReadMostly only to a particular + gpu*/ + +TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus < 2) { + SUCCEED("This test needs atleast two gpus to run." + "Hence skipping the test.\n"); + } + int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; + int *Hmm1 = NULL, DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); + // Initializing memory + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = InitVal; + } + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetReadMostly, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); +#if HT_AMD + SECTION("Launch Kernel on all other gpus") { + // launching kernel from each one of the gpus + for (int i = 1; i < Ngpus; ++i) { + DataMismatch = 0; + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); + MemAdvise3<<>>(Hmm, Hmm1, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // verifying the results + for (int j = 0; j < NumElms; ++j) { + if (Hmm1[j] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch is observed with the gpu: " << i); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm1)); + } + } + + SECTION("Launch Kernel on all other gpus and manipulate the content") { + for (int i = 0; i < Ngpus; ++i) { + DataMismatch = 0; + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetReadMostly, i)); + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + } + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + Ngpus * 10)) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + } +#endif + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); + if (data != -2) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + diff --git a/catch/unit/memory/hipPointerGetAttributes.cc b/catch/unit/memory/hipPointerGetAttributes.cc index d4ba5960ce..cbc4554b3a 100644 --- a/catch/unit/memory/hipPointerGetAttributes.cc +++ b/catch/unit/memory/hipPointerGetAttributes.cc @@ -260,8 +260,6 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { REQUIRE(attribs.devicePointer != attribs2.devicePointer); } HIP_CHECK(hipFree(A_d)); - e = hipPointerGetAttributes(&attribs, A_d); - REQUIRE(e == hipErrorInvalidValue); // Device-visible host memory printf("\nDevice-visible host memory (hipHostMalloc)\n"); @@ -272,15 +270,10 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { char* ptr1 = reinterpret_cast(attribs.hostPointer); REQUIRE((ptr1 + Nbytes / 2) == reinterpret_cast(attribs2.hostPointer)); - HIP_CHECK(hipHostFree(A_Pinned_h)); - e = hipPointerGetAttributes(&attribs, A_Pinned_h); - REQUIRE(e == hipErrorInvalidValue); // OS memory printf("\nOS-allocated memory (malloc)\n"); - e = hipPointerGetAttributes(&attribs, A_OSAlloc_h); - REQUIRE(e == hipErrorInvalidValue); } TEST_CASE("Unit_hipPointerGetAttributes_ClusterAlloc") { diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index 2f015d726a..4803737598 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -1,6 +1,12 @@ set(TEST_SRC vulkan_test.cc hipExternalMemoryGetMappedBuffer.cc + hipImportExternalMemory.cc + hipDestroyExternalMemory.cc + hipWaitExternalSemaphoresAsync.cc + hipSignalExternalSemaphoresAsync.cc + hipImportExternalSemaphore.cc + hipDestroyExternalSemaphore.cc ) find_package(Vulkan) diff --git a/catch/unit/vulkan_interop/hipDestroyExternalMemory.cc b/catch/unit/vulkan_interop/hipDestroyExternalMemory.cc new file mode 100644 index 0000000000..3f7c3cc13d --- /dev/null +++ b/catch/unit/vulkan_interop/hipDestroyExternalMemory.cc @@ -0,0 +1,45 @@ +/* +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 "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipDestroyExternalMemory_Vulkan_Negative_Parameters") { + SECTION("extMem == nullptr") { + HIP_CHECK_ERROR(hipDestroyExternalMemory(nullptr), hipErrorInvalidValue); + } + +// Segfaults in CUDA +// Disabled on AMD due to defect - EXSWHTEC-187 +#if HT_AMD && 0 + SECTION("Double free") { + VulkanTest vkt(enable_validation); + const auto storage = vkt.CreateMappedStorage(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + auto desc = vkt.BuildMemoryDescriptor(storage.memory, sizeof(*storage.host_ptr)); + hipExternalMemory_t ext_memory; + HIP_CHECK(hipImportExternalMemory(&ext_memory, &desc)); + + HIP_CHECK(hipDestroyExternalMemory(ext_memory)); + HIP_CHECK_ERROR(hipDestroyExternalMemory(ext_memory), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc b/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc new file mode 100644 index 0000000000..dcef2d7fad --- /dev/null +++ b/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc @@ -0,0 +1,40 @@ +/* +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 "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters") { + SECTION("extSem == nullptr") { + HIP_CHECK_ERROR(hipDestroyExternalSemaphore(nullptr), hipErrorInvalidValue); + } + +// Segfaults in CUDA +#if HT_AMD + SECTION("Double free") { + VulkanTest vkt(enable_validation); + const auto ext_semaphore = ImportBinarySemaphore(vkt); + HIP_CHECK(hipDestroyExternalSemaphore(ext_semaphore)); + HIP_CHECK_ERROR(hipDestroyExternalSemaphore(ext_semaphore), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipImportExternalMemory.cc b/catch/unit/vulkan_interop/hipImportExternalMemory.cc new file mode 100644 index 0000000000..5abfada9d5 --- /dev/null +++ b/catch/unit/vulkan_interop/hipImportExternalMemory.cc @@ -0,0 +1,76 @@ +/* +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 "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipImportExternalMemory_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + const auto storage = vkt.CreateMappedStorage(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + auto desc = vkt.BuildMemoryDescriptor(storage.memory, sizeof(*storage.host_ptr)); + hipExternalMemory_t ext_memory; + +// Disabled due to defect - EXSWHTEC-182 +#if HT_NVIDIA + SECTION("extMem_out == nullptr") { + HIP_CHECK_ERROR(hipImportExternalMemory(nullptr, &desc), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-183 +#if HT_NVIDIA + SECTION("memHandleDesc == nullptr") { + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, nullptr), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-185 +#if HT_NVIDIA + SECTION("memHandleDesc.size == 0") { + desc.size = 0; + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-186 +#if HT_NVIDIA + SECTION("Invalid memHandleDesc.flags") { + desc.flags = 2; + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-184 +#if HT_NVIDIA + SECTION("Invalid memHandleDesc.type") { + desc.type = static_cast(-1); + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif + +#ifdef _WIN32 + SECTION("memHandleDesc.handle == NULL") { + desc.handle.win32.handle = NULL; + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipImportExternalSemaphore.cc b/catch/unit/vulkan_interop/hipImportExternalSemaphore.cc new file mode 100644 index 0000000000..2920645459 --- /dev/null +++ b/catch/unit/vulkan_interop/hipImportExternalSemaphore.cc @@ -0,0 +1,56 @@ +/* +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 "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipImportExternalSemaphore_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + auto handle_desc = vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t ext_semaphore; + + SECTION("extSem_out == nullptr") { + HIP_CHECK_ERROR(hipImportExternalSemaphore(nullptr, &handle_desc), hipErrorInvalidValue); + } + + SECTION("semHandleDesc == nullptr") { + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, nullptr), hipErrorInvalidValue); + } + + SECTION("semHandleDesc.flags != 0") { + handle_desc.flags = 1; + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue); + } + + SECTION("Invalid semHandleDesc.type") { + handle_desc.type = static_cast(-1); + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue); + } + +#ifdef _WIN32 + SECTION("semHandleDesc.handle == NULL") { + handle_desc.handle.win32.handle = NULL; + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc b/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc new file mode 100644 index 0000000000..2c5d055e66 --- /dev/null +++ b/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc @@ -0,0 +1,204 @@ +/* +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 "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.waitSemaphoreCount = 1; + submit_info.pWaitSemaphores = &semaphore; + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + + REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); + + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 0; + HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr)); + PollStream(nullptr, hipSuccess); + + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") { + VulkanTest vkt(enable_validation); + constexpr uint64_t signal_value = 2; + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = signal_value; + + HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr)); + PollStream(nullptr, hipSuccess); + + uint64_t sem_value = 0u; + VK_CHECK_RESULT(vkGetSemaphoreCounterValue(vkt.GetDevice(), semaphore, &sem_value)); + + REQUIRE(2 == sem_value); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} +#endif + +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + +#if HT_AMD + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; +#else + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; +#endif + + const auto command_buffer = vkt.GetCommandBuffer(); + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_binary_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_binary_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc)); + + const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type); + const auto hip_timeline_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type); + hipExternalSemaphore_t hip_timeline_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc)); + + uint64_t wait_values[] = {1, 0}; + VkTimelineSemaphoreSubmitInfo timeline_submit_info = {}; + timeline_submit_info.sType = VK_STRUCTURE_TYPE_TIMELINE_SEMAPHORE_SUBMIT_INFO; + timeline_submit_info.waitSemaphoreValueCount = 2; + timeline_submit_info.pWaitSemaphoreValues = wait_values; + + VkSemaphore wait_semaphores[] = {timeline_semaphore, binary_semaphore}; + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.waitSemaphoreCount = 2; + submit_info.pWaitSemaphores = wait_semaphores; + submit_info.pNext = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? &timeline_submit_info : nullptr; + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + + REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); + + hipExternalSemaphoreSignalParams binary_signal_params = {}; + binary_signal_params.params.fence.value = 0; + hipExternalSemaphoreSignalParams timeline_signal_params = {}; + timeline_signal_params.params.fence.value = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 2 : 0; + hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; + hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params, timeline_signal_params}; + HIP_CHECK(hipSignalExternalSemaphoresAsync(ext_semaphores, signal_params, 2, nullptr)); + + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); +} + +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + + SECTION("extSemArray == nullptr") { + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(nullptr, &signal_params, 1, nullptr), + hipErrorInvalidValue); + } + + SECTION("paramsArray == nullptr") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, nullptr, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Wait params flags != 0") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + signal_params.flags = 1; + HIP_CHECK_ERROR( + hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Invalid stream") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, stream), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc b/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc new file mode 100644 index 0000000000..edebebe52a --- /dev/null +++ b/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc @@ -0,0 +1,231 @@ +/* +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 "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; + hip_ext_semaphore_wait_params.flags = 0; + hip_ext_semaphore_wait_params.params.fence.value = 0; + HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, + nullptr)); + PollStream(nullptr, hipErrorNotReady); + + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.signalSemaphoreCount = 1; + submit_info.pSignalSemaphores = &semaphore; + + *src_storage.host_ptr = 42; + + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + PollStream(nullptr, hipSuccess); + + REQUIRE(42 == *dst_storage.host_ptr); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") { + VulkanTest vkt(enable_validation); + + const auto [wait_value, signal_value] = + GENERATE(std::make_pair(2, 2), std::make_pair(2, 3), std::make_pair(3, 2)); + INFO("Wait value: " << wait_value << ", signal value: " << signal_value); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; + hip_ext_semaphore_wait_params.flags = 0; + hip_ext_semaphore_wait_params.params.fence.value = wait_value; + HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, + nullptr)); + PollStream(nullptr, hipErrorNotReady); + + VkSemaphoreSignalInfo signal_info = {}; + signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; + signal_info.semaphore = semaphore; + signal_info.value = signal_value; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + if (wait_value > signal_value) { + PollStream(nullptr, hipErrorNotReady); + signal_info.value = wait_value; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + } + PollStream(nullptr, hipSuccess); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} +#endif + +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") { + VulkanTest vkt(enable_validation); + +#if HT_AMD + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; +#else + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; +#endif + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_binary_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_binary_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc)); + + const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type); + const auto hip_timeline_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type); + hipExternalSemaphore_t hip_timeline_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc)); + + hipExternalSemaphoreWaitParams binary_semaphore_wait_params = {}; + binary_semaphore_wait_params.params.fence.value = 0; + + hipExternalSemaphoreWaitParams timeline_semaphore_wait_params = {}; + timeline_semaphore_wait_params.params.fence.value = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 0; + + hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; + hipExternalSemaphoreWaitParams wait_params[] = {binary_semaphore_wait_params, + timeline_semaphore_wait_params}; + HIP_CHECK(hipWaitExternalSemaphoresAsync(ext_semaphores, wait_params, 2, nullptr)); + + PollStream(nullptr, hipErrorNotReady); + + if (second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE) { + VkSemaphoreSignalInfo signal_info = {}; + signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; + signal_info.semaphore = timeline_semaphore; + signal_info.value = 1; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + + PollStream(nullptr, hipErrorNotReady); + } + + VkSubmitInfo submit_info = {}; + VkSemaphore signal_semaphores[] = {binary_semaphore, timeline_semaphore}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.signalSemaphoreCount = second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 2; + submit_info.pSignalSemaphores = + second_semaphore_type == VK_SEMAPHORE_TYPE_MAX_ENUM ? &binary_semaphore : signal_semaphores; + + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + PollStream(nullptr, hipSuccess); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); +} + +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + + SECTION("extSemArray == nullptr") { + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(nullptr, &wait_params, 1, nullptr), + hipErrorInvalidValue); + } + + SECTION("paramsArray == nullptr") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, nullptr, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Wait params flag != 0") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + wait_params.flags = 1; + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &wait_params, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Invalid stream") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &wait_params, 1, stream), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } +} \ No newline at end of file diff --git a/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt b/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt index f1a8bf8a82..248dd355b0 100644 --- a/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt +++ b/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt @@ -25,12 +25,12 @@ endif () if(NOT DEFINED HIP_PATH) if(NOT DEFINED ENV{HIP_PATH}) - set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH "Path to which HIP has been installed") + set(HIP_PATH ${ROCM_PATH} CACHE PATH "Path to which HIP has been installed") else() set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed") endif() endif() -set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH}) +set(CMAKE_MODULE_PATH "${HIP_PATH}/hip/cmake" ${CMAKE_MODULE_PATH}) set(CMAKE_HIP_ARCHITECTURES OFF) project(12_cmake) @@ -53,7 +53,7 @@ set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FO hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS}) # Search for rocm in common locations -list(APPEND CMAKE_PREFIX_PATH ${HIP_PATH} ${ROCM_PATH}) +list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) find_package(hip QUIET) if(TARGET hip::host) message(STATUS "Found hip::host at ${hip_DIR}") diff --git a/samples/2_Cookbook/16_assembly_to_executable/Makefile b/samples/2_Cookbook/16_assembly_to_executable/Makefile index b82ec8fdbe..0a32aa2b9c 100644 --- a/samples/2_Cookbook/16_assembly_to_executable/Makefile +++ b/samples/2_Cookbook/16_assembly_to_executable/Makefile @@ -21,15 +21,15 @@ ifeq ($(OS),Windows_NT) $(error Makefile is not supported on windows platform. Please use cmake instead to build sample.) endif ROCM_PATH?= $(wildcard /opt/rocm/) -HIP_PATH?= $(wildcard $(ROCM_PATH)/hip) +HIP_PATH?= $(ROCM_PATH) ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc -CLANG=$(HIP_PATH)/../llvm/bin/clang -LLVM_MC=$(HIP_PATH)/../llvm/bin/llvm-mc -CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/../llvm/bin/clang-offload-bundler +CLANG=$(HIP_PATH)/llvm/bin/clang +LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc +CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler SRCS=square.cpp diff --git a/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile b/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile index 835f4b2789..d68304dcac 100644 --- a/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile +++ b/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile @@ -21,17 +21,17 @@ ifeq ($(OS),Windows_NT) $(error Makefile is not supported on windows platform. Please use cmake instead to build sample.) endif ROCM_PATH?= $(wildcard /opt/rocm/) -HIP_PATH?= $(wildcard $(ROCM_PATH)/hip) +HIP_PATH?= $(ROCM_PATH) ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc -CLANG=$(HIP_PATH)/../llvm/bin/clang -LLVM_MC=$(HIP_PATH)/../llvm/bin/llvm-mc -CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/../llvm/bin/clang-offload-bundler -LLVM_AS=$(HIP_PATH)/../llvm/bin/llvm-as -LLVM_DIS=$(HIP_PATH)/../llvm/bin/llvm-dis +CLANG=$(HIP_PATH)/llvm/bin/clang +LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc +CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler +LLVM_AS=$(HIP_PATH)/llvm/bin/llvm-as +LLVM_DIS=$(HIP_PATH)/llvm/bin/llvm-dis SRCS=square.cpp