From bcdb3a7eced6bb944480282f95fb193e49ca2519 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 30 Nov 2022 12:58:13 +0100 Subject: [PATCH] EXSWHTEC-81 - Implement tests for hipGetSymbolAddress and hipGetSymbolSize (#3002) - Implement negative tests for hipGetSymbolAddress - Implement negative tests for hipGetSymbolSize - Reimplement positive test for both apis - Expand positive test with memcpy to and from symbol - Disable test sections that cause a segfault in CUDA --- catch/include/hip_array_common.hh | 84 +++++++++++++ catch/include/resource_guards.hh | 113 ++++++++++++++++- catch/include/utils.hh | 43 +++++++ catch/unit/memory/CMakeLists.txt | 2 + catch/unit/memory/hipArray3DCreate.cc | 1 + catch/unit/memory/hipArrayCommon.hh | 60 --------- catch/unit/memory/hipArrayCreate.cc | 1 + catch/unit/memory/hipFree.cc | 1 + catch/unit/memory/hipGetSymbolSizeAddress.cc | 123 +++++++++++++++++++ catch/unit/memory/hipMallocArray.cc | 1 + 10 files changed, 365 insertions(+), 64 deletions(-) create mode 100644 catch/include/hip_array_common.hh create mode 100644 catch/unit/memory/hipGetSymbolSizeAddress.cc diff --git a/catch/include/hip_array_common.hh b/catch/include/hip_array_common.hh new file mode 100644 index 0000000000..fd6f094f8d --- /dev/null +++ b/catch/include/hip_array_common.hh @@ -0,0 +1,84 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include + +template struct type_and_size_and_format { + using type = T; + static constexpr size_t size = N; + static constexpr hipArray_Format format = Format; +}; + +// Create a map of type to scalar type, vector size and scalar type format enum. +// This is useful for creating simpler function that depend on the vector size. +template struct vector_info; +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; \ No newline at end of file diff --git a/catch/include/resource_guards.hh b/catch/include/resource_guards.hh index 7e6179c81a..a9c7512a3d 100644 --- a/catch/include/resource_guards.hh +++ b/catch/include/resource_guards.hh @@ -19,6 +19,7 @@ THE SOFTWARE. #pragma once +#include #include #include @@ -80,10 +81,8 @@ template class LinearAllocGuard { } } - T* ptr() { return ptr_; }; - T* const ptr() const { return ptr_; }; - T* host_ptr() { return host_ptr_; } - T* const host_ptr() const { return host_ptr(); } + T* ptr() const { return ptr_; }; + T* host_ptr() const { return host_ptr_; } private: const LinearAllocs allocation_type_; @@ -91,6 +90,112 @@ template class LinearAllocGuard { T* host_ptr_ = nullptr; }; +template class LinearAllocGuardMultiDim { + protected: + LinearAllocGuardMultiDim(hipExtent extent) : extent_{extent} {} + + ~LinearAllocGuardMultiDim() { static_cast(hipFree(pitched_ptr_.ptr)); } + + public: + T* ptr() const { return reinterpret_cast(pitched_ptr_.ptr); }; + + size_t pitch() const { return pitched_ptr_.pitch; } + + hipExtent extent() const { return extent_; } + + hipPitchedPtr pitched_ptr() const { return pitched_ptr_; } + + size_t width() const { return extent_.width; } + + size_t width_logical() const { return extent_.width / sizeof(T); } + + size_t height() const { return extent_.height; } + + public: + hipPitchedPtr pitched_ptr_; + const hipExtent extent_; +}; + +template class LinearAllocGuard2D : public LinearAllocGuardMultiDim { + public: + LinearAllocGuard2D(const size_t width_logical, const size_t height) + : LinearAllocGuardMultiDim{make_hipExtent(width_logical * sizeof(T), height, 1)} { + HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch, + this->extent_.width, this->extent_.height)); + } + + LinearAllocGuard2D(const LinearAllocGuard2D&) = delete; + LinearAllocGuard2D(LinearAllocGuard2D&&) = delete; +}; + +template class LinearAllocGuard3D : public LinearAllocGuardMultiDim { + public: + LinearAllocGuard3D(const size_t width_logical, const size_t height, const size_t depth) + : LinearAllocGuardMultiDim{make_hipExtent(width_logical * sizeof(T), height, depth)} { + HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_)); + } + + LinearAllocGuard3D(const hipExtent extent) : LinearAllocGuardMultiDim(extent) { + HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_)); + } + + LinearAllocGuard3D(const LinearAllocGuard3D&) = delete; + LinearAllocGuard3D(LinearAllocGuard3D&&) = delete; + + size_t depth() const { return this->extent_.depth; } +}; + +template class ArrayAllocGuard { + public: + // extent should contain logical width + ArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMalloc3DArray(&ptr_, &desc, extent_, flags)); + } + + ~ArrayAllocGuard() { static_cast(hipFreeArray(ptr_)); } + + ArrayAllocGuard(const ArrayAllocGuard&) = delete; + ArrayAllocGuard(ArrayAllocGuard&&) = delete; + + hipArray_t ptr() const { return ptr_; } + + hipExtent extent() const { return extent_; } + + private: + hipArray_t ptr_ = nullptr; + const hipExtent extent_; +}; + +template class DrvArrayAllocGuard { + public: + // extent should contain width in bytes + DrvArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} { + HIP_ARRAY3D_DESCRIPTOR desc{}; + using vec_info = vector_info; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + desc.Width = extent_.width / sizeof(T); + desc.Height = extent_.height; + desc.Depth = extent_.depth; + desc.Flags = flags; + HIP_CHECK(hipArray3DCreate(&ptr_, &desc)); + } + + ~DrvArrayAllocGuard() { static_cast(hipArrayDestroy(ptr_)); } + + DrvArrayAllocGuard(const DrvArrayAllocGuard&) = delete; + DrvArrayAllocGuard(DrvArrayAllocGuard&&) = delete; + + hiparray ptr() const { return ptr_; } + + hipExtent extent() const { return extent_; } + + private: + hiparray ptr_ = nullptr; + const hipExtent extent_; +}; + enum class Streams { nullstream, perThread, created }; class StreamGuard { diff --git a/catch/include/utils.hh b/catch/include/utils.hh index 9edffc6f7c..bbab2322fe 100644 --- a/catch/include/utils.hh +++ b/catch/include/utils.hh @@ -54,6 +54,37 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele ArrayFindIfNot(array, array + num_elements, expected_value); } +template +void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height, + const size_t depth, F expected_value_generator) { + for (size_t z = 0; z < depth; ++z) { + for (size_t y = 0; y < height; ++y) { + for (size_t x = 0; x < width; ++x) { + const auto slice = reinterpret_cast(ptr) + pitch * height * z; + const auto row = slice + pitch * y; + if (reinterpret_cast(row)[x] != expected_value_generator(x, y, z)) { + INFO("Mismatch at indices: " << x << ", " << y << ", " << z); + REQUIRE(reinterpret_cast(row)[x] == expected_value_generator(x, y, z)); + } + } + } + } +} + +template +void PitchedMemorySet(T* const ptr, const size_t pitch, const size_t width, const size_t height, + const size_t depth, F expected_value_generator) { + for (size_t z = 0; z < depth; ++z) { + for (size_t y = 0; y < height; ++y) { + for (size_t x = 0; x < width; ++x) { + const auto slice = reinterpret_cast(ptr) + pitch * height * z; + const auto row = slice + pitch * y; + reinterpret_cast(row)[x] = expected_value_generator(x, y, z); + } + } + } +} + template __global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) { size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); @@ -82,6 +113,18 @@ static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { } } +template +__global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) { + const auto x = blockIdx.x * blockDim.x + threadIdx.x; + const auto y = blockIdx.y * blockDim.y + threadIdx.y; + const auto z = blockIdx.z * blockDim.z + threadIdx.z; + if (x < w && y < h && z < d) { + char* const slice = reinterpret_cast(out) + pitch * h * z; + char* const row = slice + pitch * y; + reinterpret_cast(row)[x] = z * w * h + y * w + x; + } +} + inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) { int ticks_per_ms = 0; // Clock rate is in kHz => number of clock ticks in a millisecond diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 4d2d74c033..0f50348239 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -98,6 +98,7 @@ set(TEST_SRC hipMemsetAsync.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc + hipGetSymbolSizeAddress.cc ) else() set(TEST_SRC @@ -172,6 +173,7 @@ set(TEST_SRC hipMemsetAsync.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc + hipGetSymbolSizeAddress.cc ) endif() diff --git a/catch/unit/memory/hipArray3DCreate.cc b/catch/unit/memory/hipArray3DCreate.cc index 973868eded..4cf189611b 100644 --- a/catch/unit/memory/hipArray3DCreate.cc +++ b/catch/unit/memory/hipArray3DCreate.cc @@ -20,6 +20,7 @@ THE SOFTWARE. #include #include "DriverContext.hh" #include "hipArrayCommon.hh" +#include "hip_array_common.hh" #include "hip_test_common.hh" namespace { diff --git a/catch/unit/memory/hipArrayCommon.hh b/catch/unit/memory/hipArrayCommon.hh index b40014b490..b0beeb3126 100644 --- a/catch/unit/memory/hipArrayCommon.hh +++ b/catch/unit/memory/hipArrayCommon.hh @@ -26,66 +26,6 @@ THE SOFTWARE. constexpr size_t BlockSize = 16; -template struct type_and_size_and_format { - using type = T; - static constexpr size_t size = N; - static constexpr hipArray_Format format = Format; -}; - -// Create a map of type to scalar type, vector size and scalar type format enum. -// This is useful for creating simpler function that depend on the vector size. -template struct vector_info; -template <> -struct vector_info : type_and_size_and_format {}; -template <> struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; - -template <> -struct vector_info : type_and_size_and_format {}; -template <> struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; - -template <> -struct vector_info : type_and_size_and_format {}; -template <> struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; -template <> -struct vector_info - : type_and_size_and_format {}; - // read from a texture using normalized coordinates constexpr size_t ChannelToRead = 1; template diff --git a/catch/unit/memory/hipArrayCreate.cc b/catch/unit/memory/hipArrayCreate.cc index 6cc535593a..70a8636922 100644 --- a/catch/unit/memory/hipArrayCreate.cc +++ b/catch/unit/memory/hipArrayCreate.cc @@ -27,6 +27,7 @@ hipArrayCreate API test scenarios #include #include #include +#include #include "hipArrayCommon.hh" #include "DriverContext.hh" diff --git a/catch/unit/memory/hipFree.cc b/catch/unit/memory/hipFree.cc index 1248deebc1..b29854271c 100644 --- a/catch/unit/memory/hipFree.cc +++ b/catch/unit/memory/hipFree.cc @@ -22,6 +22,7 @@ THE SOFTWARE. #include +#include #include "hipArrayCommon.hh" #include "DriverContext.hh" diff --git a/catch/unit/memory/hipGetSymbolSizeAddress.cc b/catch/unit/memory/hipGetSymbolSizeAddress.cc new file mode 100644 index 0000000000..5c011c7e81 --- /dev/null +++ b/catch/unit/memory/hipGetSymbolSizeAddress.cc @@ -0,0 +1,123 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +#include +#include +#include +#include + +namespace { +constexpr size_t kArraySize = 5; +} // anonymous namespace + +#define HIP_GET_SYMBOL_SIZE_ADDRESS_DEFINE_GLOBALS(type) \ + __device__ type type##_var = 0; \ + __device__ type type##_arr[kArraySize] = {}; \ + __global__ void type##_var_address_validation_kernel(void* ptr, bool* out) { \ + *out = static_cast(&type##_var) == ptr; \ + } \ + __global__ void type##_arr_address_validation_kernel(void* ptr, bool* out) { \ + *out = static_cast(type##_arr) == ptr; \ + } + +HIP_GET_SYMBOL_SIZE_ADDRESS_DEFINE_GLOBALS(int) +HIP_GET_SYMBOL_SIZE_ADDRESS_DEFINE_GLOBALS(float) +HIP_GET_SYMBOL_SIZE_ADDRESS_DEFINE_GLOBALS(char) +HIP_GET_SYMBOL_SIZE_ADDRESS_DEFINE_GLOBALS(double) + +template +static void HipGetSymbolSizeAddressTest(const void* symbol) { + constexpr auto size = N * sizeof(T); + + T* symbol_ptr = nullptr; + size_t symbol_size = 0; + HIP_CHECK(hipGetSymbolAddress(reinterpret_cast(&symbol_ptr), symbol)); + HIP_CHECK(hipGetSymbolSize(&symbol_size, symbol)); + REQUIRE(symbol_size == size); + REQUIRE(symbol_ptr != nullptr); + + LinearAllocGuard equal_addresses(LinearAllocs::hipMalloc, sizeof(bool)); + HIP_CHECK(hipMemset(equal_addresses.ptr(), false, sizeof(*equal_addresses.ptr()))) + validation_kernel<<<1, 1>>>(symbol_ptr, equal_addresses.ptr()); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipStreamSynchronize(nullptr)); + bool ok = false; + HIP_CHECK(hipMemcpy(&ok, equal_addresses.ptr(), sizeof(ok), hipMemcpyDeviceToHost)); + REQUIRE(ok); + + constexpr T expected_value = 42; + std::array fill_buffer; + std::fill_n(fill_buffer.begin(), N, expected_value); + HIP_CHECK(hipMemcpy(symbol_ptr, fill_buffer.data(), symbol_size, hipMemcpyHostToDevice)); + + + std::array read_buffer; + HIP_CHECK(hipMemcpy(read_buffer.data(), symbol_ptr, symbol_size, hipMemcpyDeviceToHost)); + ArrayFindIfNot(read_buffer.data(), expected_value, read_buffer.size()); +} + +#if HT_AMD +#define SYMBOL(expr) &HIP_SYMBOL(expr) +#else +#define SYMBOL(expr) HIP_SYMBOL(expr) +#endif + +#define HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(type) \ + HipGetSymbolSizeAddressTest(SYMBOL(type##_var)); \ + HipGetSymbolSizeAddressTest( \ + SYMBOL(type##_arr)); + +TEST_CASE("Unit_hipGetSymbolSizeAddress_Positive_Basic") { + SECTION("int") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(int); } + SECTION("float") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(float); } + SECTION("char") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(char); } + SECTION("double") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(double); } +} + +TEST_CASE("Unit_hipGetSymbolAddress_Negative_Parameters") { +// Causes a segfault in CUDA +#if HT_AMD + SECTION("devPtr == nullptr") { + HIP_CHECK_ERROR(hipGetSymbolAddress(nullptr, SYMBOL(int_var)), hipErrorInvalidValue); + } +#endif + + SECTION("symbolName == nullptr") { + void* ptr = nullptr; + HIP_CHECK_ERROR(hipGetSymbolAddress(&ptr, nullptr), hipErrorInvalidSymbol); + } +} + +TEST_CASE("Unit_hipGetSymbolSize_Negative_Parameters") { +// Causes a segfault in CUDA +#if HT_AMD + SECTION("size == nullptr") { + HIP_CHECK_ERROR(hipGetSymbolSize(nullptr, SYMBOL(int_var)), hipErrorInvalidValue); + } +#endif + + SECTION("symbolName == nullptr") { + size_t size = 0; + HIP_CHECK_ERROR(hipGetSymbolSize(&size, nullptr), hipErrorInvalidSymbol); + } +} \ No newline at end of file diff --git a/catch/unit/memory/hipMallocArray.cc b/catch/unit/memory/hipMallocArray.cc index b6c4939b1e..530eb11077 100644 --- a/catch/unit/memory/hipMallocArray.cc +++ b/catch/unit/memory/hipMallocArray.cc @@ -26,6 +26,7 @@ hipMallocArray API test scenarios */ #include +#include #include #include #include "hipArrayCommon.hh"