diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json index 9704e80d49..f3a751f348 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux_common.json @@ -98,6 +98,8 @@ "=== Below tests fail in stress test on 30/06/23 ===", "Unit_hipStreamValue_Wait32_Blocking_NoMask_Nor", "Unit_hipStreamValue_Write - TestParams", - "Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice" + "Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/96 ===", + "Unit_hipHostGetDevicePointer_Negative" ] } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json index 2d050a9ca4..4d7fdff960 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -185,6 +185,8 @@ "SWDEV-398981 fails in stress test", "Unit_hipStreamCreateWithPriority_MulthreadDefaultflag", "Note: UUID returned empty on some windows nodes", - "Unit_hipDeviceGetUuid_Positive" + "Unit_hipDeviceGetUuid_Positive", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/96 ===", + "Unit_hipHostGetDevicePointer_Negative" ] } diff --git a/projects/hip-tests/catch/unit/memory/hipFree.cc b/projects/hip-tests/catch/unit/memory/hipFree.cc index b46f8ae8fb..1db0ec9502 100644 --- a/projects/hip-tests/catch/unit/memory/hipFree.cc +++ b/projects/hip-tests/catch/unit/memory/hipFree.cc @@ -49,11 +49,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 @@ -62,11 +61,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 @@ -75,7 +74,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; @@ -135,7 +134,6 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo } } -#endif #endif // Freeing a invalid pointer with on device @@ -166,8 +164,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/projects/hip-tests/catch/unit/memory/hipHostGetDevicePointer.cc b/projects/hip-tests/catch/unit/memory/hipHostGetDevicePointer.cc index 7c3e689e05..7f07468935 100644 --- a/projects/hip-tests/catch/unit/memory/hipHostGetDevicePointer.cc +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/hipHostRegister.cc b/projects/hip-tests/catch/unit/memory/hipHostRegister.cc index f6964db616..5e1b10d234 100644 --- a/projects/hip-tests/catch/unit/memory/hipHostRegister.cc +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/hipHostUnregister.cc b/projects/hip-tests/catch/unit/memory/hipHostUnregister.cc index 69373133d0..ea3d018a33 100644 --- a/projects/hip-tests/catch/unit/memory/hipHostUnregister.cc +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/hipMallocPitch.cc b/projects/hip-tests/catch/unit/memory/hipMallocPitch.cc index 19c803c7c2..7b26c16034 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocPitch.cc +++ b/projects/hip-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{}; @@ -360,42 +375,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 @@ -408,6 +388,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)); } @@ -538,4 +519,3 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", "" HipTest::freeArrays(nullptr, nullptr, nullptr, A_h, B_h, C_h, false); } - diff --git a/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc b/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc index 393221da11..b0e1e7a5f8 100644 --- a/projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc +++ b/projects/hip-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"