From eafb3a23ee4957681d5e58bc39acd75dcf1abd71 Mon Sep 17 00:00:00 2001 From: Dylan Angus <61192377+dylan-angus-codeplay@users.noreply.github.com> Date: Mon, 11 Jul 2022 07:48:19 +0100 Subject: [PATCH] EXSWCPHIPT-77 - Extending tests for hipHostRegister (#2609) --- catch/include/hip_test_helper.hh | 2 +- catch/unit/memory/hipHostRegister.cc | 118 ++++++++++++++++++--------- 2 files changed, 80 insertions(+), 40 deletions(-) diff --git a/catch/include/hip_test_helper.hh b/catch/include/hip_test_helper.hh index fbbcb6cb06..9d4cbcd73f 100644 --- a/catch/include/hip_test_helper.hh +++ b/catch/include/hip_test_helper.hh @@ -38,7 +38,7 @@ static inline int getGeviceCount() { } // Get Free Memory from the system -static size_t getMemoryAmount() { +static inline size_t getMemoryAmount() { #ifdef __linux__ struct sysinfo info{}; sysinfo(&info); diff --git a/catch/unit/memory/hipHostRegister.cc b/catch/unit/memory/hipHostRegister.cc index 8312cf7f50..c2fa6ed009 100644 --- a/catch/unit/memory/hipHostRegister.cc +++ b/catch/unit/memory/hipHostRegister.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -28,51 +28,49 @@ This testfile verifies the following scenarios of hipHostRegister API */ #include -#include -#include +#include #define OFFSET 128 -static constexpr auto LEN{1024*1024}; +static constexpr auto LEN{1024 * 1024}; -template -__global__ void Inc(T* Ad) { - int tx = threadIdx.x + blockIdx.x * blockDim.x; - Ad[tx] = Ad[tx] + static_cast(1); +template __global__ void Inc(T* Ad) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tx] = Ad[tx] + static_cast(1); } template -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; +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; - size_t sizeBytes = numElements * sizeof(T); + size_t sizeBytes = numElements * sizeof(T); - if (internalRegister) { - HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); - } + if (internalRegister) { + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + } - // Reset - for (size_t i = 0; i < numElements; i++) { - A[i] = static_cast(i); - Bh[i] = 0.0f; - } + // Reset + for (size_t i = 0; i < numElements; i++) { + A[i] = static_cast(i); + Bh[i] = 0.0f; + } - HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes)); + HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes)); - HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); - // Make sure the copy worked - for (size_t i = 0; i < numElements; i++) { - REQUIRE(Bh[i] == A[i]); - } + // Make sure the copy worked + for (size_t i = 0; i < numElements; i++) { + REQUIRE(Bh[i] == A[i]); + } - if (internalRegister) { - HIP_CHECK(hipHostUnregister(A)); - } + if (internalRegister) { + HIP_CHECK(hipHostUnregister(A)); + } } + /* This testcase verifies the hipHostRegister API by 1. Allocating the memory using malloc @@ -81,9 +79,7 @@ This testcase verifies the hipHostRegister API by 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) { +TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", int, float, double) { size_t sizeBytes{LEN * sizeof(TestType)}; TestType *A, **Ad; int num_devices; @@ -118,14 +114,14 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", HIP_CHECK(hipHostUnregister(A)); free(A); - delete [] Ad; + delete[] Ad; } + /* This testcase verifies hipHostRegister API by performing memcpy on the hipHostRegistered variable. */ -TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", - int, float, double) { +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", int, float, double) { // 1 refers to hipHostRegister // 0 refers to malloc auto mem_type = GENERATE(0, 1); @@ -156,5 +152,49 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", free(A); free(Bh); - hipFree(Bd); + HIP_CHECK(hipFree(Bd)); +} + +template __global__ void fill_kernel(T* dataPtr, T value) { + size_t tid{blockIdx.x * blockDim.x + threadIdx.x}; + dataPtr[tid] = value; +} + +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { + TestType* hostPtr = nullptr; + + size_t sizeBytes = 1 * sizeof(TestType); + SECTION("hipHostRegister Negative Test - nullptr") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 1, 0), hipErrorInvalidValue); + } + + hostPtr = reinterpret_cast(malloc(sizeBytes)); + SECTION("hipHostRegister Negative Test - zero size") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); + } + +#if HT_NVIDIA + // Flags aren't used for AMD devices currently + SECTION("hipHostRegister Negative Test - invalid flag") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, 0b11111111), hipErrorInvalidValue); + } +#endif + + size_t devMemAvail{0}, devMemFree{0}; + HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail)); + auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes + REQUIRE(devMemFree > 0); + REQUIRE(devMemAvail > 0); + REQUIRE(hostMemFree > 0); + + size_t memFree = std::min(devMemFree, hostMemFree); // which is the limiter cpu or gpu + + SECTION("hipHostRegister Negative Test - invalid memory size") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, memFree, 0), hipErrorInvalidValue); + } + + free(hostPtr); + SECTION("hipHostRegister Negative Test - freed memory") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); + } }