From 528764ec78cb99b888e8f891826b57fa879150c0 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 14 Aug 2023 20:59:45 +0530 Subject: [PATCH] SWDEV-396085 - [catch2][dtest] Adding test cases for hipHostRegister() to test SVM feature (#324) Change-Id: I72bb3d1cf3410180c98f4629fdad7497698849a2 --- catch/unit/memory/CMakeLists.txt | 7 +- catch/unit/memory/hipHostRegister.cc | 831 ++++++++++++++++++++++- catch/unit/memory/hipHostRegister_exe.cc | 155 +++++ 3 files changed, 958 insertions(+), 35 deletions(-) create mode 100644 catch/unit/memory/hipHostRegister_exe.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 6f236ac229..fa340863a6 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2023 Advanced Micro Devices, Inc. All Rights Reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -88,6 +88,11 @@ hip_add_exe_to_target(NAME MemoryTest1 TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) +if(HIP_PLATFORM MATCHES "amd") + set_source_files_properties(hipHostRegister.cc PROPERTIES COMPILE_FLAGS -std=c++17) + add_executable(hipHostRegisterPerf EXCLUDE_FROM_ALL hipHostRegister_exe.cc) +endif() + set(TEST_SRC hipMemcpyFromSymbol.cc hipPtrGetAttribute.cc diff --git a/catch/unit/memory/hipHostRegister.cc b/catch/unit/memory/hipHostRegister.cc index 5e1b10d234..cb62532ae7 100644 --- a/catch/unit/memory/hipHostRegister.cc +++ b/catch/unit/memory/hipHostRegister.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -20,20 +20,40 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -This testfile verifies the following scenarios of hipHostRegister API -1. Referencing the hipHostRegister variable from kernel and performing - memset on that variable.This is verified for different datatypes. -2. hipHostRegister and perform hipMemcpy on it. -*/ +/** + * @addtogroup hipHostRegister hipHostRegister + * @{ + * @ingroup MemoryTest + * `hipError_t hipHostRegister (void *hostPtr, size_t sizeBytes, unsigned int flags)` - + * register host memory so it can be accessed from the current device. + */ #include "hip/hip_runtime_api.h" #include #include +#include +#include #include #define OFFSET 128 +#define INITIAL_VAL 1 +#define EXPECTED_VAL 2 +#define ITERATION 100 +#define ADDITIONAL_MEMORY_PERCENT 10 + static constexpr auto LEN{1024 * 1024}; +static constexpr auto LARGE_CHUNK_LEN{100 * LEN}; +static constexpr auto SMALL_CHUNK_LEN{10 * LEN}; + +#if HT_AMD +#define TEST_SKIP(arch, msg) \ + if (std::string::npos == arch.find("xnack+")) {\ + HipTest::HIP_SKIP_TEST(msg);\ + return;\ + } +#else +#define TEST_SKIP(arch, msg) +#endif template __global__ void Inc(T* Ad) { int tx = threadIdx.x + blockIdx.x * blockDim.x; @@ -41,7 +61,8 @@ template __global__ void Inc(T* Ad) { } template -void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internalRegister) { +void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, + bool internalRegister) { constexpr auto memsetval = 13.0f; A = A + offset; numElements -= offset; @@ -71,18 +92,27 @@ void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internal } } -/* -This testcase verifies the hipHostRegister API by -1. Allocating the memory using malloc -2. hipHostRegister that variable -3. Getting the corresponding device pointer of the registered varible -4. Launching kernel and access the device pointer variable -5. performing hipMemset on the device pointer variable -*/ -TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", int, float, double) { +/** + * Test Description + * ------------------------ + * - This testcase verifies the hipHostRegister API by + * 1. Allocating the memory using malloc + * 2. hipHostRegister that variable + * 3. Getting the corresponding device pointer of the registered varible + * 4. Launching kernel and access the device pointer variable + * 5. performing hipMemset on the device pointer variable + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", \ + int, float, double) { size_t sizeBytes{LEN * sizeof(TestType)}; TestType *A, **Ad; - int num_devices; + int num_devices = 0; HIP_CHECK(hipGetDeviceCount(&num_devices)); Ad = new TestType*[num_devices]; A = reinterpret_cast(malloc(sizeBytes)); @@ -118,17 +148,722 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", i delete[] Ad; } -/* -This testcase verifies hipHostRegister API by -performing memcpy on the hipHostRegistered variable. -*/ +/** + * Test Description + * ------------------------ + * - This testcase verifies that the host pointer registered by hipHostRegister API + * is accessible from current device when xnack is on. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEMPLATE_TEST_CASE("Unit_hipHostRegister_DirectReferenceFromKernel", "", \ + int, float, double) { + auto flags = GENERATE(hipHostRegisterDefault, hipHostRegisterPortable, + hipHostRegisterMapped); + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t sizeBytes{LEN * sizeof(TestType)}; + TestType *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + // Initialize buffer with data + TestType val = static_cast(1); + for (int i = 0; i < LEN; i++) { + A[i] = val; + } + HIP_CHECK(hipHostRegister(A, sizeBytes, flags)); + + // Reference the registered device pointer A from inside the kernel: + hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, 0, A); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < LEN; i++) { + REQUIRE(A[i] == (val + static_cast(1))); + } + HIP_CHECK(hipHostUnregister(A)); + free(A); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies that the host pointer registered by hipHostRegister API + is usable from multiple device when xnack is on. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEMPLATE_TEST_CASE("Unit_hipHostRegister_DirectReferenceMultGpu", "", \ + int, float, double) { + // 1 refers to doing hipHostRegister once for all devices + // 0 refers to doing hipHostRegister for each device + auto register_once = GENERATE(0, 1); + hipDeviceProp_t prop; + int numDevices = HipTest::getDeviceCount(); + size_t sizeBytes{LEN * sizeof(TestType)}; + TestType *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + // Register host memory only once for all device + if (register_once == 1) { + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + } + // Reference the registered device pointer A from inside all devices: + for (int dev = 0; dev < numDevices; dev++) { + // Initialize buffer with data + TestType val = static_cast(1); + for (int i = 0; i < LEN; i++) { + A[i] = val; + } + HIP_CHECK(hipSetDevice(dev)); + HIP_CHECK(hipGetDeviceProperties(&prop, dev)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + // Register host memory for each device + if (register_once == 0) { + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + } + hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, 0, A); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < LEN; i++) { + REQUIRE(A[i] == (val + static_cast(1))); + } + if (register_once == 0) { + HIP_CHECK(hipHostUnregister(A)); + } + } + if (register_once == 1) { + HIP_CHECK(hipHostUnregister(A)); + } + free(A); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies functionality when same host pointer is repeatedly + * registered and unregistered. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_SameChunkRepeat") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t sizeBytes{LEN * sizeof(uint8_t)}; + uint8_t *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + for (int iter = 0; iter < ITERATION; iter++) { + // Initialize buffer with data + memset(A, INITIAL_VAL, sizeBytes); + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + + // Reference the registered device pointer A from inside the kernel: + hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, 0, A); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < LEN; i++) { + REQUIRE(A[i] == EXPECTED_VAL); + } + HIP_CHECK(hipHostUnregister(A)); + } + free(A); +} + +/** + * Test Description + * ------------------------ + * - Allocate a large chunk of host memory. Divide the memory into smaller chunks. + * Register each smaller chunk in one attempt. Access all the chunks in Kernel. Verify + * results. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_Chunks_SingleAttempt") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t sizeBytes{LARGE_CHUNK_LEN * sizeof(uint8_t)}; + size_t sizeBytesChunk{SMALL_CHUNK_LEN * sizeof(uint8_t)}; + uint8_t *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + // Initialize buffer with data + memset(A, INITIAL_VAL, sizeBytes); + uint8_t *arrPtr[LARGE_CHUNK_LEN / SMALL_CHUNK_LEN]; + for (int cnt = 0; cnt < (LARGE_CHUNK_LEN / SMALL_CHUNK_LEN); cnt++) { + arrPtr[cnt] = A + (cnt*sizeBytesChunk); + HIP_CHECK(hipHostRegister(arrPtr[cnt], sizeBytesChunk, 0)); + } + // Reference each registered chunk inside the kernel: + for (int cnt = 0; cnt < (LARGE_CHUNK_LEN / SMALL_CHUNK_LEN); cnt++) { + uint8_t *ptrA = arrPtr[cnt]; + hipLaunchKernelGGL(Inc, dim3(SMALL_CHUNK_LEN / 32), dim3(32), 0, 0, ptrA); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < SMALL_CHUNK_LEN; i++) { + REQUIRE(ptrA[i] == EXPECTED_VAL); + } + } + for (int cnt = 0; cnt < (LARGE_CHUNK_LEN / SMALL_CHUNK_LEN); cnt++) { + HIP_CHECK(hipHostUnregister(arrPtr[cnt])); + } + free(A); +} + +/** + * Test Description + * ------------------------ + * - Allocate a large chunk of host memory. Divide the memory into smaller chunks. + * Register each smaller chunk, access the chunk in Kernel and unregister the chunk. + * Verify results. Perform this series of operation in a round robin manner for + * all chunks. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_Chunks_RoundRobin") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t sizeBytes{LARGE_CHUNK_LEN * sizeof(uint8_t)}; + size_t sizeBytesChunk{SMALL_CHUNK_LEN * sizeof(uint8_t)}; + uint8_t *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + // Initialize buffer with data + memset(A, INITIAL_VAL, sizeBytes); + for (int cnt = 0; cnt < (LARGE_CHUNK_LEN / SMALL_CHUNK_LEN); cnt++) { + uint8_t *ptrA = A + (cnt*sizeBytesChunk); + HIP_CHECK(hipHostRegister(ptrA, sizeBytesChunk, 0)); + hipLaunchKernelGGL(Inc, dim3(SMALL_CHUNK_LEN / 32), dim3(32), 0, 0, ptrA); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < SMALL_CHUNK_LEN; i++) { + REQUIRE(ptrA[i] == EXPECTED_VAL); + } + HIP_CHECK(hipHostUnregister(ptrA)); + } + free(A); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies that the host pointer registered by hipHostRegister API + * can be memset using hipMemset. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_Perform_hipMemset") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t sizeBytes{LEN * sizeof(uint8_t)}; + uint8_t *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + // Register the host pointer + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + // Memset the registered pointer + HIP_CHECK(hipMemset(A, INITIAL_VAL, sizeBytes)); + // Reference the registered device pointer A from inside the kernel: + hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, 0, A); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < LEN; i++) { + REQUIRE(A[i] == EXPECTED_VAL); + } + HIP_CHECK(hipHostUnregister(A)); + free(A); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies that the host pointer registered by hipHostRegister API + * can be used with hipMemcpy. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_Perform_hipMemcpy") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t sizeBytes{LEN * sizeof(uint8_t)}; + uint8_t *A, *B; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + B = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(B != nullptr); + memset(B, INITIAL_VAL, sizeBytes); + // Register the host pointer + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + // Memcpy from B to A + HIP_CHECK(hipMemcpy(A, B, sizeBytes, hipMemcpyDefault)); + // Reference the registered device pointer A from inside the kernel: + hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, 0, A); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + // Verify if we can Memcpy from A to B + HIP_CHECK(hipMemcpy(B, A, sizeBytes, hipMemcpyDefault)); + for (int i = 0; i < LEN; i++) { + REQUIRE(B[i] == EXPECTED_VAL); + } + HIP_CHECK(hipHostUnregister(A)); + free(A); + free(B); +} + +/** + * Test Description + * ------------------------ + * - Oversubscription: This testcase allocates host memory of size > total + * GPU memory. Register the memory and try accessing it from kernel. Verify + * the behaviour. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_Oversubscription") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + TEST_SKIP(arch, "Xnack+ is not supported. Skipping the test ...") + size_t maxGpuMem = 0, availableMem = 0; + // Get available GPU memory and total GPU memory + HIP_CHECK(hipMemGetInfo(&availableMem, &maxGpuMem)); + size_t allocsize = maxGpuMem + + ((maxGpuMem*ADDITIONAL_MEMORY_PERCENT)/100); + // Get free host In bytes + size_t hostMemFree = HipTest::getMemoryAmount() * 1024 * 1024; + // Ensure that allocsize < hostMemFree + if (allocsize >= hostMemFree) { + HipTest::HIP_SKIP_TEST("Available Host Memory is not sufficient ..."); + return; + } + uint8_t* A = reinterpret_cast(malloc(allocsize)); + REQUIRE(A != nullptr); + size_t used_size = LEN; + // Inititalize only the first used_size bytes chunk + memset(A, INITIAL_VAL, used_size); + // Inititalize only the last used_size bytes chunk + memset((A + allocsize - used_size), INITIAL_VAL, used_size); + // Register the entire host memory chunk + HIP_CHECK(hipHostRegister(A, allocsize, 0)); + // Reference only the first used_size bytes + hipLaunchKernelGGL(Inc, dim3(used_size / 32), dim3(32), 0, 0, A); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < used_size; i++) { + REQUIRE(A[i] == EXPECTED_VAL); + } + // Reference only the last used_size bytes chunk + uint8_t* B = (A + allocsize - used_size); + hipLaunchKernelGGL(Inc, dim3(used_size / 32), dim3(32), 0, 0, B); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + for (int i = 0; i < used_size; i++) { + REQUIRE(B[i] == EXPECTED_VAL); + } + HIP_CHECK(hipHostUnregister(A)); + free(A); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies that the host pointer registered by hipHostRegister API + * can be used with Async APIs (hipMemsetAsync, hipMemcpyAsync and kernel) on a user + * defined stream. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_AsyncApis") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + bool useRegPtrInDev = false; +#if HT_AMD + if (std::string::npos == arch.find("xnack+")) { + useRegPtrInDev = false; + } else { + useRegPtrInDev = true; + } +#else + useRegPtrInDev = GENERATE(true, false); +#endif + size_t sizeBytes{LEN * sizeof(uint32_t)}; + uint32_t *A, *B, *dPtr; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + B = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(B != nullptr); + for (int i = 0; i < LEN; i++) { + B[i] = i; + } + // Register the host pointer + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + if (useRegPtrInDev) { + dPtr = A; + } else { + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&dPtr), A, 0)); + } + hipStream_t strm{nullptr}; + HIP_CHECK(hipStreamCreate(&strm)); + // Memcpy from B to A + HIP_CHECK(hipMemcpyAsync(dPtr, B, sizeBytes, hipMemcpyHostToDevice, strm)); + // Reference the registered device pointer A from inside the kernel: + hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, strm, dPtr); + HIP_CHECK(hipMemcpyAsync(B, dPtr, sizeBytes, hipMemcpyDeviceToHost, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + for (int i = 0; i < LEN; i++) { + REQUIRE(B[i] == (i + 1)); + } + HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipHostUnregister(A)); + free(A); + free(B); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies the behaviour of host registered memory when + * used with hipGraph. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_Graphs") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + bool useRegPtrInDev = false; +#if HT_AMD + if (std::string::npos == arch.find("xnack+")) { + useRegPtrInDev = false; + } else { + useRegPtrInDev = true; + } +#else + useRegPtrInDev = GENERATE(true, false); +#endif + size_t sizeBytes{LEN * sizeof(uint32_t)}; + uint32_t *A, *B, *dPtr; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + B = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(B != nullptr); + for (int i = 0; i < LEN; i++) { + B[i] = i; + } + // Register the host pointer + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + if (useRegPtrInDev) { + dPtr = A; + } else { + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&dPtr), A, 0)); + } + // Use dPtr in graphs + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D, memcpyD2H; + hipGraphNode_t kernel_vecInc; + void* kernelArgs1[] = {&dPtr}; + hipKernelNodeParams kernelNodeParams{}; + kernelNodeParams.func = reinterpret_cast(Inc); + kernelNodeParams.gridDim = dim3(LEN / 32); + kernelNodeParams.blockDim = dim3(32); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecInc, graph, nullptr, 0, + &kernelNodeParams)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, dPtr, B, + sizeBytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nullptr, 0, B, dPtr, + sizeBytes, hipMemcpyDeviceToHost)); + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D, &kernel_vecInc, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecInc, &memcpyD2H, 1)); + // Instantiate and execute Graph + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + // Verify Result + for (int i = 0; i < LEN; i++) { + REQUIRE(B[i] == (i + 1)); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipHostUnregister(A)); + free(A); + free(B); +} + +#if HT_AMD +/** + * Test Description + * ------------------------ + * - This testcase measures performance when same memory chunk is repeatedly + * registered and unregistered. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_RegUnreg_Perf_SameChunk") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string arch = prop.gcnArchName; + if (std::string::npos == arch.find("xnack+")) { + HipTest::HIP_SKIP_TEST("Xnack+ is not supported. Skipping the test ..."); + return; + } + hip::SpawnProc proc("hipHostRegisterPerf", true); + REQUIRE(proc.run("svm_enable 1") == 0); + float perf_svm_enable = std::stof(proc.getOutput()); + INFO("perf_svm_enable: " << perf_svm_enable); + REQUIRE(proc.run("svm_disable 1") == 0); + float perf_svm_disable = std::stof(proc.getOutput()); + INFO("perf_svm_disable: " << perf_svm_disable); + REQUIRE(perf_svm_enable <= perf_svm_disable); +} + +/** + * Test Description + * ------------------------ + * - This testcase measures performance when different memory chunks + * are repeatedly registered and unregistered. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_RegUnreg_Perf_DiffChunk") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string arch = prop.gcnArchName; + if (std::string::npos == arch.find("xnack+")) { + HipTest::HIP_SKIP_TEST("Xnack+ is not supported. Skipping the test ..."); + return; + } + hip::SpawnProc proc("hipHostRegisterPerf", true); + REQUIRE(proc.run("svm_enable 0") == 0); + float perf_svm_enable = std::stof(proc.getOutput()); + INFO("perf_svm_enable: " << perf_svm_enable); + REQUIRE(proc.run("svm_disable 0") == 0); + float perf_svm_disable = std::stof(proc.getOutput()); + INFO("perf_svm_disable: " << perf_svm_disable); + REQUIRE(perf_svm_enable <= perf_svm_disable); +} + +/** + * Test Description + * ------------------------ + * - This testcase measures performance when same memory chunk is repeatedly + * registered and unregistered on multiple GPUs. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_RegUnreg_Perf_SameChunk_MGPU") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string arch = prop.gcnArchName; + if (std::string::npos == arch.find("xnack+")) { + HipTest::HIP_SKIP_TEST("Xnack+ is not supported. Skipping the test ..."); + return; + } + int dev_count = HipTest::getDeviceCount(); + if (dev_count < 2) { + HipTest::HIP_SKIP_TEST("Only 1 GPU available. Skipping this test ..."); + return; + } + hip::SpawnProc proc("hipHostRegisterPerf", true); + REQUIRE(proc.run("svm_enable 2") == 0); + float perf_svm_enable = std::stof(proc.getOutput()); + INFO("perf_svm_enable: " << perf_svm_enable); + REQUIRE(proc.run("svm_disable 2") == 0); + float perf_svm_disable = std::stof(proc.getOutput()); + INFO("perf_svm_disable: " << perf_svm_disable); + REQUIRE(perf_svm_enable <= perf_svm_disable); +} + +/** + * Test Description + * ------------------------ + * - This testcase verifies whether hipMemAdvise can be used with + * host memory registered with hipHostRegister. + * registered and unregistered. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +TEST_CASE("Unit_hipHostRegister_MemAdvise_SetGet") { + // Execute the test only if xnack is supported + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string arch = prop.gcnArchName; + if ((std::string::npos == arch.find("xnack+")) || + (prop.concurrentManagedAccess == 0)) { + const char *msg = "Xnack/ConcurrentAccess not supported. Skipping test"; + HipTest::HIP_SKIP_TEST(msg); + return; + } + int numDevices = HipTest::getDeviceCount(); + size_t sizeBytes{LEN * sizeof(uint8_t)}; + uint8_t *A; + A = reinterpret_cast(malloc(sizeBytes)); + REQUIRE(A != nullptr); + memset(A, INITIAL_VAL, sizeBytes); + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + int out = 0; + SECTION("Attribute = hipMemAdviseSetReadMostly") { + HIP_CHECK(hipMemAdvise(A, sizeBytes, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, + A, sizeBytes)); + REQUIRE(out == 1); + HIP_CHECK(hipMemAdvise(A, sizeBytes, hipMemAdviseUnsetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, + A, sizeBytes)); + REQUIRE(out == 0); + } + SECTION("Attribute = hipMemAdviseSetPreferredLocation") { + HIP_CHECK(hipMemAdvise(A, sizeBytes, + hipMemAdviseSetPreferredLocation, hipCpuDeviceId)); + HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), + hipMemRangeAttributePreferredLocation, A, sizeBytes)); + REQUIRE(out == hipCpuDeviceId); + for (int dev = 0; dev < numDevices; dev++) { + HIP_CHECK(hipMemAdvise(A, sizeBytes, + hipMemAdviseSetPreferredLocation, dev)); + HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), + hipMemRangeAttributePreferredLocation, A, sizeBytes)); + REQUIRE(out == dev); + } + HIP_CHECK(hipMemAdvise(A, sizeBytes, + hipMemAdviseUnsetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), + hipMemRangeAttributePreferredLocation, A, sizeBytes)); + REQUIRE(out == hipInvalidDeviceId); + } + SECTION("Attribute = hipMemAdviseSetAccessedBy") { + size_t size = numDevices*sizeof(int); + int *chkOut = reinterpret_cast(malloc(size)); + HIP_CHECK(hipMemAdvise(A, sizeBytes, + hipMemAdviseSetAccessedBy, hipCpuDeviceId)); + for (int dev = 0; dev < numDevices; dev++) { + HIP_CHECK(hipMemAdvise(A, sizeBytes, + hipMemAdviseSetAccessedBy, dev)); + } + HIP_CHECK(hipMemRangeGetAttribute(chkOut, size, + hipMemRangeAttributeAccessedBy, A, sizeBytes)); + for (int dev = 0; dev < numDevices; dev++) { + REQUIRE(chkOut[dev] == dev); + } + free(chkOut); + } + HIP_CHECK(hipHostUnregister(A)); + free(A); +} +#endif +/** + * Test Description + * ------------------------ + * - This testcase verifies hipHostRegister API by performing memcpy + * on the hipHostRegistered variable. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", int, float, double) { // 1 refers to hipHostRegister // 0 refers to malloc auto mem_type = GENERATE(0, 1); HIP_CHECK(hipSetDevice(0)); - size_t sizeBytes = LEN * sizeof(TestType); TestType* A = reinterpret_cast(malloc(sizeBytes)); @@ -161,6 +896,17 @@ template __global__ void fill_kernel(T* dataPtr, T value) { dataPtr[tid] = value; } +/** + * Test Description + * ------------------------ + * - This testcase verifies all the supported flags of hipHostRegister. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Flags", "", int, float, double) { size_t sizeBytes = 1 * sizeof(TestType); TestType* hostPtr = reinterpret_cast(malloc(sizeBytes)); @@ -171,25 +917,40 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Flags", "", int, float, double) { bool valid; }; - /* EXSWCPHIPT-29 - 0x08 is hipHostRegisterReadOnly which currently doesn't have a definition in the headers */ - /* hipHostRegisterIoMemory is a valid flag but requires access to I/O mapped memory to be tested */ - FlagType flags = GENERATE( - FlagType{hipHostRegisterDefault, true}, FlagType{hipHostRegisterPortable, true}, - FlagType{0x08, true}, FlagType{hipHostRegisterPortable | hipHostRegisterMapped, true}, - FlagType{hipHostRegisterPortable | hipHostRegisterMapped | 0x08, true}, FlagType{0xF0, false}, - FlagType{0xFFF2, false}, FlagType{0xFFFFFFFF, false}); + /* EXSWCPHIPT-29 - 0x08 is hipHostRegisterReadOnly which currently doesn't + have a definition in the headers */ + /* hipHostRegisterIoMemory is a valid flag but requires access to I/O mapped + memory to be tested */ + FlagType flags = GENERATE(FlagType{hipHostRegisterDefault, true}, + FlagType{hipHostRegisterPortable, true}, + FlagType{0x08, true}, + FlagType{hipHostRegisterPortable | hipHostRegisterMapped, true}, + FlagType{hipHostRegisterPortable | hipHostRegisterMapped | 0x08, true}, + FlagType{0xF0, false}, + FlagType{0xFFF2, false}, FlagType{0xFFFFFFFF, false}); INFO("Testing hipHostRegister flag: " << flags.value); if (flags.valid) { HIP_CHECK(hipHostRegister(hostPtr, sizeBytes, flags.value)); HIP_CHECK(hipHostUnregister(hostPtr)); } else { - HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, flags.value), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, flags.value), + hipErrorInvalidValue); } - free(hostPtr); } +/** + * Test Description + * ------------------------ + * - These negative tests checks invalid parameter values. + * Test source + * ------------------------ + * - catch\unit\memory\hipHostRegister.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { TestType* hostPtr = nullptr; @@ -205,12 +966,14 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { size_t devMemAvail{0}, devMemFree{0}; HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail)); - auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes + auto hostMemFree = + HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes REQUIRE(devMemFree > 0); REQUIRE(devMemAvail > 0); REQUIRE(hostMemFree > 0); - size_t memFree = (std::max)(devMemFree, hostMemFree); // which is the limiter cpu or gpu + // which is the limiter cpu or gpu + size_t memFree = (std::max)(devMemFree, hostMemFree); SECTION("hipHostRegister Negative Test - invalid memory size") { HIP_CHECK_ERROR(hipHostRegister(hostPtr, memFree, 0), hipErrorInvalidValue); diff --git a/catch/unit/memory/hipHostRegister_exe.cc b/catch/unit/memory/hipHostRegister_exe.cc new file mode 100644 index 0000000000..40c42a82bb --- /dev/null +++ b/catch/unit/memory/hipHostRegister_exe.cc @@ -0,0 +1,155 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include // NOLINT +#include "hip/hip_runtime_api.h" + +#define ITERATION 1000 +#define SIZE (64*1024*1024) +#define ARRAY_SIZE 20 + +static bool UNSETENV(std::string var) { + int result = -1; +#ifdef __unix__ + result = unsetenv(var.c_str()); +#else + result = _putenv((var + '=').c_str()); +#endif + return (result == 0) ? true: false; +} + +static bool SETENV(std::string var, std::string value, int overwrite) { + int result = -1; +#ifdef __unix__ + result = setenv(var.c_str(), value.c_str(), overwrite); +#else + result = _putenv((var + '=' + value).c_str()); +#endif + return (result == 0) ? true: false; +} + +/** + Expects 2 command line arg, first command is flag svm_enable = 1/0 + and second command is test number: 0 = Register/Unregister different + chunks of host memory, 1 = Register/Unregister the same chunk of host + memory repeatedly, 2 = Register/Unregister the same chunk of host + memory repeatedly on multiple GPUs. +*/ +int main(int argc, char** argv) { + if (argc != 3) { + std::cerr << "Invalid number of args passed.\n" + << "argc : " << argc << std::endl; + return -1; + } + std::string env_flag = argv[1]; + int test = std::stoi(argv[2]); + // disable SVM feature using HSA_USE_SVM=0 env from shell + UNSETENV("HSA_USE_SVM"); + if (env_flag == "svm_enable") { + SETENV("HSA_USE_SVM", "1", 1); + } else { + SETENV("HSA_USE_SVM", "0", 1); + } + if (test == 0) { + uint8_t *A[ARRAY_SIZE]; + for (int i = 0; i < ARRAY_SIZE; i++) { + A[i] = reinterpret_cast(malloc(SIZE)); + if (A[i] == nullptr) { + return -1; + } + } + auto t1 = std::chrono::high_resolution_clock::now(); + for (int count = 0; count < ITERATION; count++) { + // Register the host pointer + if (hipSuccess != hipHostRegister(A[count%ARRAY_SIZE], SIZE, 0)) { + return -1; + } + // Unregister the host pointer + if (hipSuccess != hipHostUnregister(A[count%ARRAY_SIZE])) { + return -1; + } + } + auto t2 = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < ARRAY_SIZE; i++) { + free(A[i]); + } + std::chrono::duration fp_ms = t2 - t1; + std::cout << fp_ms.count() << std::endl; + } else if (test == 1) { + uint8_t *A; + A = reinterpret_cast(malloc(SIZE)); + if (A == nullptr) { + return -1; + } + auto t1 = std::chrono::high_resolution_clock::now(); + for (int count = 0; count < ITERATION; count++) { + // Register the host pointer + if (hipSuccess != hipHostRegister(A, SIZE, 0)) { + return -1; + } + // Unregister the host pointer + if (hipSuccess != hipHostUnregister(A)) { + return -1; + } + } + auto t2 = std::chrono::high_resolution_clock::now(); + free(A); + std::chrono::duration fp_ms = t2 - t1; + std::cout << fp_ms.count() << std::endl; + } else if (test == 2) { + uint8_t *A; + A = reinterpret_cast(malloc(SIZE)); + if (A == nullptr) { + return -1; + } + int dev_count = 0; + if (hipSuccess != hipGetDeviceCount(&dev_count)) { + return -1; + } + auto t1 = std::chrono::high_resolution_clock::now(); + for (int dev = 0; dev < dev_count; dev++) { + if (hipSuccess != hipSetDevice(dev)) { + return -1; + } + for (int count = 0; count < ITERATION; count++) { + // Register the host pointer + if (hipSuccess != hipHostRegister(A, SIZE, 0)) { + return -1; + } + // Unregister the host pointer + if (hipSuccess != hipHostUnregister(A)) { + return -1; + } + } + } + auto t2 = std::chrono::high_resolution_clock::now(); + free(A); + std::chrono::duration fp_ms = t2 - t1; + std::cout << fp_ms.count() << std::endl; + } else { + // Undefined test + } + UNSETENV("HSA_USE_SVM"); + return 0; +}