diff --git a/docs/markdown/hip_debugging.md b/docs/markdown/hip_debugging.md index 1ebb83352e..8215db17c8 100644 --- a/docs/markdown/hip_debugging.md +++ b/docs/markdown/hip_debugging.md @@ -127,7 +127,7 @@ Breakpoint 1, main () ``` ### Other Debugging Tools -There are also other debugging tools availble online developers can google and choose the one best suits the debugging requirements. +There are also other debugging tools available online developers can google and choose the one best suits the debugging requirements. ## Debugging HIP Applications @@ -192,7 +192,7 @@ AMD_SERIALIZE_COPY, for serializing copies. So HIP runtime can wait for GPU idle before/after any GPU command depending on the environment setting. ### Making Device visible -For system with multiple devices, it's possible to make only certain device(s) visible to HIP via setting environment varible, +For system with multiple devices, it's possible to make only certain device(s) visible to HIP via setting environment variable, HIP_VISIBLE_DEVICES, only devices whose index is present in the sequence are visible to HIP. For example, @@ -200,7 +200,7 @@ For example, $ HIP_VISIBLE_DEVICES=0,1 ``` -or in the appliation, +or in the application, ``` if (totalDeviceNum > 2) { setenv("HIP_VISIBLE_DEVICES", "0,1,2", 1); @@ -210,11 +210,11 @@ if (totalDeviceNum > 2) { ``` ### Dump code object -Developers can dump code object to anylize compiler related issues via setting environment variable, +Developers can dump code object to analyze compiler related issues via setting environment variable, GPU_DUMP_CODE_OBJECT ### HSA related environment variables -HSA provides some environment varibles help to analize issues in driver or hardware, for example, +HSA provides some environment variables help to analyze issues in driver or hardware, for example, HSA_ENABLE_SDMA=0 It causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. @@ -225,8 +225,24 @@ HSA_ENABLE_INTERRUPT=0 Causes completion signals to be detected with memory-based polling rather than interrupts. This environment variable can be useful to diagnose interrupt storm issues in the driver. +### Summary of environment variables in HIP + +The following is the summary of the most useful environment variables in HIP. + +| **Environment variable** | **Default value** | **Usage** | +| ---------------------------------------------------------------------------------------------------------------| ----------------- | --------- | +| AMD_LOG_LEVEL
Enable HIP log on different Level. | 0 | 0: Disable log.
1: Enable log on error level.
2: Enable log on warning and below levels.
0x3: Enable log on information and below levels.
0x4: Decode and display AQL packets. | +| AMD_LOG_MASK
Enable HIP log on different Level. | 0x7FFFFFFF | 0x1: Log API calls.
0x02: Kernel and Copy Commands and Barriers.
0x4: Synchronization and waiting for commands to finish.
0x8: Enable log on information and below levels.
0x20: Queue commands and queue contents.
0x40:Signal creation, allocation, pool.
0x80: Locks and thread-safety code.
0x100: Copy debug.
0x200: Detailed copy debug.
0x400: Resource allocation, performance-impacting events.
0x800: Initialization and shutdown.
0x1000: Misc debug, not yet classified.
0x2000: Show raw bytes of AQL packet.
0x4000: Show code creation debug.
0x8000: More detailed command info, including barrier commands.
0x10000: Log message location.
0xFFFFFFFF: Log always even mask flag is zero. | +| HIP_VISIBLE_DEVICES
Only devices whose index is present in the sequence are visible to HIP. | | 0,1,2: Depending on the number of devices on the system. | +| GPU_DUMP_CODE_OBJECT
Dump code object. | 0 | 0: Disable.
1: Enable. | +| AMD_SERIALIZE_KERNEL
Serialize kernel enqueue. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | +| AMD_SERIALIZE_COPY
Serialize copies. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | +| HIP_HOST_COHERENT
Coherent memory in hipHostMalloc. | 0 | 0: memory is not coherent between host and GPU.
1: memory is coherent with host. | +| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 0 | 0: Disable.
1: Enable. | + + ## General Debugging Tips -- 'gdb --args' can be used to conviently pass the executable and arguments to gdb. +- 'gdb --args' can be used to conveniently pass the executable and arguments to gdb. - From inside GDB, you can set environment variables "set env". Note the command does not use an '=' sign: ``` diff --git a/docs/markdown/hip_programming_guide.md b/docs/markdown/hip_programming_guide.md index 08a2374aa8..94183bb6f9 100644 --- a/docs/markdown/hip_programming_guide.md +++ b/docs/markdown/hip_programming_guide.md @@ -66,12 +66,15 @@ ROCm defines two coherency options for host memory: - Non-coherent memory : Can be cached by GPU, but cannot support synchronization while the kernel is running.  Non-coherent memory can be optionally synchronized only at command (end-of-kernel or copy command) boundaries.  This memory is appropriate for high-performance access when fine-grain synchronization is not required. HIP provides the developer with controls to select which type of memory is used via allocation flags passed to hipHostMalloc and the HIP_HOST_COHERENT environment variable. By default, the environment variable HIP_HOST_COHERENT is set to 0 in HIP. -- hipHostMallocCoherent=0, hipHostMallocNonCoherent=0: Use HIP_HOST_COHERENT environment variable, +The control logic in the current version of HIP is as follows: +- No flags are passed in: the host memory allocation is coherent, the HIP_HOST_COHERENT environment variable is ignored. +- hipHostMallocCoherent=1: The host memory allocation will be coherent, the HIP_HOST_COHERENT environment variable is ignored. +- hipHostMallocMapped=1: The host memory allocation will be coherent, the HIP_HOST_COHERENT environment variable is ignored. +- hipHostMallocNonCoherent=1, hipHostMallocCoherent=0, and hipHostMallocMapped=0: The host memory will be non-coherent, the HIP_HOST_COHERENT environment variable is ignored. +- hipHostMallocCoherent=0, hipHostMallocNonCoherent=0, hipHostMallocMapped=0, but one of the other HostMalloc flags is set: - If HIP_HOST_COHERENT is defined as 1, the host memory allocation is coherent. - If HIP_HOST_COHERENT is not defined, or defined as 0, the host memory allocation is non-coherent. -- hipHostMallocCoherent=1, hipHostMallocNonCoherent=0: The host memory allocation will be coherent.  HIP_HOST_COHERENT env variable is ignored. -- hipHostMallocCoherent=0, hipHostMallocNonCoherent=1: The host memory allocation will be non-coherent.  HIP_HOST_COHERENT env variable is ignored. -- hipHostMallocCoherent=1, hipHostMallocNonCoherent=1: Illegal. +- hipHostMallocCoherent=1, hipHostMallocNonCoherent=1: Illegal. ### Visibility of Zero-Copy Host Memory Coherent host memory is automatically visible at synchronization points. diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index 0fddc8f346..434dc92b2b 100644 --- a/tests/catch/CMakeLists.txt +++ b/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/tests/catch/TypeQualifiers/CMakeLists.txt b/tests/catch/TypeQualifiers/CMakeLists.txt new file mode 100644 index 0000000000..3c60bdf411 --- /dev/null +++ b/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/tests/catch/unit/memory/hipManagedKeyword.cc b/tests/catch/TypeQualifiers/hipManagedKeyword.cc similarity index 100% rename from tests/catch/unit/memory/hipManagedKeyword.cc rename to tests/catch/TypeQualifiers/hipManagedKeyword.cc diff --git a/tests/catch/hipTestMain/CMakeLists.txt b/tests/catch/hipTestMain/CMakeLists.txt index 2b3cb3f60c..a70169efb0 100644 --- a/tests/catch/hipTestMain/CMakeLists.txt +++ b/tests/catch/hipTestMain/CMakeLists.txt @@ -69,3 +69,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/tests/catch/include/hip_test_helper.hh b/tests/catch/include/hip_test_helper.hh index 1d0a57b776..0ea01021e3 100644 --- a/tests/catch/include/hip_test_helper.hh +++ b/tests/catch/include/hip_test_helper.hh @@ -23,10 +23,46 @@ THE SOFTWARE. #pragma once #include "hip_test_common.hh" +#ifdef __linux__ +#include +#endif + namespace HipTest { static inline int getGeviceCount() { int dev = 0; - HIPCHECK(hipGetDeviceCount(&dev)); + HIP_CHECK(hipGetDeviceCount(&dev)); return dev; } + +// Get Free Memory from the system +static size_t getMemoryAmount() { +#ifdef __linux__ + struct sysinfo info{}; + sysinfo(&info); + return info.freeram / (1024 * 1024); // MB +#elif defined(_WIN32) + MEMORYSTATUSEX statex; + statex.dwLength = sizeof(statex); + GlobalMemoryStatusEx(&statex); + return (statex.ullAvailPhys / (1024 * 1024)); // MB +#endif +} + +static size_t getHostThreadCount(const size_t memPerThread, + const size_t maxThreads) { + if (memPerThread == 0) return 0; + auto memAmount = getMemoryAmount(); + const auto processor_count = std::thread::hardware_concurrency(); + if (processor_count == 0 || memAmount == 0) return 0; + size_t thread_count = 0; + if ((processor_count * memPerThread) < memAmount) + thread_count = processor_count; + else + thread_count = reinterpret_cast(memAmount / memPerThread); + if (maxThreads > 0) { + return (thread_count > maxThreads) ? maxThreads : thread_count; + } + return thread_count; +} + } // namespace HipTest diff --git a/tests/catch/include/hip_test_smi.hh b/tests/catch/include/hip_test_smi.hh new file mode 100644 index 0000000000..8962d4f230 --- /dev/null +++ b/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/tests/catch/multiproc/CMakeLists.txt b/tests/catch/multiproc/CMakeLists.txt index 4d2c69b114..3e19d38ef9 100644 --- a/tests/catch/multiproc/CMakeLists.txt +++ b/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/tests/catch/multiproc/hipHostMallocTestsMproc.cc b/tests/catch/multiproc/hipHostMallocTestsMproc.cc new file mode 100644 index 0000000000..cca3c396aa --- /dev/null +++ b/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/tests/catch/multiproc/hipMallocConcurrency.cc b/tests/catch/multiproc/hipMallocConcurrency.cc deleted file mode 100644 index f6520403d5..0000000000 --- a/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/tests/catch/multiproc/hipMallocConcurrencyMproc.cc b/tests/catch/multiproc/hipMallocConcurrencyMproc.cc new file mode 100644 index 0000000000..9018f694a2 --- /dev/null +++ b/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/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index d9e688d167..a4619410b6 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -32,7 +32,18 @@ set(TEST_SRC hipHostGetFlags.cc hipMemoryAllocateCoherent.cc hipMallocManaged_MultiScenario.cc - hipManagedKeyword.cc + hipMemsetInvalidPtr.cc + hipMemset.cc + hipMemsetAsyncMultiThread.cc + hipMemsetAsyncAndKernel.cc + hipMemset3D.cc + hipMemset2D.cc + hipMemset2DAsyncMultiThreadAndKernel.cc + hipHostMallocTests.cc + hipMallocConcurrency.cc + hipMemset3DFunctional.cc + hipMemset3DNegative.cc + hipMemset3DRegressMultiThread.cc ) else() set(TEST_SRC @@ -65,7 +76,18 @@ set(TEST_SRC hipHostGetFlags.cc hipMemoryAllocateCoherent.cc hipMallocManaged_MultiScenario.cc - hipManagedKeyword.cc + hipMemsetInvalidPtr.cc + hipMemset.cc + hipMemsetAsyncMultiThread.cc + hipMemsetAsyncAndKernel.cc + 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/tests/catch/unit/memory/hipHostMallocTests.cc b/tests/catch/unit/memory/hipHostMallocTests.cc new file mode 100644 index 0000000000..de445e2e24 --- /dev/null +++ b/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/tests/catch/unit/memory/hipMallocConcurrency.cc b/tests/catch/unit/memory/hipMallocConcurrency.cc new file mode 100644 index 0000000000..5083ea1c25 --- /dev/null +++ b/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/tests/catch/unit/memory/hipMemset.cc b/tests/catch/unit/memory/hipMemset.cc new file mode 100644 index 0000000000..600a008436 --- /dev/null +++ b/tests/catch/unit/memory/hipMemset.cc @@ -0,0 +1,281 @@ +/* + * 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 hipMemset small size buffers with unique memset values. + 2) Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique + number of elements and memset values. + 3) Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async + apis with unique number of elements and memset values. + 4) Test two memset async operations at the same time. +*/ + + +#include + + +// Table with unique number of elements and memset values. +// (N, memsetval, memsetD32val, memsetD16val, memsetD8val) +typedef std::tuple tupletype; +static constexpr std::initializer_list tableItems { + std::make_tuple((4*1024*1024), 0x42, 0xDEADBEEF, 0xDEAD, 0xDE), + std::make_tuple((10) , 0x42, 0x101 , 0x10, 0x1), + std::make_tuple((10013) , 0x5a, 0xDEADBEEF, 0xDEAD, 0xDE), + std::make_tuple((256*1024*1024), 0xa6, 0xCAFEBABE, 0xCAFE, 0xCA) + }; + +enum MemsetType { + hipMemsetTypeDefault, + hipMemsetTypeD8, + hipMemsetTypeD16, + hipMemsetTypeD32 +}; + +template +static bool testhipMemset(T *A_h, T *A_d, T memsetval, enum MemsetType type, + size_t numElements) { + size_t Nbytes = numElements * sizeof(T); + bool testResult = true; + constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr. + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + A_h = reinterpret_cast (malloc(Nbytes)); + REQUIRE(A_h != nullptr); + + for (int offset = MAX_OFFSET; offset >= 0; offset --) { + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset(A_d + offset, memsetval, numElements - offset)); + + } else if (type == hipMemsetTypeD8) { + HIP_CHECK(hipMemsetD8((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset)); + + } else if (type == hipMemsetTypeD16) { + HIP_CHECK(hipMemsetD16((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset)); + + } else if (type == hipMemsetTypeD32) { + HIP_CHECK(hipMemsetD32((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset)); + } + + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = offset; i < numElements; i++) { + if (A_h[i] != memsetval) { + testResult = false; + CAPTURE(i, A_h[i], memsetval); + break; + } + } + } + + HIP_CHECK(hipFree(A_d)); + free(A_h); + return testResult; +} + + +template +static bool testhipMemsetAsync(T *A_h, T *A_d, T memsetval, + enum MemsetType type, size_t numElements) { + size_t Nbytes = numElements * sizeof(T); + bool testResult = true; + constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr. + hipStream_t stream; + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + A_h = reinterpret_cast (malloc(Nbytes)); + REQUIRE(A_h != nullptr); + + for (int offset = MAX_OFFSET; offset >= 0; offset --) { + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset, + stream)); + + } else if (type == hipMemsetTypeD8) { + HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset, stream)); + + } else if (type == hipMemsetTypeD16) { + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset, stream)); + + } else if (type == hipMemsetTypeD32) { + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)(A_d + offset), memsetval, + numElements - offset, stream)); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = offset; i < numElements; i++) { + if (A_h[i] != memsetval) { + testResult = false; + CAPTURE(i, A_h[i], memsetval); + break; + } + } + } + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + return testResult; +} + + +/** + * Test hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32 apis with unique + * number of elements and memset values. + */ +TEST_CASE("Unit_hipMemset_SetMemoryWithOffset") { + char memsetval; + int memsetD32val; + int16_t memsetD16val; + char memsetD8val; + size_t N; + bool ret; + + std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) = + GENERATE(table(tableItems)); + + + SECTION("Memset with hipMemsetTypeDefault") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemset(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD32") { + int32_t *iA_d{nullptr}, *iA_h{nullptr}; + ret = testhipMemset(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD16") { + int16_t *siA_d{nullptr}, *siA_h{nullptr}; + ret = testhipMemset(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD8") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemset(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N); + REQUIRE(ret == true); + } +} + + +/** + * Test hipMemsetAsync, hipMemsetD8Async, hipMemsetD16Async, hipMemsetD32Async + * apis with unique number of elements and memset values. + */ +TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") { + char memsetval; + int memsetD32val; + int16_t memsetD16val; + char memsetD8val; + size_t N; + bool ret; + + std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) = + GENERATE(table(tableItems)); + + + SECTION("Memset with hipMemsetTypeDefault") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemsetAsync(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD32") { + int32_t *iA_d{nullptr}, *iA_h{nullptr}; + ret = testhipMemsetAsync(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD16") { + int16_t *siA_d{nullptr}, *siA_h{nullptr}; + ret = testhipMemsetAsync(siA_h, siA_d, memsetD16val, hipMemsetTypeD16, N); + REQUIRE(ret == true); + } + + SECTION("Memset with hipMemsetTypeD8") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemsetAsync(cA_h, cA_d, memsetD8val, hipMemsetTypeD8, N); + REQUIRE(ret == true); + } +} + +/** + * Test hipMemset small size buffers with unique memset values. + */ +TEST_CASE("Unit_hipMemset_SmallBufferSizes") { + char *A_d, *A_h; + constexpr int memsetval = 0x24; + + auto numElements = GENERATE(range(1, 4)); + int numBytes = numElements * sizeof(char); + + HIP_CHECK(hipMalloc(&A_d, numBytes)); + A_h = reinterpret_cast (malloc(numBytes)); + + HIP_CHECK(hipMemset(A_d, memsetval, numBytes)); + HIP_CHECK(hipMemcpy(A_h, A_d, numBytes, hipMemcpyDeviceToHost)); + + for (int i = 0; i < numBytes; i++) { + if (A_h[i] != memsetval) { + INFO("Mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(false); + } + } + + HIP_CHECK(hipFree(A_d)); + free(A_h); +} + + +/** + * Test two memset async operations at the same time. + */ +TEST_CASE("Unit_hipMemset_2AsyncOperations") { + std::vector v; + v.resize(2048); + float* p2, *p3; + hipMalloc(reinterpret_cast(&p2), 4096 + 4096*2); + p3 = p2+2048; + hipStream_t s; + hipStreamCreate(&s); + hipMemsetAsync(p2, 0, 32*32*4, s); + hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s); + hipStreamSynchronize(s); + for (int i = 0; i < 256; ++i) { + hipMemsetAsync(p2, 0, 32*32*4, s); + hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s); + } + hipStreamSynchronize(s); + hipDeviceSynchronize(); + hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost); + hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost); + + REQUIRE(v[0] == 0); + REQUIRE(v[1024] == 1.75f); +} diff --git a/tests/catch/unit/memory/hipMemset2D.cc b/tests/catch/unit/memory/hipMemset2D.cc new file mode 100644 index 0000000000..e48b8931f0 --- /dev/null +++ b/tests/catch/unit/memory/hipMemset2D.cc @@ -0,0 +1,175 @@ +/* + * 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. + */ + +/** + Testcase Scenarios : + 1) hipMemset2D api with basic functionality. + 2) hipMemset2DAsync api with basic functionality. + 3) hipMemset2D api with partial memset and unique width/height. +*/ + + +#include + + +// Table with unique width/height and memset values. +// (width2D, height2D, memsetWidth, memsetHeight) +typedef std::tuple tupletype; + +static constexpr std::initializer_list tableItems { + std::make_tuple(20, 20, 20, 20), + std::make_tuple(10, 10, 4, 4), + std::make_tuple(100, 100, 20, 40), + std::make_tuple(256, 256, 39, 19), + std::make_tuple(100, 100, 20, 0), + std::make_tuple(100, 100, 0, 20), + std::make_tuple(100, 100, 0, 0), + }; + + + +/** + * Basic Functionality of hipMemset2D + */ +TEST_CASE("Unit_hipMemset2D_BasicFunctional") { + constexpr int memsetval = 0x24; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH; + size_t elements = numW * numH; + char *A_d, *A_h; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, + numH)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, + hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("Memset2D mismatch at index:" << i << " computed:" + << A_h[i] << " memsetval:" << memsetval); + REQUIRE(false); + } + } + + hipFree(A_d); + free(A_h); +} + + +/** + * Basic Functionality of hipMemset2DAsync + */ +TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") { + constexpr int memsetval = 0x26; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH; + size_t elements = numW * numH; + char *A_d, *A_h; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width, numH)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, + hipMemcpyDeviceToHost)); + + for (size_t i=0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("Memset2DAsync mismatch at index:" << i << " computed:" + << A_h[i] << " memsetval:" << memsetval); + REQUIRE(false); + } + } + + hipFree(A_d); + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); +} + + +/** + * Memset partial buffer with unique Width and Height + */ +TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") { + int width2D, height2D; + int memsetWidth, memsetHeight; + char *A_d, *A_h; + size_t pitch_A; + constexpr int memsetval = 0x26; + + std::tie(width2D, height2D, memsetWidth, memsetHeight) = + GENERATE(table(tableItems)); + + size_t width = width2D * sizeof(char); + size_t sizeElements = width * height2D; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width, height2D)); + + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t index = 0; index < sizeElements; index++) { + A_h[index] = 'c'; + } + + INFO("2D Dimension: Width:" << width2D << " Height:" << height2D << + " MemsetWidth:" << memsetWidth << " MemsetHeight:" << memsetHeight); + + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight)); + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, width2D, height2D, + hipMemcpyDeviceToHost)); + + for (int row = 0; row < memsetHeight; row++) { + for (int column = 0; column < memsetWidth; column++) { + if (A_h[(row * width) + column] != memsetval) { + INFO("A_h[" << row << "][" << column << "]" << + " didnot match " << memsetval); + REQUIRE(false); + } + } + } + + hipFree(A_d); + free(A_h); +} + diff --git a/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc b/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc new file mode 100644 index 0000000000..04240a4104 --- /dev/null +++ b/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc @@ -0,0 +1,185 @@ +/* + * 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) Order of execution of device kernel and hipMemset2DAsync api + 2) hipMemSet2DAsync execution in multiple threads +*/ + +#include +#include +#include +#include + + +/* Defines */ +#define NUM_THREADS 1000 +#define ITER 100 +#define NUM_H 256 +#define NUM_W 256 + + + +void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch, + size_t width, hipStream_t stream) { + constexpr int memsetval = 0x22; + HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream)); + HIPCHECK(hipMemcpy2DAsync(A_h, width, A_d, pitch, NUM_W, NUM_H, + hipMemcpyDeviceToHost, stream)); +} + + +/** + * Order of execution of device kernel and hipMemset2DAsync api. + */ +TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { + constexpr auto N = 4 * 1024 * 1024; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + constexpr int memsetval = 0x22; + char *A_d, *A_h, *B_d, *B_h, *C_d; + size_t pitch_A, pitch_B, pitch_C; + size_t width = NUM_W * sizeof(char); + size_t sizeElements = width * NUM_H; + size_t elements = NUM_W * NUM_H; + unsigned blocks{}; + int validateCount{}; + hipStream_t stream; + + blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, + width, NUM_H)); + + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + B_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(B_h != nullptr); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, + width, NUM_H)); + + for (size_t i = 0; i < elements; i++) { + B_h[i] = i; + } + HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipStreamCreate(&stream)); + + + for (size_t k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, B_d, C_d, elements); + + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); + + for (size_t p = 0 ; p < elements ; p++) { + if (A_h[p] == memsetval) { + validateCount+= 1; + } + } + } + + REQUIRE(static_cast(validateCount) == (ITER * elements)); + + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); + free(A_h); free(B_h); + HIP_CHECK(hipStreamDestroy(stream)); +} + + +/** + * hipMemSet2DAsync execution in multiple threads. + */ +TEST_CASE("Unit_hipMemset2DAsync_MultiThread") { + constexpr auto N = 4 * 1024 * 1024; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + constexpr auto memPerThread = 200; + constexpr int memsetval = 0x22; + char *A_d, *A_h, *B_d, *B_h, *C_d; + size_t pitch_A, pitch_B, pitch_C; + size_t width = NUM_W * sizeof(char); + size_t sizeElements = width * NUM_H; + size_t elements = NUM_W * NUM_H; + unsigned blocks{}; + int validateCount{}; + hipStream_t stream; + + blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + auto thread_count = HipTest::getHostThreadCount(memPerThread, NUM_THREADS); + if (thread_count == 0) { + WARN("Resources not available for thread creation"); + return; + } + + std::thread *t = new std::thread[thread_count]; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, + width, NUM_H)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + B_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(B_h != nullptr); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, + width, NUM_H)); + + for (size_t i = 0 ; i < elements ; i++) { + B_h[i] = i; + } + HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipStreamCreate(&stream)); + + for (int i = 0 ; i < ITER ; i++) { + for (size_t k = 0 ; k < thread_count; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, A_h, pitch_A, + width, stream); + } else { + t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, B_h, pitch_A, + width, stream); + } + } + for (size_t j = 0 ; j < thread_count; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t k = 0 ; k < elements ; k++) { + if ((A_h[k] == memsetval) && (B_h[k] == memsetval)) { + validateCount+= 1; + } + } + } + + REQUIRE(static_cast(validateCount) == (ITER * elements)); + + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); + free(A_h); free(B_h); + HIP_CHECK(hipStreamDestroy(stream)); + + delete[] t; +} diff --git a/tests/catch/unit/memory/hipMemset3D.cc b/tests/catch/unit/memory/hipMemset3D.cc new file mode 100644 index 0000000000..0fc1a83818 --- /dev/null +++ b/tests/catch/unit/memory/hipMemset3D.cc @@ -0,0 +1,128 @@ +/* +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. +*/ + +/** + Functional test for Memset3D and Memset3DAsync + */ + + +#include + + +/** + * Basic Functional test of hipMemset3D + */ +TEST_CASE("Unit_hipMemset3D_BasicFunctional") { + constexpr int memsetval = 0x22; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + constexpr size_t depth = 10; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH * depth; + size_t elements = numW * numH * depth; + char *A_h; + + hipExtent extent = make_hipExtent(width, numH, depth); + hipPitchedPtr devPitchedPtr; + + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + 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("Memset3D mismatch at index:" << i << " computed:" + << A_h[i] << " memsetval:" << memsetval); + REQUIRE(false); + } + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); +} + +/** + * Basic Functional test of hipMemset3DAsync + */ +TEST_CASE("Unit_hipMemset3DAsync_BasicFunctional") { + constexpr int memsetval = 0x22; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + constexpr size_t depth = 10; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH * depth; + size_t elements = numW * numH * depth; + hipExtent extent = make_hipExtent(width, numH, depth); + hipPitchedPtr devPitchedPtr; + char *A_h; + + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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("Memset3DAsync mismatch at index:" << i << " computed:" + << A_h[i] << " memsetval:" << memsetval); + REQUIRE(false); + } + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); +} diff --git a/tests/catch/unit/memory/hipMemset3DFunctional.cc b/tests/catch/unit/memory/hipMemset3DFunctional.cc new file mode 100644 index 0000000000..421e066ab4 --- /dev/null +++ b/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/tests/catch/unit/memory/hipMemset3DNegative.cc b/tests/catch/unit/memory/hipMemset3DNegative.cc new file mode 100644 index 0000000000..798c90e1ac --- /dev/null +++ b/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/tests/catch/unit/memory/hipMemset3DRegressMultiThread.cc b/tests/catch/unit/memory/hipMemset3DRegressMultiThread.cc new file mode 100644 index 0000000000..b3a65f2104 --- /dev/null +++ b/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)); +} diff --git a/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc b/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc new file mode 100644 index 0000000000..adbd4a3964 --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc @@ -0,0 +1,193 @@ +/* + * 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. +*/ + +/* + * Test for checking order of execution of device kernel and + * hipMemsetAsync apis on all gpus + */ + +#include +#include +#include + +#define ITER 6 +#define N 1024 * 1024 + +constexpr auto blocksPerCU = 6; // to hide latency +constexpr auto threadsPerBlock = 256; +static unsigned blocks = 0; + + +template +class MemSetKernelTest { + public: + T *A_h, *B_d, *B_h, *C_d; + T memSetVal; + size_t Nbytes; + bool testResult = true; + int validateCount = 0; + hipStream_t stream; + + void memAllocate(T memSetValue) { + memSetVal = memSetValue; + Nbytes = N * sizeof(T); + + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(A_h != nullptr); + HIP_CHECK(hipMalloc(&B_d , Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(B_h != nullptr); + HIP_CHECK(hipMalloc(&C_d , Nbytes)); + + for (int i = 0 ; i < N ; i++) { + B_h[i] = i; + } + HIP_CHECK(hipMemcpy(B_d , B_h , Nbytes , hipMemcpyHostToDevice)); + HIP_CHECK(hipStreamCreate(&stream)); + } + + void memDeallocate() { + HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); + free(B_h); free(A_h); + HIP_CHECK(hipStreamDestroy(stream)); + } + + void validateExecutionOrder() { + for (int p = 0 ; p < N ; p++) { + if (A_h[p] == memSetVal) { + validateCount+= 1; + } + } + } + + bool resultAfterAllIterations() { + testResult = (validateCount == (ITER * N)) ? true : false; + memDeallocate(); + return testResult; + } +}; + +static bool testhipMemsetAsyncWithKernel() { + MemSetKernelTest obj; + constexpr char memsetval = 0x42; + + obj.memAllocate(memsetval); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +static bool testhipMemsetD32AsyncWithKernel() { + MemSetKernelTest obj; + constexpr int memsetD32val = 0xDEADBEEF; + + obj.memAllocate(memsetD32val); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal, + N, obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +static bool testhipMemsetD16AsyncWithKernel() { + MemSetKernelTest obj; + constexpr int16_t memsetD16val = 0xDEAD; + + obj.memAllocate(memsetD16val); + for (int k = 0 ; k < ITER ; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)obj.C_d , obj.memSetVal, + N, obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + +static bool testhipMemsetD8AsyncWithKernel() { + MemSetKernelTest obj; + constexpr char memsetD8val = 0xDE; + + obj.memAllocate(memsetD8val); + for (int k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); + HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)obj.C_d, obj.memSetVal, + N, obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); + + obj.validateExecutionOrder(); + } + return obj.resultAfterAllIterations(); +} + + +/* + * Test for checking order of execution of device kernel and + * hipMemsetAsync apis on all gpus + */ +TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") { + int numDevices = 0; + bool ret; + + blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGetDeviceCount(&numDevices)); + REQUIRE(numDevices > 0); + + auto devNum = GENERATE_COPY(range(0, numDevices)); + HIP_CHECK(hipSetDevice(devNum)); + + SECTION("hipMemsetAsync With Kernel") { + ret = testhipMemsetAsyncWithKernel(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD32Async With Kernel") { + ret = testhipMemsetD32AsyncWithKernel(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD16Async With Kernel") { + ret = testhipMemsetD16AsyncWithKernel(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD8Async With Kernel") { + ret = testhipMemsetD8AsyncWithKernel(); + REQUIRE(ret == true); + } +} diff --git a/tests/catch/unit/memory/hipMemsetAsyncMultiThread.cc b/tests/catch/unit/memory/hipMemsetAsyncMultiThread.cc new file mode 100644 index 0000000000..68deacda92 --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetAsyncMultiThread.cc @@ -0,0 +1,243 @@ +/* + * 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. + */ + +/* + * Test that validates functionality of hipmemsetAsync apis over multi threads + */ + +#include +#include + + +#define NUM_THREADS 20 +#define ITER 10 +#define N (4*1024*1024) + + +template +class MemSetAsyncMthreadTest { + public: + T *A_h, *A_d, *B_h; + T memSetVal; + size_t Nbytes; + bool testResult = true; + int validateCount = 0; + hipStream_t stream; + + void memAllocate(T memSetValue) { + memSetVal = memSetValue; + Nbytes = N * sizeof(T); + + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(A_h != nullptr); + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(B_h != nullptr); + + HIP_CHECK(hipStreamCreate(&stream)); + } + + void threadCompleteStatus() { + for (int k = 0 ; k < N ; k++) { + if ((A_h[k] == memSetVal) && (B_h[k] == memSetVal)) { + validateCount+= 1; + } + } + } + + bool resultAfterAllIterations() { + memDeallocate(); + testResult = (validateCount == (ITER * N)) ? true: false; + return testResult; + } + + void memDeallocate() { + HIP_CHECK(hipFree(A_d)); + free(A_h); + free(B_h); + HIP_CHECK(hipStreamDestroy(stream)); + } +}; + +template +void queueJobsForhipMemsetAsync(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetAsync(A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +template +void queueJobsForhipMemsetD32Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +template +void queueJobsForhipMemsetD16Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +template +void queueJobsForhipMemsetD8Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes, + hipStream_t stream) { + HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memSetVal, N, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); +} + +/* Queue hipMemsetAsync jobs on multiple threads and verify they all + * finished on all threads successfully + */ +bool testhipMemsetAsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr char memsetval = 0x42; + obj.memAllocate(memsetval); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetAsync, obj.A_d, obj.A_h, + obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetAsync, obj.A_d, obj.B_h, + obj.memSetVal, obj.Nbytes, obj.stream); + } + } + + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD32AsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr int memsetD32val = 0xDEADBEEF; + obj.memAllocate(memsetD32val); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetD32Async, obj.A_d, + obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetD32Async, obj.A_d, + obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); + } + } + + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD16AsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr int16_t memsetD16val = 0xDEAD; + obj.memAllocate(memsetD16val); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetD16Async, obj.A_d, + obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetD16Async, obj.A_d, + obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); + } + } + + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + +bool testhipMemsetD8AsyncWithMultiThread() { + MemSetAsyncMthreadTest obj; + constexpr char memsetD8val = 0xDE; + obj.memAllocate(memsetD8val); + std::thread t[NUM_THREADS]; + + for (int i = 0 ; i < ITER ; i++) { + for (int k = 0 ; k < NUM_THREADS ; k++) { + if (k%2) { + t[k] = std::thread(queueJobsForhipMemsetD8Async, obj.A_d, + obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream); + } else { + t[k] = std::thread(queueJobsForhipMemsetD8Async, obj.A_d, + obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream); + } + } + for (int j = 0 ; j < NUM_THREADS ; j++) { + t[j].join(); + } + + HIP_CHECK(hipStreamSynchronize(obj.stream)); + obj.threadCompleteStatus(); + } + return obj.resultAfterAllIterations(); +} + + +/* + * Test that validates functionality of hipmemsetAsync apis over multi threads + */ +TEST_CASE("Unit_hipMemsetAsync_QueueJobsMultithreaded") { + bool ret; + + SECTION("hipMemsetAsync With MultiThread") { + ret = testhipMemsetAsyncWithMultiThread(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD32Async With MultiThread") { + ret = testhipMemsetD32AsyncWithMultiThread(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD16Async With MultiThread") { + ret = testhipMemsetD16AsyncWithMultiThread(); + REQUIRE(ret == true); + } + + SECTION("hipMemsetD8Async With MultiThread") { + ret = testhipMemsetD8AsyncWithMultiThread(); + REQUIRE(ret == true); + } +} diff --git a/tests/catch/unit/memory/hipMemsetInvalidPtr.cc b/tests/catch/unit/memory/hipMemsetInvalidPtr.cc new file mode 100644 index 0000000000..8c87c66ea8 --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetInvalidPtr.cc @@ -0,0 +1,127 @@ +/* + * 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 hipMemset apis with invalid pointer and invalid 2D pitch. + 2) Test hipMemsetAsync apis with invalid pointer and invalid 2D pitch. +*/ + + +#include + +#define N 50 +#define MEMSETVAL 0x42 + +/** + * Testcase validates hipMemset apis behavior with + * invalid pointer and invalid 2D pitch value. + */ +TEST_CASE("Unit_hipMemset_InvalidPtrTests") { + hipError_t ret; + constexpr int Nbytes = N*sizeof(char); + char *A_d; + + SECTION("hipMemset with null") { + ret = hipMemset(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemset with hostptr") { + char *A_h; + A_h = reinterpret_cast(malloc(Nbytes)); + + ret = hipMemset(A_h, MEMSETVAL, Nbytes); + REQUIRE(ret != hipSuccess); + + free(A_h); + } + + SECTION("hipMemsetD32 with null") { + ret = hipMemsetD32(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD16 with null") { + ret = hipMemsetD16(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD8 with null") { + ret = hipMemsetD8(NULL, MEMSETVAL , Nbytes); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemset2D with null") { + constexpr size_t NUM_H = 256, NUM_W = 256; + size_t pitch_A; + size_t width = NUM_W * sizeof(char); + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width , NUM_H)); + ret = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H); + REQUIRE(ret != hipSuccess); + + hipFree(A_d); + } +} + + +/** + * Testcase validates hipMemsetAsync apis behavior with + * invalid pointer and invalid 2D pitch value. + */ +TEST_CASE("Unit_hipMemsetAsync_InvalidPtrTests") { + hipError_t ret; + constexpr int Nbytes = N*sizeof(char); + char *A_d; + + SECTION("hipMemsetAsync with null") { + ret = hipMemsetAsync(NULL, MEMSETVAL, Nbytes , 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD32Async with null") { + ret = hipMemsetD32Async(NULL, MEMSETVAL , Nbytes, 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD16Async with null") { + ret = hipMemsetD16Async(NULL, MEMSETVAL , Nbytes, 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemsetD8Async with null") { + ret = hipMemsetD8Async(NULL, MEMSETVAL , Nbytes, 0); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipMemset2DAsync with null") { + constexpr size_t NUM_H = 256, NUM_W = 256; + size_t pitch_A; + size_t width = NUM_W * sizeof(char); + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, + width , NUM_H)); + ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0); + REQUIRE(ret != hipSuccess); + + hipFree(A_d); + } +} diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 14c05e7f74..31157f7da6 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -9,6 +9,8 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipStreamWithCUMask.cc + hipStreamGetCUMask.cc + hipAPIStreamDisable.cc ) else() set(TEST_SRC @@ -20,6 +22,7 @@ set(TEST_SRC hipStreamAddCallback.cc hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc + hipAPIStreamDisable.cc ) endif() diff --git a/tests/catch/unit/stream/hipAPIStreamDisable.cc b/tests/catch/unit/stream/hipAPIStreamDisable.cc new file mode 100644 index 0000000000..4da259074b --- /dev/null +++ b/tests/catch/unit/stream/hipAPIStreamDisable.cc @@ -0,0 +1,68 @@ +/* +Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include "hip/math_functions.h" + +#define NUM_STREAMS 8 + +namespace hipAPIStreamDisableTest { +const int NN = 1 << 21; + +__global__ void kernel(float* x, float* y, int n) { + int tid = threadIdx.x; + if (tid < 1) { + for (int i = 0; i < n; i++) { + x[i] = sqrt(powf(3.14159, i)); + } + y[tid] = y[tid] + 1.0f; + } +} + +__global__ void nKernel(float* y) { + int tid = threadIdx.x; + y[tid] = y[tid] + 1.0f; +} +} // namespace hipAPIStreamDisableTest + +/** + * Validate basic multistream functionalities + */ +TEST_CASE("Unit_hipStreamCreate_MultistreamBasicFunctionalities") { + hipStream_t streams[NUM_STREAMS]; + float *data[NUM_STREAMS], *yd, *xd; + float y = 1.0f, x = 1.0f; + HIP_CHECK(hipMalloc(reinterpret_cast(&yd), sizeof(float))); + HIP_CHECK(hipMalloc(reinterpret_cast(&xd), sizeof(float))); + HIP_CHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice)); + + for (int i = 0; i < NUM_STREAMS; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + HIP_CHECK(hipMalloc(&data[i], + (hipAPIStreamDisableTest::NN)*sizeof(float))); + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::kernel), + dim3(1), dim3(1), 0, streams[i], data[i], xd, + hipAPIStreamDisableTest::NN); + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipAPIStreamDisableTest::nKernel), + dim3(1), dim3(1), 0, 0, yd); + } + HIP_CHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost)); + REQUIRE(x == y); +} diff --git a/tests/catch/unit/stream/hipStreamGetCUMask.cc b/tests/catch/unit/stream/hipStreamGetCUMask.cc new file mode 100644 index 0000000000..49f21bdbda --- /dev/null +++ b/tests/catch/unit/stream/hipStreamGetCUMask.cc @@ -0,0 +1,177 @@ +/* +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. +*/ + +/** +Testcase Scenarios : +1) Test to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask. +2) Test to verify hipExtStreamGetCUMask api returns custom mask set. +3) Negative tests for hipExtStreamGetCUMask api. +*/ + +#include +#include + + +/** + * Scenario to verify hipExtStreamGetCUMask api returning default CU Mask or global CU Mask. + * Scenario to verify hipExtStreamGetCUMask api returns custom mask set. + */ +TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { + constexpr int maxNum = 6; + std::vector cuMask(maxNum); + hipDeviceProp_t props; + std::stringstream ss; + char* gCUMask{nullptr}; + std::string globalCUMask(""); + std::vector defaultCUMask; + + int nGpu = 0; + HIP_CHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + INFO("info: didn't find any GPU! skipping the test!"); + return; + } + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + INFO("info: running on bus " << "0x" << props.pciBusID << " " << + props.name << " with " << props.multiProcessorCount << " CUs"); + + // Get global CU Mask if exists + gCUMask = getenv("ROC_GLOBAL_CU_MASK"); + if (gCUMask != nullptr && gCUMask[0] != '\0') { + globalCUMask.assign(gCUMask); + + for_each(globalCUMask.begin(), globalCUMask.end(), [](char & c) { + c = ::tolower(c); + }); + } + + // Create default CU Mask + uint32_t temp = 0; + uint32_t bit_index = 0; + for (uint32_t i = 0; i < (uint32_t)props.multiProcessorCount; i++) { + temp |= 1UL << bit_index; + if (bit_index >= 32) { + defaultCUMask.push_back(temp); + temp = 0; + bit_index = 0; + temp |= 1UL << bit_index; + } + bit_index += 1; + } + if (bit_index != 0) { + defaultCUMask.push_back(temp); + } + + SECTION("Verify with default CU Mask or global CU Mask") { + // make a default CU mask bit-array where all CUs are active + // this default mask is expected to be returned when there is no + // custom or global CU mask defined + + HIP_CHECK(hipExtStreamGetCUMask(0, cuMask.size(), &cuMask[0])); + + ss << std::hex; + for (int i = cuMask.size() - 1; i >= 0; i--) { + ss << cuMask[i]; + } + + // remove extra 0 from ss if any present + size_t found = ss.str().find_first_not_of("0"); + if (found != std::string::npos) { + ss.str(ss.str().substr(found, ss.str().length())); + } + + INFO("info: CU mask for the default stream is: 0x" << ss.str().c_str()); + if (globalCUMask.size() > 0) { + if (ss.str().compare(globalCUMask) != 0) { + INFO("Expected CU mask:" << globalCUMask.c_str() << + ", api returned:" << ss.str().c_str()); + REQUIRE(false); + } + } else { + for (int i = 0 ; i < min(cuMask.size(), defaultCUMask.size()); i++) { + if (cuMask[i] != defaultCUMask[i]) { + INFO("Expected CU mask " << defaultCUMask[i] << + ", api returned:" << cuMask[i]); + REQUIRE(false); + } + } + } + } + + SECTION("Verify with custom mask set") { + std::vector customMask(defaultCUMask); + hipStream_t stream; + customMask[0] = 0xe; + + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, customMask.size(), + customMask.data())); + ss.str(""); + for (int i = customMask.size() - 1; i >= 0; i--) { + ss << customMask[i]; + } + INFO("info: setting a custom CU mask 0x" << ss.str()); + + HIP_CHECK(hipExtStreamGetCUMask(stream, cuMask.size(), &cuMask[0])); + ss.str(""); + for (int i = cuMask.size() - 1; i >= 0; i--) { + ss << cuMask[i]; + } + + size_t found = ss.str().find_first_not_of("0"); + if (found != std::string::npos) { + ss.str(ss.str().substr(found, ss.str().length())); + } + + INFO("info: reading back CU mask 0x" << ss.str() << + " for stream " << stream); + + if (!gCUMask) { + for (size_t i = 0; i < customMask.size(); i++) { + if (customMask[i] != cuMask[i]) { + INFO("Error! expected CU mask:" << customMask[i] + << ", Received CU mask:" << cuMask[i]); + REQUIRE(false); + } + } + } + + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +/** + * Negative tests for hipExtStreamGetCUMask. + */ +TEST_CASE("Unit_hipExtStreamGetCUMask_Negative") { + hipError_t ret; + constexpr int maxNum = 6; + std::vector cuMask(maxNum); + + SECTION("cuMask is nullptr") { + ret = hipExtStreamGetCUMask(0, cuMask.size(), nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("cuMaskSize is 0") { + ret = hipExtStreamGetCUMask(0, 0, &cuMask[0]); + REQUIRE(ret == hipErrorInvalidValue); + } +} diff --git a/tests/catch/unit/stream/hipStreamGetFlags.cc b/tests/catch/unit/stream/hipStreamGetFlags.cc index 583abe5068..9cbe37b6de 100644 --- a/tests/catch/unit/stream/hipStreamGetFlags.cc +++ b/tests/catch/unit/stream/hipStreamGetFlags.cc @@ -16,16 +16,24 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +/** +Testcase Scenarios : +1) Test flag value of stream created with hipStreamCreateWithFlags/ + /hipStreamCreate/hipStreamCreateWithPriority. +2) Negative tests for hipStreamGetFlags api. +3) Test flag value when streams created with CUMask. +*/ + #include -TEST_CASE("Unit_hipStreamGetFlags_Negative") { - // Get flags for uninitialized stream + +/** + * Test flag value of stream created with various types. + */ +TEST_CASE("Unit_hipStreamGetFlags_BasicFunctionalities") { hipStream_t stream; - HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); - REQUIRE(hipStreamGetFlags(stream, nullptr) == hipErrorInvalidValue); -} -TEST_CASE("Unit_hipStreamGetFlags") { - hipStream_t stream; - unsigned int flags; + unsigned int flags; + // Check flag value of stream created with hipStreamCreateWithFlags + SECTION("Check flag value of streams hipStreamCreateWithFlags") { HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); HIP_CHECK(hipStreamGetFlags(stream, &flags)); REQUIRE(flags == hipStreamDefault); @@ -34,4 +42,50 @@ TEST_CASE("Unit_hipStreamGetFlags") { HIP_CHECK(hipStreamGetFlags(stream, &flags)); REQUIRE(flags == hipStreamNonBlocking); HIP_CHECK(hipStreamDestroy(stream)); + } + // Check flag value of stream created with hipStreamCreate + SECTION("Check flag value of streams hipStreamCreate") { + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); + } + // Check flag value of stream created with hipStreamCreateWithPriority + SECTION("Check flag value of streams hipStreamCreateWithPriority") { + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, 0)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, 0)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamNonBlocking); + HIP_CHECK(hipStreamDestroy(stream)); + } } + +/** + * Negative Scenarios + */ +TEST_CASE("Unit_hipStreamGetFlags_Negative") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + // nullptr check + REQUIRE(hipStreamGetFlags(stream, nullptr) != hipSuccess); + REQUIRE(hipStreamGetFlags(nullptr, hipStreamDefault) != hipSuccess); + HIP_CHECK(hipStreamDestroy(stream)); +} + +#if HT_AMD +/** + * Test flag value when streams created with CUMask. + */ +TEST_CASE("Unit_hipStreamGetFlags_StreamsCreatedWithCUMask") { + hipStream_t stream; + unsigned int flags; + const uint32_t cuMask = 0xffffffff; + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask)); + HIP_CHECK(hipStreamGetFlags(stream, &flags)); + REQUIRE(flags == hipStreamDefault); + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif diff --git a/tests/catch/unit/stream/hipStreamGetPriority.cc b/tests/catch/unit/stream/hipStreamGetPriority.cc index 3578fd9b5a..29f182ecdf 100644 --- a/tests/catch/unit/stream/hipStreamGetPriority.cc +++ b/tests/catch/unit/stream/hipStreamGetPriority.cc @@ -16,11 +16,28 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +/** +Testcase Scenarios : +1) Negative tests for hipStreamGetPriority api. +2) Create stream and check default priority of stream is within range. +3) Create stream with high or low priority and check priority is set as expected. +4) Create stream with higher priority or lower priority for the priority range returned. +5) Create stream with CUMask and check priority is returned as expected. +*/ + #include + +/** + * Negative tests for hipStreamGetPriority api. + */ TEST_CASE("Unit_hipStreamGetPriority_Negative") { hipStream_t stream = 0; REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue); } + +/** + * Create stream and check default priority of stream is within range. + */ TEST_CASE("Unit_hipStreamGetPriority_default") { int priority_low = 0; int priority_high = 0; @@ -37,6 +54,10 @@ TEST_CASE("Unit_hipStreamGetPriority_default") { REQUIRE(priority >= priority_high); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with high priority and check priority is set as expected. + */ TEST_CASE("Unit_hipStreamGetPriority_high") { int priority_low = 0; int priority_high = 0; @@ -44,12 +65,17 @@ TEST_CASE("Unit_hipStreamGetPriority_high") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + priority_high)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority == priority_high); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with higher priority for the priority range returned. + */ TEST_CASE("Unit_hipStreamGetPriority_higher") { int priority_low = 0; int priority_high = 0; @@ -57,12 +83,17 @@ TEST_CASE("Unit_hipStreamGetPriority_higher") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, + priority_high-1)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority == priority_high); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with low priority and check priority is set as expected. + */ TEST_CASE("Unit_hipStreamGetPriority_low") { int priority_low = 0; int priority_high = 0; @@ -70,12 +101,17 @@ TEST_CASE("Unit_hipStreamGetPriority_low") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, + priority_low)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority_low == priority); HIP_CHECK(hipStreamDestroy(stream)); } + +/** + * Create stream with lower priority for the priority range returned. + */ TEST_CASE("Unit_hipStreamGetPriority_lower") { int priority_low = 0; int priority_high = 0; @@ -83,9 +119,35 @@ TEST_CASE("Unit_hipStreamGetPriority_lower") { HIP_CHECK(hipSetDevice(devID)); HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); hipStream_t stream; - HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1)); + HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, + priority_low+1)); int priority = 0; HIP_CHECK(hipStreamGetPriority(stream, &priority)); REQUIRE(priority_low == priority); HIP_CHECK(hipStreamDestroy(stream)); } + +#if HT_AMD +/** + * Create stream with CUMask and check priority is returned as expected. + */ +TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") { + hipStream_t stream; + int priority = 0; + int priority_normal = 0; + int priority_low = 0; + int priority_high = 0; + // Test is to get the Stream Priority Range + HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + priority_normal = priority_low + priority_high; + // Check if priorities are indeed supported + REQUIRE_FALSE(priority_low == priority_high); + // Creating a stream with hipExtStreamCreateWithCUMask and checking + // priority. + const uint32_t cuMask = 0xffffffff; + HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, 1, &cuMask)); + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + REQUIRE_FALSE(priority_normal != priority); + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif