diff --git a/tests/catch/include/resource_guards.hh b/tests/catch/include/resource_guards.hh index 7e6179c81a..b3ef7813f7 100644 --- a/tests/catch/include/resource_guards.hh +++ b/tests/catch/include/resource_guards.hh @@ -80,10 +80,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 +89,65 @@ 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; } +}; + enum class Streams { nullstream, perThread, created }; class StreamGuard { diff --git a/tests/catch/include/utils.hh b/tests/catch/include/utils.hh index 9edffc6f7c..05eecea79f 100644 --- a/tests/catch/include/utils.hh +++ b/tests/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 (int z = 0; z < depth; ++z) { + for (int y = 0; y < height; ++y) { + for (int 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 (int z = 0; z < depth; ++z) { + for (int y = 0; y < height; ++y) { + for (int 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/tests/catch/unit/memory/hipFree.cc b/tests/catch/unit/memory/hipFree.cc index 1248deebc1..018b95b9c3 100644 --- a/tests/catch/unit/memory/hipFree.cc +++ b/tests/catch/unit/memory/hipFree.cc @@ -48,11 +48,10 @@ using namespace std::chrono_literals; const std::chrono::duration delay = 50ms; constexpr size_t numAllocs = 10; -#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ -TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float4) { - TestType* devPtr{}; +TEST_CASE("Unit_hipFreeImplicitSyncDev") { + int* devPtr{}; size_t size_mult = GENERATE(1, 32, 64, 128, 256); - HIP_CHECK(hipMalloc(&devPtr, sizeof(TestType) * size_mult)); + HIP_CHECK(hipMalloc(&devPtr, sizeof(*devPtr) * size_mult)); HipTest::runKernelForDuration(delay); // make sure device is busy @@ -61,11 +60,11 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float HIP_CHECK(hipStreamQuery(nullptr)); } -TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, float4) { - TestType* hostPtr{}; +TEST_CASE("Unit_hipFreeImplicitSyncHost") { + int* hostPtr{}; size_t size_mult = GENERATE(1, 32, 64, 128, 256); - HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(TestType) * size_mult)); + HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(*hostPtr) * size_mult)); HipTest::runKernelForDuration(delay); // make sure device is busy @@ -74,7 +73,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, floa HIP_CHECK(hipStreamQuery(nullptr)); } -#if HT_NVIDIA // Meaningless at the moment, since we are not running wait kernel on nvidia. +#if HT_NVIDIA TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, float4) { using vec_info = vector_info; DriverContext ctx; @@ -134,7 +133,6 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo } } -#endif #endif // Freeing a invalid pointer with on device @@ -165,8 +163,6 @@ TEST_CASE("Unit_hipFreeNegativeHost") { #if HT_NVIDIA TEST_CASE("Unit_hipFreeNegativeArray") { DriverContext ctx; - hipArray_t arrayPtr{}; - hiparray cuArrayPtr{}; SECTION("ArrayFree") { HIP_CHECK(hipFreeArray(nullptr)); } SECTION("ArrayDestroy") { diff --git a/tests/catch/unit/memory/hipHostGetDevicePointer.cc b/tests/catch/unit/memory/hipHostGetDevicePointer.cc index 7c3e689e05..7f07468935 100644 --- a/tests/catch/unit/memory/hipHostGetDevicePointer.cc +++ b/tests/catch/unit/memory/hipHostGetDevicePointer.cc @@ -21,11 +21,19 @@ THE SOFTWARE. */ #include +#include TEST_CASE("Unit_hipHostGetDevicePointer_Negative") { int* hPtr{nullptr}; + int* dPtr{nullptr}; HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); + if (!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) { + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 0), + hipErrorNotSupported); + return; + } + SECTION("Nullptr as device") { HIP_CHECK_ERROR(hipHostGetDevicePointer(nullptr, hPtr, 0), hipErrorInvalidValue); } @@ -36,13 +44,29 @@ TEST_CASE("Unit_hipHostGetDevicePointer_Negative") { hipErrorInvalidValue); } - // Not adding check for flags since CUDA spec states that there might be more values added to it + SECTION("Non pinned memory as host") { + int* hPtr = reinterpret_cast(malloc(sizeof(*hPtr))); + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 0), + hipErrorInvalidValue); + free(hPtr); + } + + SECTION("Flags non-zero") { + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 1), + hipErrorInvalidValue); + } + HIP_CHECK(hipHostFree(hPtr)); } template __global__ void set(T* ptr, T val) { *ptr = val; } TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") { + if(!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) { + HipTest::HIP_SKIP_TEST("Device does not support mapping host memory"); + return; + } + int* hPtr{nullptr}; HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); @@ -71,8 +95,8 @@ TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") { HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipHostUnregister(&res)); - REQUIRE(value == 10); + REQUIRE(res == value); } HIP_CHECK(hipHostFree(hPtr)); -} +} \ No newline at end of file diff --git a/tests/catch/unit/memory/hipHostRegister.cc b/tests/catch/unit/memory/hipHostRegister.cc index f6964db616..5e1b10d234 100644 --- a/tests/catch/unit/memory/hipHostRegister.cc +++ b/tests/catch/unit/memory/hipHostRegister.cc @@ -27,9 +27,10 @@ This testfile verifies the following scenarios of hipHostRegister API 2. hipHostRegister and perform hipMemcpy on it. */ +#include "hip/hip_runtime_api.h" #include #include -#include "hip/hip_runtime_api.h" +#include #define OFFSET 128 static constexpr auto LEN{1024 * 1024}; @@ -63,9 +64,7 @@ void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internal 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]); - } + ArrayMismatch(A, Bh, numElements); if (internalRegister) { HIP_CHECK(hipHostUnregister(A)); diff --git a/tests/catch/unit/memory/hipHostUnregister.cc b/tests/catch/unit/memory/hipHostUnregister.cc index 69373133d0..ea3d018a33 100644 --- a/tests/catch/unit/memory/hipHostUnregister.cc +++ b/tests/catch/unit/memory/hipHostUnregister.cc @@ -68,6 +68,12 @@ TEST_CASE("Unit_hipHostUnregister_NullPtr") { HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue); } +TEST_CASE("Unit_hipHostUnregister_Ptr_Different_Than_Specified_To_Register") { + std::vector alloc(2); + HIP_CHECK(hipHostRegister(alloc.data(), alloc.size(), 0)); + HIP_CHECK_ERROR(hipHostUnregister(&alloc.data()[1]), hipErrorHostMemoryNotRegistered); +} + TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") { auto x = std::unique_ptr(new int); HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); diff --git a/tests/catch/unit/memory/hipMallocPitch.cc b/tests/catch/unit/memory/hipMallocPitch.cc index 5a20671e14..b84b45087b 100644 --- a/tests/catch/unit/memory/hipMallocPitch.cc +++ b/tests/catch/unit/memory/hipMallocPitch.cc @@ -228,6 +228,21 @@ TEST_CASE("Unit_hipMallocPitch_Negative") { } } +TEST_CASE("Unit_hipMallocPitch_Zero_Dims") { + void* ptr = nullptr; + size_t pitch = 0; + + SECTION("width == 0") { + HIP_CHECK(hipMallocPitch(&ptr, &pitch, 0, 1)); + REQUIRE(ptr == nullptr); + } + + SECTION("height == 0") { + HIP_CHECK(hipMallocPitch(&ptr, &pitch, 1, 0)); + REQUIRE(ptr == nullptr); + } +} + TEST_CASE("Unit_hipMemAllocPitch_Negative") { size_t pitch = 0; hipDeviceptr_t ptr{}; @@ -366,42 +381,7 @@ static void MemoryAllocDiffSizes(int gpu) { static void threadFunc(int gpu) { MemoryAllocDiffSizes(gpu); } -/* - * This testcase verifies the negative scenarios of hipMallocPitch API - */ -#if 0 //TODO: Review, fix and re-enable test -TEST_CASE("Unit_hipMallocPitch_Negative") { - float* A_d; - size_t pitch_A = 0; - size_t width{NUM_W * sizeof(float)}; -#if HT_NVIDIA - SECTION("NullPtr to Pitched Ptr") { - REQUIRE(hipMallocPitch(nullptr, - &pitch_A, width, NUM_H) != hipSuccess); - } - - SECTION("nullptr to pitch") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - nullptr, width, NUM_H) != hipSuccess); - } -#endif - SECTION("Width 0 in hipMallocPitch") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, 0, NUM_H) == hipSuccess); - } - SECTION("Height 0 in hipMallocPitch") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, width, 0) == hipSuccess); - } - - SECTION("Max int values") { - REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, std::numeric_limits::max(), - std::numeric_limits::max()) != hipSuccess); - } -} -#endif /* * This testcase verifies the basic scenario of * hipMallocPitch API for different datatypes @@ -414,6 +394,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic", size_t width{NUM_W * sizeof(TestType)}; REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, NUM_H) == hipSuccess); + REQUIRE(width <= pitch_A); HIP_CHECK(hipFree(A_d)); } diff --git a/tests/catch/unit/memory/hipPointerGetAttribute.cc b/tests/catch/unit/memory/hipPointerGetAttribute.cc index 393221da11..b0e1e7a5f8 100644 --- a/tests/catch/unit/memory/hipPointerGetAttribute.cc +++ b/tests/catch/unit/memory/hipPointerGetAttribute.cc @@ -316,9 +316,8 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") { == hipErrorInvalidValue); } SECTION("Pass invalid attribute") { - hipPointer_attribute attr{HIP_POINTER_ATTRIBUTE_DEVICE_POINTER}; - REQUIRE(hipPointerGetAttribute(&data, attr, - reinterpret_cast(A_h)) == hipErrorInvalidValue); + REQUIRE(hipPointerGetAttribute(&data, static_cast(-1), + reinterpret_cast(A_h)) == hipErrorInvalidValue); } #if HT_AMD SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE"