diff --git a/projects/hip/tests/catch/CMakeLists.txt b/projects/hip/tests/catch/CMakeLists.txt index 0fddc8f346..434dc92b2b 100644 --- a/projects/hip/tests/catch/CMakeLists.txt +++ b/projects/hip/tests/catch/CMakeLists.txt @@ -69,6 +69,7 @@ add_subdirectory(unit) add_subdirectory(ABM) add_subdirectory(hipTestMain) add_subdirectory(stress) +add_subdirectory(TypeQualifiers) if(UNIX) add_subdirectory(multiproc) diff --git a/projects/hip/tests/catch/TypeQualifiers/CMakeLists.txt b/projects/hip/tests/catch/TypeQualifiers/CMakeLists.txt new file mode 100644 index 0000000000..3c60bdf411 --- /dev/null +++ b/projects/hip/tests/catch/TypeQualifiers/CMakeLists.txt @@ -0,0 +1,12 @@ +# Common Tests +set(TEST_SRC + hipManagedKeyword.cc +) + + +# Create shared lib of all tests +add_library(TypeQualifiers SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) + +# Add dependency on build_tests to build it on this custom target +add_dependencies(build_tests TypeQualifiers) + diff --git a/projects/hip/tests/catch/unit/memory/hipManagedKeyword.cc b/projects/hip/tests/catch/TypeQualifiers/hipManagedKeyword.cc similarity index 100% rename from projects/hip/tests/catch/unit/memory/hipManagedKeyword.cc rename to projects/hip/tests/catch/TypeQualifiers/hipManagedKeyword.cc diff --git a/projects/hip/tests/catch/hipTestMain/CMakeLists.txt b/projects/hip/tests/catch/hipTestMain/CMakeLists.txt index 5b62dbbb49..7fe98c71d9 100644 --- a/projects/hip/tests/catch/hipTestMain/CMakeLists.txt +++ b/projects/hip/tests/catch/hipTestMain/CMakeLists.txt @@ -71,3 +71,14 @@ target_link_libraries(StressTest PRIVATE memory stdc++fs) add_dependencies(build_stress_test StressTest) add_custom_target(stress_test COMMAND StressTest) +# Space Specifiers/Qualifiers exe +add_executable(TypeQualifierTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc) +if(HIP_PLATFORM MATCHES "amd") + set_property(TARGET TypeQualifierTests PROPERTY CXX_STANDARD 17) +else() + target_compile_options(TypeQualifierTests PUBLIC -std=c++17) +endif() +target_link_libraries(TypeQualifierTests PRIVATE TypeQualifiers stdc++fs) + +catch_discover_tests(TypeQualifierTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") +add_dependencies(build_tests TypeQualifierTests) diff --git a/projects/hip/tests/catch/include/hip_test_smi.hh b/projects/hip/tests/catch/include/hip_test_smi.hh new file mode 100644 index 0000000000..8962d4f230 --- /dev/null +++ b/projects/hip/tests/catch/include/hip_test_smi.hh @@ -0,0 +1,88 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +/** + * @brief Error codes retured by rocm_smi_lib functions + */ +typedef enum { + RSMI_STATUS_SUCCESS = 0x0, //!< Operation was successful + RSMI_STATUS_INVALID_ARGS, //!< Passed in arguments are not valid + RSMI_STATUS_NOT_SUPPORTED, //!< The requested information or + //!< action is not available for the + //!< given input, on the given system + RSMI_STATUS_FILE_ERROR, //!< Problem accessing a file. This + //!< may because the operation is not + //!< supported by the Linux kernel + //!< version running on the executing + //!< machine + RSMI_STATUS_PERMISSION, //!< Permission denied/EACCESS file + //!< error. Many functions require + //!< root access to run. + RSMI_STATUS_OUT_OF_RESOURCES, //!< Unable to acquire memory or other + //!< resource + RSMI_STATUS_INTERNAL_EXCEPTION, //!< An internal exception was caught + RSMI_STATUS_INPUT_OUT_OF_BOUNDS, //!< The provided input is out of + //!< allowable or safe range + RSMI_STATUS_INIT_ERROR, //!< An error occurred when rsmi + //!< initializing internal data + //!< structures + RSMI_INITIALIZATION_ERROR = RSMI_STATUS_INIT_ERROR, + RSMI_STATUS_NOT_YET_IMPLEMENTED, //!< The requested function has not + //!< yet been implemented in the + //!< current system for the current + //!< devices + RSMI_STATUS_NOT_FOUND, //!< An item was searched for but not + //!< found + RSMI_STATUS_INSUFFICIENT_SIZE, //!< Not enough resources were + //!< available for the operation + RSMI_STATUS_INTERRUPT, //!< An interrupt occurred during + //!< execution of function + RSMI_STATUS_UNEXPECTED_SIZE, //!< An unexpected amount of data + //!< was read + RSMI_STATUS_NO_DATA, //!< No data was found for a given + //!< input + RSMI_STATUS_UNEXPECTED_DATA, //!< The data read or provided to + //!< function is not what was expected + RSMI_STATUS_BUSY, //!< A resource or mutex could not be + //!< acquired because it is already + //!< being used + RSMI_STATUS_REFCOUNT_OVERFLOW, //!< An internal reference counter + //!< exceeded INT32_MAX + + RSMI_STATUS_UNKNOWN_ERROR = 0xFFFFFFFF, //!< An unknown error occurred +} rsmi_status_t; + + +/** + * @brief Types of memory + */ +typedef enum { + RSMI_MEM_TYPE_FIRST = 0, + + RSMI_MEM_TYPE_VRAM = RSMI_MEM_TYPE_FIRST, //!< VRAM memory + RSMI_MEM_TYPE_VIS_VRAM, //!< VRAM memory that is visible + RSMI_MEM_TYPE_GTT, //!< GTT memory + + RSMI_MEM_TYPE_LAST = RSMI_MEM_TYPE_GTT +} rsmi_memory_type_t; diff --git a/projects/hip/tests/catch/multiproc/CMakeLists.txt b/projects/hip/tests/catch/multiproc/CMakeLists.txt index 4d2c69b114..3e19d38ef9 100644 --- a/projects/hip/tests/catch/multiproc/CMakeLists.txt +++ b/projects/hip/tests/catch/multiproc/CMakeLists.txt @@ -1,6 +1,5 @@ -# AMD Tests +# Common Tests set(LINUX_TEST_SRC - hipMallocConcurrency.cc childMalloc.cc hipDeviceComputeCapabilityMproc.cc hipDeviceGetPCIBusIdMproc.cc @@ -10,12 +9,16 @@ set(LINUX_TEST_SRC hipGetDevicePropertiesMproc.cc hipSetGetDeviceMproc.cc hipIpcMemAccessTest.cc + hipHostMallocTestsMproc.cc + hipMallocConcurrencyMproc.cc ) if(UNIX) # Create shared lib of all tests add_library(MultiProc SHARED EXCLUDE_FROM_ALL ${LINUX_TEST_SRC}) + target_link_libraries(MultiProc ${CMAKE_DL_LIBS}) + # Add dependency on build_tests to build it on this custom target add_dependencies(build_tests MultiProc) endif() diff --git a/projects/hip/tests/catch/multiproc/hipHostMallocTestsMproc.cc b/projects/hip/tests/catch/multiproc/hipHostMallocTestsMproc.cc new file mode 100644 index 0000000000..cca3c396aa --- /dev/null +++ b/projects/hip/tests/catch/multiproc/hipHostMallocTestsMproc.cc @@ -0,0 +1,314 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + 1) Validate page table allocation of hipHostMalloc() api when + HIP_VISIBLE_DEVICES set to single device. + 2) Validate page table allocation of hipHostMalloc() api when + HIP_VISIBLE_DEVICES set to list of multiple devices. +*/ + +#include +#include + +#if defined(__linux__) +#include +#include +#include +#include + +#if HT_AMD +/* AMD specific as below tests use rocm-smi */ + +/** + * Defines + */ +#define LIB_ROCMSMI "librocm_smi64.so" +#define ALLOC_SIZE (1*1024*1024) + +/** + * Global variables + */ +static rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t, + rsmi_memory_type_t, + uint64_t *); +static rsmi_status_t (*rsmi_init_fp)(uint64_t); +static rsmi_status_t (*rsmi_shut_down_fp)(); +static void *rocm_smi_h; + + +/** + * Fetches Gpu device count + */ +static void getDeviceCount(int *pdevCnt) { + int fd[2], val = 0; + pid_t childpid; + + // create pipe descriptors + pipe(fd); + + // disable visible_devices env from shell + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); + + childpid = fork(); + + if (childpid > 0) { // Parent + close(fd[1]); + // parent will wait to read the device cnt + read(fd[0], &val, sizeof(val)); + + // close the read-descriptor + close(fd[0]); + + // wait for child exit + wait(nullptr); + + *pdevCnt = val; + } else if (!childpid) { // Child + int devCnt = 1; + // writing only, no need for read-descriptor + close(fd[0]); + + HIP_CHECK(hipGetDeviceCount(&devCnt)); + // send the value on the write-descriptor: + write(fd[1], &devCnt, sizeof(devCnt)); + + // close the write descriptor: + close(fd[1]); + exit(0); + } else { // failure + *pdevCnt = 0; + return; + } +} + +/** + * Initializes rocm smi library handles + */ +static bool rocm_smi_init() { + // Open ROCm SMI Library + if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) { + printf("Error opening rocm smi library!\n"); + return false; + } + + void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get"); + if (!fnsym) { + printf("Error getting rsmi_dev_memory_usage_get() function\n"); + dlclose(rocm_smi_h); + return false; + } + rsmi_dev_memory_usage_get_fp = reinterpret_cast(fnsym); + + fnsym = dlsym(rocm_smi_h, "rsmi_init"); + if (!fnsym) { + printf("Error getting rsmi_init() function\n"); + dlclose(rocm_smi_h); + return false; + } + rsmi_init_fp = reinterpret_cast(fnsym); + + fnsym = dlsym(rocm_smi_h, "rsmi_shut_down"); + if (!fnsym) { + printf("Error getting rsmi_shut_down() function\n"); + dlclose(rocm_smi_h); + return false; + } + rsmi_shut_down_fp = reinterpret_cast(fnsym); + + uint64_t init_flags = 0; + rsmi_status_t retsmi_init; + retsmi_init = rsmi_init_fp(init_flags); + if (RSMI_STATUS_SUCCESS != retsmi_init) { + printf("Error when initializing rocm_smi\n"); + dlclose(rocm_smi_h); + return false; + } + + return true; +} + +/** + * Exits rocm smi library + */ +static void rocm_smi_exit() { + rsmi_shut_down_fp(); + dlclose(rocm_smi_h); +} + +/** + * Validates page table memory allocations + * by setting visible devices selected. + */ +static bool validatePageTableAllocations(const char *devList, int visDevCnt) { + int fd[2]; + bool testResult = false; + pid_t pid; + int numdev = 0; + + getDeviceCount(&numdev); + REQUIRE(numdev > 0); + + if (pipe(fd) < 0) { + printf("Pipe system call failed\n"); + return false; + } + + pid = fork(); + + if (!pid) { // Child process + rsmi_status_t ret; + std::vector prev, current; + uint64_t used = 0; + int tmpdev = 0, changeCnt = 0, indx = 0; + char *ptr = nullptr; + bool testPassed = true; + + // Disable visible_devices env from shell + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); + + setenv("HIP_VISIBLE_DEVICES", devList, 1); + + // First Call to initialize hip api + hipGetDeviceCount(&tmpdev); + + + // Get memory snapshot before hostmalloc + for (indx = 0; indx < numdev; indx++) { + ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used); + if (RSMI_STATUS_SUCCESS != ret) { + printf("Error while running rsmi_dev_memory_usage_get func\n"); + dlclose(rocm_smi_h); + rsmi_shut_down_fp(); + return false; + } + prev.push_back(used); + } + + HIP_CHECK(hipHostMalloc(&ptr, ALLOC_SIZE)); + + // Get memory snapshot after hostmalloc + for (indx = 0; indx < numdev; indx++) { + ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used); + if (RSMI_STATUS_SUCCESS != ret) { + printf("Error while running rsmi_dev_memory_usage_get func\n"); + dlclose(rocm_smi_h); + rsmi_shut_down_fp(); + hipHostFree(ptr); + return false; + } + current.push_back(used); + } + + for (indx = 0; indx < numdev; indx++) { + if (current[indx] - prev[indx]) + changeCnt++; + } + + // Check if memory allocation happened only for visible devices + if (changeCnt == visDevCnt) { + testPassed = true; + } else { + testPassed = false; + } + + hipHostFree(ptr); + + // writing only, no need for read-descriptor + close(fd[0]); + + // send the value on the write-descriptor: + write(fd[1], &testPassed, sizeof(testPassed)); + + // close the write descriptor: + close(fd[1]); + exit(0); + } else if (pid > 0) { // parent + close(fd[1]); + read(fd[0], &testResult, sizeof(testResult)); + close(fd[0]); + wait(nullptr); + } else { + printf("fork() failed\n"); + HIP_ASSERT(false); + } + + return testResult; +} + +/** + * Test page table allocation when HIP_VISIBLE_DEVICES set to + * single device + */ +TEST_CASE("Unit_hipHostMalloc_SingleVisibleDevicePageAlloc") { + int devCnt; + std::string str; + + if (!rocm_smi_init()) { + WARN("Testcase skipped as rocm smi not initialized/present"); + return; + } + + getDeviceCount(&devCnt); + REQUIRE(devCnt > 0); + + // Select single visible device and validate memory usage + for (int i = 0; i < devCnt; i++) { + str = std::to_string(i); + REQUIRE(validatePageTableAllocations(str.c_str(), 1) == true); + } + + rocm_smi_exit(); +} + +/** + * Test page table allocation when HIP_VISIBLE_DEVICES set to + * multiple devices + */ +TEST_CASE("Unit_hipHostMalloc_MultipleVisibleDevicesPageAlloc") { + int devCnt = 0, vdCnt = 0; + std::string str; + + if (!rocm_smi_init()) { + WARN("Testcase skipped as rocm smi not initialized/present"); + return; + } + + getDeviceCount(&devCnt); + REQUIRE(devCnt > 0); + + // Select multiple visible devices and validate memory usage + for (int i = 0; i < devCnt; i++) { + if (i == 0) + str += std::to_string(i); + else + str += "," + std::to_string(i); + + vdCnt++; + REQUIRE(validatePageTableAllocations(str.c_str(), vdCnt) == true); + } + + rocm_smi_exit(); +} +#endif // HT_AMD +#endif // __linux__ diff --git a/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc b/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc deleted file mode 100644 index f6520403d5..0000000000 --- a/projects/hip/tests/catch/multiproc/hipMallocConcurrency.cc +++ /dev/null @@ -1,171 +0,0 @@ -#include -#include -#include -#ifdef __linux__ -#include -#include -#endif -#include -#include -#include -#include - -#include - - -static constexpr size_t N = 4 * 1024 * 1024; -static constexpr unsigned blocksPerCU = 6; // to hide latency -static constexpr unsigned threadsPerBlock = 256; -/** - * Validates data consitency on supplied gpu - */ -bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { - size_t Nbytes = N * sizeof(int); - int *A_d, *B_d, *C_d; - int *A_h, *B_h, *C_h; - size_t prevAvl, prevTot, curAvl, curTot; - bool TestPassed = true; - - HIP_CHECK(hipSetDevice(gpu)); - HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); - printf("tgs allocating..\n"); - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); - - HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, - static_cast(A_d), static_cast(B_d), C_d, N); - - HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - - if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) { - printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid()); - } else { - printf("%s : Validation FAILED for gpu %d from pid %d\n", __func__, gpu, getpid()); - TestPassed &= false; - } - - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); - - if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { - // In concurrent calls on one GPU, we cannot verify leaking in this way - printf( - "%s : Memory allocation mismatch observed." - "Possible memory leak.\n", - __func__); - TestPassed &= false; - } - - return TestPassed; -} - - -#if 1 -/** - * Fetches Gpu device count - */ -void getDeviceCount1(int* pdevCnt) { -#ifdef __linux__ - int fd[2], val = 0; - pid_t childpid; - - // create pipe descriptors - pipe(fd); - - // disable visible_devices env from shell - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); - - childpid = fork(); - - if (childpid > 0) { // Parent - close(fd[1]); - // parent will wait to read the device cnt - read(fd[0], &val, sizeof(val)); - - // close the read-descriptor - close(fd[0]); - - // wait for child exit - wait(NULL); - - *pdevCnt = val; - } else if (!childpid) { // Child - int devCnt = 1; - // writing only, no need for read-descriptor - close(fd[0]); - - HIP_CHECK(hipGetDeviceCount(&devCnt)); - // send the value on the write-descriptor: - write(fd[1], &devCnt, sizeof(devCnt)); - - // close the write descriptor: - close(fd[1]); - exit(0); - } else { // failure - *pdevCnt = 1; - return; - } - -#else - HIP_CHECK(hipGetDeviceCount(pdevCnt)); -#endif -} -#endif - - -TEST_CASE("hipMallocChild_Concurrency_MultiGpu") { - bool TestPassed = false; -#ifdef __linux__ - // Parallel execution on multiple gpus from different child processes - int devCnt = 1, pid = 0; - - // Get GPU count - getDeviceCount1(&devCnt); - - // Spawn child for each GPU - for (int gpu = 0; gpu < devCnt; gpu++) { - if ((pid = fork()) < 0) { - INFO("Child_Concurrency_MultiGpu : fork() returned error" << pid); - REQUIRE(false); - - } else if (!pid) { // Child process - bool TestPassedChild = false; - TestPassedChild = validateMemoryOnGPU(gpu); - - if (TestPassedChild) { - printf("returning exit(1) for success\n"); - exit(1); // child exit with success status - } else { - printf("Child_Concurrency_MultiGpu : childpid %d failed\n", getpid()); - exit(2); // child exit with failure status - } - } - } - - // Parent shall wait for child to complete - int cnt = 0; - - for (int i = 0; i < devCnt; i++) { - int pidwait = 0, exitStatus; - pidwait = wait(&exitStatus); - - printf("exitStatus for iter:%d is %d\n", i, exitStatus); - if (pidwait < 0) { - break; - } - - if (WEXITSTATUS(exitStatus) == 1) cnt++; - } - - if (cnt && (cnt == devCnt)) TestPassed = true; - -#else - INFO("Test hipMallocChild_Concurrency_MultiGpu skipped on non-linux"); -#endif - REQUIRE(TestPassed == true); -} diff --git a/projects/hip/tests/catch/multiproc/hipMallocConcurrencyMproc.cc b/projects/hip/tests/catch/multiproc/hipMallocConcurrencyMproc.cc new file mode 100644 index 0000000000..9018f694a2 --- /dev/null +++ b/projects/hip/tests/catch/multiproc/hipMallocConcurrencyMproc.cc @@ -0,0 +1,228 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + 1) Run hipMalloc() api/kernel code on same gpu parallely from parent and child + processes, validate the results. + + 2) Execute hipMalloc() api simultaneously on all the gpus by spawning multiple + child processes. Validate buffers allocated after running kernel code. + +*/ + +#include +#include +#include + +#ifdef __linux__ +#include +#include +#include + +/** + * Fetches Gpu device count + */ +static void getDeviceCount(int* pdevCnt) { + int fd[2], val = 0; + pid_t childpid; + + // create pipe descriptors + pipe(fd); + + // disable visible_devices env from shell +#ifdef HT_NVIDIA + unsetenv("CUDA_VISIBLE_DEVICES"); +#else + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); +#endif + + childpid = fork(); + + if (childpid > 0) { // Parent + close(fd[1]); + // parent will wait to read the device cnt + read(fd[0], &val, sizeof(val)); + + // close the read-descriptor + close(fd[0]); + + // wait for child exit + wait(nullptr); + + *pdevCnt = val; + } else if (!childpid) { // Child + int devCnt = 1; + // writing only, no need for read-descriptor + close(fd[0]); + + HIP_CHECK(hipGetDeviceCount(&devCnt)); + // send the value on the write-descriptor: + write(fd[1], &devCnt, sizeof(devCnt)); + + // close the write descriptor: + close(fd[1]); + exit(0); + } else { // failure + *pdevCnt = 0; + return; + } +} + +/** + * Validates data consistency on supplied gpu + */ +static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + size_t prevAvl, prevTot, curAvl, curTot; + bool TestPassed = true; + constexpr auto N = 4 * 1024 * 1024; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + size_t Nbytes = N * sizeof(int); + + HIP_CHECK(hipSetDevice(gpu)); + HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, N); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) { + printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid()); + } else { + printf("Validation FAILED for gpu %d from pid %d\n", gpu, getpid()); + TestPassed = false; + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { + // In concurrent calls on one GPU, we cannot verify leaking in this way + printf( + "%s : Memory allocation mismatch observed." + "Possible memory leak.\n", + __func__); + TestPassed = false; + } + + return TestPassed; +} + +/** + * Parallel execution of parent and child on gpu0 + */ +TEST_CASE("Unit_hipMalloc_ChildConcurrencyDefaultGpu") { + int devCnt = 0, pid = 0; + constexpr auto resSuccess = 1, resFailure = 2; + bool TestPassed = true; + + // Get GPU count + getDeviceCount(&devCnt); + REQUIRE(devCnt > 0); + + if ((pid = fork()) < 0) { + INFO("Child_Concurrency_DefaultGpu : fork() returned error : " << pid); + HIP_ASSERT(false); + + } else if (!pid) { // Child process + bool TestPassedChild = false; + + // Allocates and validates memory on Gpu0 simultaneously with parent + TestPassedChild = validateMemoryOnGPU(0, true); + + if (TestPassedChild) { + exit(resSuccess); // child exit with success status + } else { + exit(resFailure); // child exit with failure status + } + + } else { // Parent process + int exitStatus; + + // Allocates and validates memory on Gpu0 simultaneously with child + TestPassed = validateMemoryOnGPU(0, true); + + // Wait and get result from child + pid = wait(&exitStatus); + if ((WEXITSTATUS(exitStatus) == resFailure) || (pid < 0)) + TestPassed = false; + } + + REQUIRE(TestPassed == true); +} + +/** + * Parallel execution of api on multiple gpus from + * different child processes. + */ +TEST_CASE("Unit_hipMalloc_ChildConcurrencyMultiGpu") { + int devCnt = 0, pid = 0; + constexpr auto resSuccess = 1, resFailure = 2; + + // Get GPU count + getDeviceCount(&devCnt); + REQUIRE(devCnt > 0); + + // Spawn child for each GPU + for (int gpu = 0; gpu < devCnt; gpu++) { + if ((pid = fork()) < 0) { + INFO("Child_Concurrency_MultiGpu : fork() returned error : " << pid); + REQUIRE(false); + + } else if (!pid) { // Child process + bool TestPassedChild = false; + TestPassedChild = validateMemoryOnGPU(gpu); + + if (TestPassedChild) { + exit(resSuccess); // child exit with success status + } else { + exit(resFailure); // child exit with failure status + } + } + } + + // Parent shall wait for child to complete + int passCnt = 0; + for (int i = 0; i < devCnt; i++) { + int pidwait = 0, exitStatus; + pidwait = wait(&exitStatus); + + printf("exitStatus for dev:%d is %d\n", i, WEXITSTATUS(exitStatus)); + if (pidwait < 0) { + break; + } + + if (WEXITSTATUS(exitStatus) == resSuccess) passCnt++; + } + REQUIRE(passCnt == devCnt); +} +#endif // __linux__ diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index 773abef797..61571884a1 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -40,6 +40,11 @@ set(TEST_SRC hipMemset3D.cc hipMemset2D.cc hipMemset2DAsyncMultiThreadAndKernel.cc + hipHostMallocTests.cc + hipMallocConcurrency.cc + hipMemset3DFunctional.cc + hipMemset3DNegative.cc + hipMemset3DRegressMultiThread.cc ) else() set(TEST_SRC @@ -80,6 +85,11 @@ set(TEST_SRC hipMemset3D.cc hipMemset2D.cc hipMemset2DAsyncMultiThreadAndKernel.cc + hipHostMallocTests.cc + hipMallocConcurrency.cc + hipMemset3DFunctional.cc + hipMemset3DNegative.cc + hipMemset3DRegressMultiThread.cc ) endif() # Create shared lib of all tests diff --git a/projects/hip/tests/catch/unit/memory/hipHostMallocTests.cc b/projects/hip/tests/catch/unit/memory/hipHostMallocTests.cc new file mode 100644 index 0000000000..de445e2e24 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipHostMallocTests.cc @@ -0,0 +1,61 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + 1) Test hipHostMalloc() api with ptr as nullptr and check for return value. + 2) Test hipHostMalloc() api with size as max(size_t) and check for OOM error. + 3) Test hipHostMalloc() api with flags as max(unsigned int) and validate + return value. + 4) Pass size as zero for hipHostMalloc() api and check ptr is reset with + with return value success. +*/ + +#include + +/** + * Performs argument validation of hipHostMalloc api. + */ +TEST_CASE("Unit_hipHostMalloc_ArgValidation") { + hipError_t ret; + constexpr size_t allocSize = 1000; + char *ptr; + + SECTION("Pass ptr as nullptr") { + ret = hipHostMalloc(static_cast(nullptr), allocSize); + REQUIRE(ret != hipSuccess); + } + + SECTION("Size as max(size_t)") { + ret = hipHostMalloc(&ptr, std::numeric_limits::max()); + REQUIRE(ret != hipSuccess); + } + + SECTION("Flags as max(uint)") { + ret = hipHostMalloc(&ptr, allocSize, + std::numeric_limits::max()); + REQUIRE(ret != hipSuccess); + } + + SECTION("Pass size as zero and check ptr reset") { + HIP_CHECK(hipHostMalloc(&ptr, 0)); + REQUIRE(ptr == nullptr); + } +} diff --git a/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc b/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc new file mode 100644 index 0000000000..5083ea1c25 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc @@ -0,0 +1,412 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + 1) Test hipMalloc() api passing zero size and confirming *ptr returning + nullptr. Also pass nullptr to hipFree() api. + + 2) Pass maximum value of size_t for hipMalloc() api and make sure appropriate + error is returned. + + 3) Check for hipMalloc() error code, passing invalid/null pointer. + + 4) Regress hipMalloc()/hipFree() in loop for bigger chunk of allocation + with adequate number of iterations and later test for kernel execution on + default gpu. + + 5) Regress hipMalloc()/hipFree() in loop while allocating smaller chunks + keeping maximum number of iterations and then run kernel code on default + gpu, perfom data validation. + + 6) Check hipMalloc() api adaptability when app creates small chunks of memory + continuously, stores it for later use and then frees it at later point + of time. + + 7) Multithread Scenario : Exercise hipMalloc() api parellely on all gpus from + multiple threads and regress the api. + + 8) Validate memory usage with hipMemGetInfo() while regressing hipMalloc() + api. Check for any possible memory leaks. +*/ + +#include +#include +#include + + +#include +#include +#include + + +/* Buffer size for bigger chunks in alloc/free cycles */ +static constexpr auto BuffSizeBC = 5*1024*1024; + +/* Buffer size for smaller chunks in alloc/free cycles */ +static constexpr auto BuffSizeSC = 16; + +/* You may change it for individual test. + * But default 100 is for quick return in Jenkin Build */ +static constexpr auto NumDiv = 100; + +/* Max alloc/free iterations for smaller chunks */ +static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv); + +/* Max alloc/free iterations for bigger chunks */ +static constexpr auto MaxAllocFree_BigChunks = 10000; + +/* Max alloc and pool iterations */ +static constexpr auto MaxAllocPoolIter = (2000000/NumDiv); + +/* Test status shared across threads */ +static std::atomic g_thTestPassed{true}; + + + +/** + * Validates data consistency on supplied gpu + */ +static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + size_t prevAvl, prevTot, curAvl, curTot; + bool TestPassed = true; + constexpr auto N = 4 * 1024 * 1024; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + size_t Nbytes = N * sizeof(int); + + HIP_CHECK(hipSetDevice(gpu)); + HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, N); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) { + UNSCOPED_INFO("Validation PASSED for gpu " << gpu); + } else { + UNSCOPED_INFO("Validation FAILED for gpu " << gpu); + TestPassed = false; + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { + // In concurrent calls on one GPU, we cannot verify leaking in this way + UNSCOPED_INFO( + "validateMemoryOnGPU : Memory allocation mismatch observed." + << "Possible memory leak."); + TestPassed = false; + } + + return TestPassed; +} + + +/** + * Regress memory allocation and free in loop + */ +static bool regressAllocInLoop(int gpu) { + bool TestPassed = true; + size_t tot, avail, ptot, pavail, numBytes; + int i = 0; + int *ptr; + + HIP_CHECK(hipSetDevice(gpu)); + numBytes = BuffSizeBC; + + // Exercise allocation in loop with bigger chunks + for (i = 0; i < MaxAllocFree_BigChunks; i++) { + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + HIP_CHECK(hipMalloc(&ptr, numBytes)); + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK(hipFree(ptr)); + + if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes + UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << + numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" << + pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" << + tot << ", pavail-avail=" << pavail-avail); + TestPassed = false; + break; + } + } + + // Exercise allocation in loop with smaller chunks and maximum iters + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + numBytes = BuffSizeSC; + + for (i = 0; i < MaxAllocFree_SmallChunks; i++) { + HIP_CHECK(hipMalloc(&ptr, numBytes)); + + HIP_CHECK(hipFree(ptr)); + } + + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + if ((pavail != avail) || (ptot != tot)) { + UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." << + "Possible memory leak."); + TestPassed &= false; + } + + return TestPassed; +} + +/** + * Validates data consistency on supplied gpu + * In Multithreaded Environment + */ +static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + size_t prevAvl, prevTot, curAvl, curTot; + bool TestPassed = true; + constexpr auto N = 4 * 1024 * 1024; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + size_t Nbytes = N * sizeof(int); + HIPCHECK(hipSetDevice(gpu)); + HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, N); + + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) { + UNSCOPED_INFO("Validation PASSED for gpu " << gpu); + } else { + UNSCOPED_INFO("Validation FAILED for gpu " << gpu); + TestPassed = false; + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIPCHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { + // In concurrent calls on one GPU, we cannot verify leaking in this way + UNSCOPED_INFO( + "validateMemoryOnGpuMThread : Memory allocation mismatch observed." + "Possible memory leak."); + TestPassed = false; + } + + return TestPassed; +} + +/** + * Regress memory allocation and free in loop + * In Multithreaded Environment + */ +static bool regressAllocInLoopMthread(int gpu) { + bool TestPassed = true; + size_t tot, avail, ptot, pavail, numBytes; + int i = 0; + int *ptr; + + HIPCHECK(hipSetDevice(gpu)); + numBytes = BuffSizeBC; + + // Exercise allocation in loop with bigger chunks + for (i = 0; i < MaxAllocFree_BigChunks; i++) { + HIPCHECK(hipMemGetInfo(&pavail, &ptot)); + HIPCHECK(hipMalloc(&ptr, numBytes)); + HIPCHECK(hipMemGetInfo(&avail, &tot)); + HIPCHECK(hipFree(ptr)); + + if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes + UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << + numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" << + pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" << + tot << ", pavail-avail=" << pavail-avail); + TestPassed = false; + break; + } + } + + // Exercise allocation in loop with smaller chunks and maximum iters + HIPCHECK(hipMemGetInfo(&pavail, &ptot)); + numBytes = BuffSizeSC; + + for (i = 0; i < MaxAllocFree_SmallChunks; i++) { + HIPCHECK(hipMalloc(&ptr, numBytes)); + + HIPCHECK(hipFree(ptr)); + } + + HIPCHECK(hipMemGetInfo(&avail, &tot)); + + if ((pavail != avail) || (ptot != tot)) { + UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." << + "Possible memory leak."); + TestPassed &= false; + } + + return TestPassed; +} + +/* + * Thread func to regress alloc and check data consistency + */ +static void threadFunc(int gpu) { + g_thTestPassed = regressAllocInLoopMthread(gpu); + g_thTestPassed = g_thTestPassed & validateMemoryOnGpuMThread(gpu); + + UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << + g_thTestPassed.load()); +} + + +/* Performs Argument Validation of api */ +TEST_CASE("Unit_hipMalloc_ArgumentValidation") { + int *ptr; + hipError_t ret; + + SECTION("hipMalloc() when size(0)") { + HIP_CHECK(hipMalloc(&ptr, 0)); + // ptr expected to be reset to null ptr + REQUIRE(ptr == nullptr); + } + + SECTION("hipFree() when freeing nullptr ") { + ptr = nullptr; + // api should return success and shudnt crash + HIP_CHECK(hipFree(ptr)); + } + + SECTION("hipMalloc() with invalid argument") { + constexpr auto sizeBytes = 100; + ret = hipMalloc(nullptr, sizeBytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMalloc() with max size_t") { + ret = hipMalloc(&ptr, std::numeric_limits::max()); + REQUIRE(ret != hipSuccess); + } +} + +/** + * Regress hipMalloc()/hipFree() in loop for bigger chunks and + * smaller chunks of memory allocation + */ +TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") { + int devCnt = 0; + + // Get GPU count + HIP_CHECK(hipGetDeviceCount(&devCnt)); + REQUIRE(devCnt > 0); + + CHECK(regressAllocInLoop(0) == true); + CHECK(validateMemoryOnGPU(0) == true); +} + +/** + * Application Behavior Modelling. + * Check hipMalloc() api adaptability when app creates small chunks of memory + * continuously, stores it for later use and then frees it at later point + * of time. + */ +TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { + size_t avail, tot, pavail, ptot; + bool ret; + hipError_t err; + std::vector ptrlist; + constexpr auto BuffSize = 10; + int devCnt, *ptr; + + // Get GPU count + HIP_CHECK(hipGetDeviceCount(&devCnt)); + REQUIRE(devCnt > 0); + + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + + // Allocate small chunks of memory million times + for (int i = 0; i < MaxAllocPoolIter ; i++) { + if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) { + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + INFO("Loop regression pool allocation failure. " << + "Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " << + avail/(1024.0*1024.0) << " iter " << i << " error " + << hipGetErrorString(err)); + + REQUIRE(false); + } + + // Store pointers allocated to emulate memory pool of app + ptrlist.push_back(ptr); + } + + // Free ptrs at later point of time + for ( auto &t : ptrlist ) { + HIP_CHECK(hipFree(t)); + } + + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + ret = validateMemoryOnGPU(0); + REQUIRE(ret == true); + REQUIRE(pavail == avail); + REQUIRE(ptot == tot); +} + + +/** + * Exercise hipMalloc() api parellely on all gpus from + * multiple threads and regress the api. + */ +TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") { + std::vector threadlist; + int devCnt; + + // Get GPU count + HIP_CHECK(hipGetDeviceCount(&devCnt)); + REQUIRE(devCnt > 0); + + for (int i = 0; i < devCnt; i++) { + threadlist.push_back(std::thread(threadFunc, i)); + } + + for (auto &t : threadlist) { + t.join(); + } + + REQUIRE(g_thTestPassed == true); +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemset3DFunctional.cc b/projects/hip/tests/catch/unit/memory/hipMemset3DFunctional.cc new file mode 100644 index 0000000000..421e066ab4 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemset3DFunctional.cc @@ -0,0 +1,515 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + + 1) Passing width as 0 in extent, verify hipMemset3D api returns success and + doesn't modify the buffer passed. + 2) Passing width as 0 in extent, verify hipMemset3DAsync api returns success + and doesn't modify the buffer passed. + + 3) Passing height as 0 in extent, verify hipMemset3D api returns success and + doesn't modify the buffer passed. + 4) Passing height as 0 in extent, verify hipMemset3DAsync api returns success + and doesn't modify the buffer passed. + + 5) Passing depth as 0 in extent, verify hipMemset3D api returns success and + doesn't modify the buffer passed. + 6) Passing depth as 0 in extent, verify hipMemset3DAsync api returns success + and doesn't modify the buffer passed. + + 7) When extent passed with width, height and depth all as zeroes, verify + hipMemset3D api returns success and doesn't modify the buffer passed. + 8) When extent passed with width, height and depth all as zeroes, verify + hipMemset3DAsync api returns success and doesn't modify the buffer passed. + + 9) Validate data after performing memory set operation with max memset value + for hipMemset3D api. + 10) Validate data after performing memory set operation with max memset value + for hipMemset3DAsync api. + + 11) Select random slice of 3d array and Memset complete slice with + hipMemset3D api. + 12) Select random slice of 3d array and Memset complete slice with + hipMemset3DAsync api. + + 13) Seek device pitched ptr to desired portion of 3d array and memset the + portion with hipMemset3D api. + 14) Seek device pitched ptr to desired portion of 3d array and memset the + portion with hipMemset3DAsync api. +*/ + +#include + +/* + * Defines + */ +#define MEMSETVAL 1 +#define TESTVAL 2 +#define NUMH_EXT 256 +#define NUMW_EXT 100 +#define DEPTH_EXT 10 +#define NUMH_MAX 256 +#define NUMW_MAX 256 +#define DEPTH_MAX 10 +#define ZSIZE_S 32 +#define YSIZE_S 32 +#define XSIZE_S 32 +#define ZSIZE_P 30 +#define YSIZE_P 30 +#define XSIZE_P 30 +#define ZPOS_START 10 +#define ZSET_LEN 10 +#define ZPOS_END 19 +#define YPOS_START 10 +#define YSET_LEN 10 +#define YPOS_END 19 +#define XPOS_START 10 +#define XSET_LEN 10 +#define XPOS_END 19 + + +/** + * Memset with extent passed and verify data to be intact + */ +static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) { + hipPitchedPtr devPitchedPtr; + hipError_t ret; + char *A_h; + size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + memset(A_h, 0, sizeElements); + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + if (bAsync) { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height + << "," << extent.depth << ") memset " + << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + + ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height + << "," << tstExtent.depth << ") memset " + << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + ret = hipMemset3D(devPitchedPtr, MEMSETVAL, extent); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height + << "," << extent.depth << ") memset " + << MEMSETVAL << ",ret : " << ret); + REQUIRE(ret == hipSuccess); + + ret = hipMemset3D(devPitchedPtr, TESTVAL, tstExtent); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height + << "," << tstExtent.depth << ") memset " + << TESTVAL << ",ret : " << ret); + REQUIRE(ret == hipSuccess); + } + + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (size_t i = 0; i < elements; i++) { + if (A_h[i] != MEMSETVAL) { + INFO("testMemsetWithExtent: index:" << i << ",computed:" + << std::hex << static_cast(A_h[i]) << ",memsetval:" + << std::hex << MEMSETVAL); + REQUIRE(false); + } + } + + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); +} + + +/** + * Validates data after performing memory set operation with max memset value + */ +static void testMemsetMaxValue(bool bAsync) { + hipPitchedPtr devPitchedPtr; + unsigned char *A_h; + int memsetval = std::numeric_limits::max(); + size_t numH = NUMH_MAX, numW = NUMW_MAX, depth = DEPTH_MAX; + size_t width = numW * sizeof(unsigned char); + hipExtent extent = make_hipExtent(width, numH, depth); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + + A_h = reinterpret_cast (malloc(sizeElements)); + REQUIRE(A_h != nullptr); + memset(A_h, 0, sizeElements); + + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + if (bAsync) { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + } + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (size_t i = 0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("testMemsetMaxValue: index:" << i << ",computed:" + << std::hex << static_cast(A_h[i]) << ",memsetval:" + << std::hex << memsetval); + REQUIRE(false); + } + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); +} + +/** + * Function seeks device ptr to random slice and performs Memset operation + * on the slice selected. + */ +static void seekAndSet3DArraySlice(bool bAsync) { + char array3D[ZSIZE_S][YSIZE_S][XSIZE_S]{}; + dim3 arr_dimensions = dim3(ZSIZE_S, YSIZE_S, XSIZE_S); + hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x, + arr_dimensions.y, arr_dimensions.z); + hipPitchedPtr devicePitchedPointer; + int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL; + + HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent)); + HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); + + // select random slice for memset + unsigned int seed = time(nullptr); + int slice_index = rand_r(&seed) % ZSIZE_S; + + INFO("memset3d for sliceindex " << slice_index); + + // Get attributes from device pitched pointer + size_t pitch = devicePitchedPointer.pitch; + size_t slicePitch = pitch * extent.height; + + // Point devptr to selected slice + char *devPtrSlice = (reinterpret_cast(devicePitchedPointer.ptr)) + + slice_index * slicePitch; + hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x, + arr_dimensions.y, 1); + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrSlice, pitch, + arr_dimensions.x, arr_dimensions.y); + + if (bAsync) { + // Memset selected slice (Async) + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, + extentSlice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + // Memset selected slice + HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice)); + } + + // Copy result back to host buffer + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, + arr_dimensions.x, arr_dimensions.y); + myparms.srcPtr = devicePitchedPointer; + myparms.extent = extent; +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (int z = 0; z < ZSIZE_S; z++) { + for (int y = 0; y < YSIZE_S; y++) { + for (int x = 0; x < XSIZE_S; x++) { + if (z == slice_index) { + if (array3D[z][y][x] != memsetval4seeked) { + INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z + << "," << y << "," << x << ") " << "computed:" << std::hex + << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval4seeked); + REQUIRE(false); + } + } else { + if (array3D[z][y][x] != memsetval) { + INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z + << "," << y << "," << x << ") " << "computed:" << std::hex + << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval); + REQUIRE(false); + } + } + } + } + } + + HIP_CHECK(hipFree(devicePitchedPointer.ptr)); +} + +/** + * Function seeks device ptr to selected portion of 3d array + * and performs Memset operation on the portion. + */ +static void seekAndSet3DArrayPortion(bool bAsync) { + char array3D[ZSIZE_P][YSIZE_P][XSIZE_P]{}; + dim3 arr_dimensions = dim3(ZSIZE_P, YSIZE_P, XSIZE_P); + hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x, + arr_dimensions.y, arr_dimensions.z); + hipPitchedPtr devicePitchedPointer; + int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL; + + HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent)); + HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); + + // For memsetting extent/size(10,10,10) in the mid portion of cube(30,30,30), + // seek device ptr to (10,10,10) and then memset 10 bytes across x,y,z axis. + size_t pitch = devicePitchedPointer.pitch; + size_t slicePitch = pitch * extent.height; + int slice_index = ZPOS_START, y = YPOS_START, x = XPOS_START; + + // Select 10th slice + char *devPtrSlice = (reinterpret_cast(devicePitchedPointer.ptr)) + + slice_index * slicePitch; + + // Now select row at height as 10 + char *current_row = reinterpret_cast(devPtrSlice + y * pitch); + + // Now select index of selected row as 10 + char *devPtrIndexed = ¤t_row[x]; + + // Make dev Pitchedptr, extent + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrIndexed, pitch, + arr_dimensions.x, arr_dimensions.y); + hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN, + ZSET_LEN); + + if (bAsync) { + // Memset selected portion (Async) + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, + setExtent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + // Memset selected portion + HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent)); + } + + // Copy result back to host buffer + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, + arr_dimensions.x, arr_dimensions.y); + myparms.srcPtr = devicePitchedPointer; + myparms.extent = extent; +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (int z = 0; z < ZSIZE_P; z++) { + for (int y = 0; y < YSIZE_P; y++) { + for (int x = 0; x < XSIZE_P; x++) { + if ((z >= ZPOS_START && z <= ZPOS_END) && + (y >= YPOS_START && y <= YPOS_END) && + (x >= XPOS_START && x <= XPOS_END)) { + if (array3D[z][y][x] != memsetval4seeked) { + INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z + << "," << y << "," << x << ") " << "computed:" << std::hex + << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval4seeked); + REQUIRE(false); + } + } else { + if (array3D[z][y][x] != memsetval) { + INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z + << "," << y << "," << x << ") " << "computed:" << std::hex + << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval); + REQUIRE(false); + } + } + } + } + } + + HIP_CHECK(hipFree(devicePitchedPointer.ptr)); +} + + + +/** + * Test Memset3D with different combinations of extent + * taking zero and non-zero fields. + */ +TEST_CASE("Unit_hipMemset3D_MemsetWithExtent") { + hipExtent testExtent; + size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; + + SECTION("Memset with extent width(0)") { + // Memset with extent width(0) and verify data to be intact + testExtent = make_hipExtent(0, numH, depth); + testMemsetWithExtent(0, testExtent); + } + + SECTION("Memset with extent height(0)") { + // Memset with extent height(0) and verify data to be intact + testExtent = make_hipExtent(numW, 0, depth); + testMemsetWithExtent(0, testExtent); + } + + SECTION("Memset with extent depth(0)") { + // Memset with extent depth(0) and verify data to be intact + testExtent = make_hipExtent(numW, numH, 0); + testMemsetWithExtent(0, testExtent); + } + + SECTION("Memset with extent width,height,depth as 0") { + // Memset with extent width,height,depth as 0 and verify data to be intact + testExtent = make_hipExtent(0, 0, 0); + testMemsetWithExtent(0, testExtent); + } +} + + +/** + * Test Memset3DAsync with different combinations of extent + * taking zero and non-zero fields. + */ +TEST_CASE("Unit_hipMemset3DAsync_MemsetWithExtent") { + hipExtent testExtent; + size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; + + SECTION("Memset with extent width(0)") { + // Memset with extent width(0) and verify data to be intact + testExtent = make_hipExtent(0, numH, depth); + testMemsetWithExtent(1, testExtent); + } + + SECTION("Memset with extent height(0)") { + // Memset with extent height(0) and verify data to be intact + testExtent = make_hipExtent(numW, 0, depth); + testMemsetWithExtent(1, testExtent); + } + + SECTION("Memset with extent depth(0)") { + // Memset with extent depth(0) and verify data to be intact + testExtent = make_hipExtent(numW, numH, 0); + testMemsetWithExtent(1, testExtent); + } + + SECTION("Memset with extent width,height,depth as 0") { + // Memset with extent width,height,depth as 0 and verify data to be intact + testExtent = make_hipExtent(0, 0, 0); + testMemsetWithExtent(1, testExtent); + } +} + +/** + * Memset3D with max unsigned char and verify memset operation is success + */ +TEST_CASE("Unit_hipMemset3D_MemsetMaxValue") { + testMemsetMaxValue(0); +} + +/** + * Memset3DAsync with max unsigned char and verify memset operation is success + */ +TEST_CASE("Unit_hipMemset3DAsync_MemsetMaxValue") { + testMemsetMaxValue(1); +} + +/** + * Seek and set random slice of 3d array, verify memset is success + */ +TEST_CASE("Unit_hipMemset3D_SeekSetSlice") { + seekAndSet3DArraySlice(0); +} + +/** + * Seek and set random slice of 3d array with async, verify memset is success + */ +TEST_CASE("Unit_hipMemset3DAsync_SeekSetSlice") { + seekAndSet3DArraySlice(1); +} + +/** + * Memset3D selected portion of 3d array + */ +TEST_CASE("Unit_hipMemset3D_SeekSetArrayPortion") { + seekAndSet3DArrayPortion(0); +} + +/** + * Memset3DAsync selected portion of 3d array + */ +TEST_CASE("Unit_hipMemset3DAsync_SeekSetArrayPortion") { + seekAndSet3DArrayPortion(1); +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc new file mode 100644 index 0000000000..798c90e1ac --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc @@ -0,0 +1,236 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + 1) Test hipMemset3D() with uninitialized devPitchedPtr. + 2) Test hipMemset3DAsync() with uninitialized devPitchedPtr. + + 3) Reset devPitchedPtr to zero and check return value for hipMemset3D(). + 4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync(). + + 5) Test hipMemset3D() with extent.width as max size_t and keeping height, + depth as valid values. + 6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height, + depth as valid values. + 7) Test hipMemset3D() with extent.height as max size_t and keeping width, + depth as valid values. + 8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width, + depth as valid values. + 9) Test hipMemset3D() with extent.depth as max size_t and keeping height, + width as valid values. +10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height, + width as valid values. + +11) Device Ptr out bound and extent(0) passed for hipMemset3D(). +12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync(). + +13) Device Ptr out bound and valid extent passed for hipMemset3D(). +14) Device Ptr out bound and valid extent passed for hipMemset3DAsync(). +*/ + +#include + +TEST_CASE("Unit_hipMemset3D_Negative") { + hipError_t ret; + hipPitchedPtr devPitchedPtr; + constexpr int memsetval = 1; + constexpr size_t numH = 256, numW = 256; + constexpr size_t depth = 10; + constexpr size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + + SECTION("Using uninitialized devpitched ptr") { + hipPitchedPtr devPitchedUnPtr; + + ret = hipMemset3D(devPitchedUnPtr, memsetval, extent); + REQUIRE(ret != hipSuccess); + } + + SECTION("Reset devPitchedPtr to zero") { + hipPitchedPtr rdevPitchedPtr{}; + + ret = hipMemset3D(rdevPitchedPtr, memsetval, extent); + REQUIRE(ret != hipSuccess); + } + + SECTION("Pass extent fields as max size_t") { + hipExtent extMW = make_hipExtent(std::numeric_limits::max(), + numH, + depth); + hipExtent extMH = make_hipExtent(width, + std::numeric_limits::max(), + depth); + hipExtent extMD = make_hipExtent(width, + numH, + std::numeric_limits::max()); + + ret = hipMemset3D(devPitchedPtr, memsetval, extMW); + REQUIRE(ret != hipSuccess); + + ret = hipMemset3D(devPitchedPtr, memsetval, extMH); + REQUIRE(ret != hipSuccess); + + if ((TestContext::get()).isAmd()) { + ret = hipMemset3D(devPitchedPtr, memsetval, extMD); + REQUIRE(ret != hipSuccess); + } else { + WARN("Test is skipped for max depth." + << "Cuda doesn't check the maximum depth of extent field"); + } + } + + SECTION("Device Ptr out bound and extent(0) passed for memset") { + size_t pitch = devPitchedPtr.pitch; + size_t slicePitch = pitch * extent.height; + constexpr auto advanceOffset = 10; + + // Point devptr to end of allocated memory + char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) + + depth * slicePitch; + + // Advance devptr further to go out of boundary + devPtrMod = devPtrMod + advanceOffset; + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, + numW * sizeof(char), numH); + hipExtent extent0{}; + ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0); + + // api expected to check extent0 and return success before going for ptr. + REQUIRE(ret == hipSuccess); + } + + SECTION("Device Ptr out bound and valid extent passed for memset") { + size_t pitch = devPitchedPtr.pitch; + size_t slicePitch = pitch * extent.height; + constexpr auto advanceOffset = 10; + + // Point devptr to end of allocated memory + char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) + + depth * slicePitch; + + // Advance devptr further to go out of boundary + devPtrMod = devPtrMod + advanceOffset; + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, + numW * sizeof(char), numH); + ret = hipMemset3D(modDevPitchedPtr, memsetval, extent); + + REQUIRE(ret != hipSuccess); + } + + HIP_CHECK(hipFree(devPitchedPtr.ptr)); +} + +TEST_CASE("Unit_hipMemset3DAsync_Negative") { + hipError_t ret; + hipPitchedPtr devPitchedPtr; + hipStream_t stream; + constexpr int memsetval = 1; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + constexpr size_t depth = 10; + constexpr size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + + SECTION("Using uninitialized devpitched ptr") { + hipPitchedPtr devPitchedUnPtr; + + ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream); + REQUIRE(ret != hipSuccess); + } + + SECTION("Reset devPitchedPtr to zero") { + hipPitchedPtr rdevPitchedPtr{}; + + ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream); + REQUIRE(ret != hipSuccess); + } + + SECTION("Pass extent fields as max size_t") { + hipExtent extMW = make_hipExtent(std::numeric_limits::max(), + numH, + depth); + hipExtent extMH = make_hipExtent(width, + std::numeric_limits::max(), + depth); + hipExtent extMD = make_hipExtent(width, + numH, + std::numeric_limits::max()); + + ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream); + REQUIRE(ret != hipSuccess); + + ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream); + REQUIRE(ret != hipSuccess); + + if ((TestContext::get()).isAmd()) { + ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream); + REQUIRE(ret != hipSuccess); + } else { + WARN("Test is skipped for max depth." + << "Cuda doesn't check the maximum depth of extent field"); + } + } + + SECTION("Device Ptr out bound and extent(0) passed for memset") { + size_t pitch = devPitchedPtr.pitch; + size_t slicePitch = pitch * extent.height; + constexpr auto advanceOffset = 10; + + // Point devptr to end of allocated memory + char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) + + depth * slicePitch; + + // Advance devptr further to go out of boundary + devPtrMod = devPtrMod + advanceOffset; + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, + numW * sizeof(char), numH); + hipExtent extent0{}; + ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream); + + // api expected to check extent0 and return success before going for ptr. + REQUIRE(ret == hipSuccess); + } + + SECTION("Device Ptr out bound and valid extent passed for memset") { + size_t pitch = devPitchedPtr.pitch; + size_t slicePitch = pitch * extent.height; + constexpr auto advanceOffset = 10; + + // Point devptr to end of allocated memory + char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) + + depth * slicePitch; + + // Advance devptr further to go out of boundary + devPtrMod = devPtrMod + advanceOffset; + hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, + numW * sizeof(char), numH); + ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream); + + REQUIRE(ret != hipSuccess); + } + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(devPitchedPtr.ptr)); +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemset3DRegressMultiThread.cc b/projects/hip/tests/catch/unit/memory/hipMemset3DRegressMultiThread.cc new file mode 100644 index 0000000000..b3a65f2104 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemset3DRegressMultiThread.cc @@ -0,0 +1,266 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + 1) Validate Async behavior of hipMemset3DAsync with commands queued + concurrently from multiple threads. + 2) Validate hipMemset3DAsync behavior when api is queued along with kernel + function operating on same memory. + 3) Perform regression of hipMemset3D api in loop with device memory allocated + on different gpus. + 4) Perform regression of hipMemset3DAsync api in loop with device memory + allocated on different gpus. +*/ + +#include + + +/* + * Defines + */ +#define MAX_REGRESS_ITERS 2 +#define MAX_THREADS 10 + +/** + * kernel function sets device memory with value passed + */ +static __global__ void func_set_value(hipPitchedPtr devicePitchedPointer, + hipExtent extent, + unsigned char val) { + // Index Calculation + size_t x = threadIdx.x + blockDim.x * blockIdx.x; + size_t y = threadIdx.y + blockDim.y * blockIdx.y; + size_t z = threadIdx.z + blockDim.z * blockIdx.z; + + // Get attributes from device pitched pointer + char *devicePointer = reinterpret_cast(devicePitchedPointer.ptr); + size_t pitch = devicePitchedPointer.pitch; + size_t slicePitch = pitch * extent.height; + + // Loop over the device buffer + if (z < extent.depth) { + char *current_slice_index = devicePointer + z * slicePitch; + if (y < extent.height) { + // Get data array containing all elements from the current row + char *current_row = reinterpret_cast(current_slice_index + + y * pitch); + if (x < extent.width) { + current_row[x] = val; + } + } + } +} + + +/** + * Thread function queues kernel function and memset cmds + */ +static void threadFunc(hipStream_t stream, hipPitchedPtr devpPtr, + int memsetval, int testval, hipExtent extent, hipMemcpy3DParms myparms) { + // Kernel Launch Configuration + constexpr auto size = 8; + dim3 threadsPerBlock = dim3(size, size, size); + dim3 blocks; + blocks = dim3((extent.width + threadsPerBlock.x - 1) / threadsPerBlock.x, + (extent.height + threadsPerBlock.y - 1) / threadsPerBlock.y, + (extent.depth + threadsPerBlock.z - 1) / threadsPerBlock.z); + + hipLaunchKernelGGL(func_set_value, dim3(blocks), dim3(threadsPerBlock), 0, + stream, devpPtr, extent, memsetval); + HIPCHECK(hipMemset3DAsync(devpPtr, testval, extent, stream)); + HIPCHECK(hipMemcpy3DAsync(&myparms, stream)); +} + + +/** + * Performs api regression in loop + */ +bool loopRegression(bool bAsync) { + bool testPassed = true; + char *A_h; + constexpr int memsetval = 1; + constexpr size_t numH = 256, numW = 100, depth = 10; + int numGpu = 0, hasPeerAccess = 0; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + std::vector devPitchedPtrlist; + hipPitchedPtr pitchedPtr, devpPtr; + + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + memset(A_h, 0, sizeElements); + + // Populate hipMemcpy3D parameters + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.extent = extent; + +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + HIP_CHECK(hipGetDeviceCount(&numGpu)); + REQUIRE(numGpu > 0); + + // Alloc 3D arrays in all GPUs + for (int j = 0; j < numGpu; j++) { + HIP_CHECK(hipSetDevice(j)); + HIP_CHECK(hipMalloc3D(&pitchedPtr, extent)); + devPitchedPtrlist.push_back(pitchedPtr); + } + + for (int itern = 0; itern < MAX_REGRESS_ITERS; itern++) { + // Validate hipMemset3D data consistency in multiple iters + for (int i = 0; i < numGpu; i++) { + for (int j = 0; j < numGpu; j++) { + HIP_CHECK(hipDeviceCanAccessPeer(&hasPeerAccess, i, j)); + if (!hasPeerAccess) { + // Skip and continue if no peer access + continue; + } + + HIP_CHECK(hipSetDevice(i)); + devpPtr = devPitchedPtrlist[j]; + HIP_CHECK(hipDeviceEnablePeerAccess(j, 0)); + HIP_CHECK(hipMemset3D(devpPtr, 0, extent)); + + if (bAsync) { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(devpPtr, memsetval, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + HIP_CHECK(hipMemset3D(devpPtr, memsetval, extent)); + } + + myparms.srcPtr = devpPtr; + memset(A_h, 0, sizeElements); + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (size_t indx = 0; indx < elements; indx++) { + if (A_h[indx] != memsetval) { + testPassed = false; + printf("RegressIter : mismatch at index:%d computed:%02x, " + "memsetval:%02x\n", static_cast(indx), + static_cast(A_h[indx]), static_cast(memsetval)); + break; + } + } + } + } + } + + for (int j = 0; j < numGpu; j++) { + HIP_CHECK(hipFree(devPitchedPtrlist[j].ptr)); + } + + free(A_h); + return testPassed; +} + +/** + * Perform regression of hipMemset3D api with device memory allocated + * on different gpus. + */ +TEST_CASE("Unit_hipMemset3D_RegressInLoop") { + bool TestPassed = false; + + TestPassed = loopRegression(0); + REQUIRE(TestPassed == true); +} + +/** + * Perform regression of hipMemset3DAsync api with device memory allocated + * on different gpus. + */ +TEST_CASE("Unit_hipMemset3DAsync_RegressInLoop") { + bool TestPassed = false; + + TestPassed = loopRegression(1); + REQUIRE(TestPassed == true); +} + +/** + * Async commands queued concurrently and executed + */ +TEST_CASE("Unit_hipMemset3DAsync_ConcurrencyMthread") { + char *A_h; + constexpr int memsetval = 1, testval = 2; + constexpr size_t numH = 256, numW = 100, depth = 10; + size_t width = numW * sizeof(char); + hipExtent extent = make_hipExtent(width, numH, depth); + size_t sizeElements = width * numH * depth; + size_t elements = numW* numH* depth; + hipPitchedPtr devpPtr; + hipStream_t stream; + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMalloc3D(&devpPtr, extent)); + + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + memset(A_h, 0, sizeElements); + + // Populate hipMemcpy3D parameters + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.srcPtr = devpPtr; + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.extent = extent; + +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + + std::vector threadlist; + + // Queue cmds concurrently from multiple threads on same stream + for (int i = 0; i < MAX_THREADS; i++) { + threadlist.push_back(std::thread(threadFunc, stream, devpPtr, memsetval, + testval, extent, myparms)); + } + + for (auto &t : threadlist) { + t.join(); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + + for (size_t k = 0 ; k < elements ; k++) { + if (A_h[k] != testval) { + CAPTURE(A_h[k], testval, k); + REQUIRE(false); + } + } + + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + HIP_CHECK(hipFree(devpPtr.ptr)); +}