From b64d0cff29b2c433e86ab6f9f500698810d211f4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Wed, 25 May 2022 14:37:20 +0100 Subject: [PATCH] hipMemset Negative Testing (#2639) --- catch/unit/memory/CMakeLists.txt | 6 +- catch/unit/memory/hipMemset3DNegative.cc | 236 ----------------------- catch/unit/memory/hipMemsetInvalidPtr.cc | 127 ------------ catch/unit/memory/hipMemsetNegative.cc | 229 ++++++++++++++++++++++ 4 files changed, 231 insertions(+), 367 deletions(-) delete mode 100644 catch/unit/memory/hipMemset3DNegative.cc delete mode 100644 catch/unit/memory/hipMemsetInvalidPtr.cc create mode 100644 catch/unit/memory/hipMemsetNegative.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index d4e821463c..3ce80dd780 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -42,14 +42,13 @@ set(TEST_SRC hipPointerGetAttributes.cc hipHostGetFlags.cc hipMallocManaged_MultiScenario.cc - hipMemsetInvalidPtr.cc + hipMemsetNegative.cc hipMemset.cc hipMemsetAsyncMultiThread.cc hipMemset3D.cc hipMemset2D.cc hipHostMallocTests.cc hipMemset3DFunctional.cc - hipMemset3DNegative.cc hipMemset3DRegressMultiThread.cc hipMallocManagedFlagsTst.cc hipMemPrefetchAsyncExtTsts.cc @@ -108,14 +107,13 @@ set(TEST_SRC hipHostRegister.cc hipHostGetFlags.cc hipMallocManaged_MultiScenario.cc - hipMemsetInvalidPtr.cc + hipMemsetNegative.cc hipMemset.cc hipMemsetAsyncMultiThread.cc hipMemset3D.cc hipMemset2D.cc hipHostMallocTests.cc hipMemset3DFunctional.cc - hipMemset3DNegative.cc hipMemset3DRegressMultiThread.cc hipMallocManagedFlagsTst.cc hipMemPrefetchAsyncExtTsts.cc diff --git a/catch/unit/memory/hipMemset3DNegative.cc b/catch/unit/memory/hipMemset3DNegative.cc deleted file mode 100644 index 798c90e1ac..0000000000 --- a/catch/unit/memory/hipMemset3DNegative.cc +++ /dev/null @@ -1,236 +0,0 @@ -/* -Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/** -Testcase Scenarios : - 1) Test hipMemset3D() with uninitialized devPitchedPtr. - 2) Test hipMemset3DAsync() with uninitialized devPitchedPtr. - - 3) Reset devPitchedPtr to zero and check return value for hipMemset3D(). - 4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync(). - - 5) Test hipMemset3D() with extent.width as max size_t and keeping height, - depth as valid values. - 6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height, - depth as valid values. - 7) Test hipMemset3D() with extent.height as max size_t and keeping width, - depth as valid values. - 8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width, - depth as valid values. - 9) Test hipMemset3D() with extent.depth as max size_t and keeping height, - width as valid values. -10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height, - width as valid values. - -11) Device Ptr out bound and extent(0) passed for hipMemset3D(). -12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync(). - -13) Device Ptr out bound and valid extent passed for hipMemset3D(). -14) Device Ptr out bound and valid extent passed for hipMemset3DAsync(). -*/ - -#include - -TEST_CASE("Unit_hipMemset3D_Negative") { - hipError_t ret; - hipPitchedPtr devPitchedPtr; - constexpr int memsetval = 1; - constexpr size_t numH = 256, numW = 256; - constexpr size_t depth = 10; - constexpr size_t width = numW * sizeof(char); - hipExtent extent = make_hipExtent(width, numH, depth); - - HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); - - SECTION("Using uninitialized devpitched ptr") { - hipPitchedPtr devPitchedUnPtr; - - ret = hipMemset3D(devPitchedUnPtr, memsetval, extent); - REQUIRE(ret != hipSuccess); - } - - SECTION("Reset devPitchedPtr to zero") { - hipPitchedPtr rdevPitchedPtr{}; - - ret = hipMemset3D(rdevPitchedPtr, memsetval, extent); - REQUIRE(ret != hipSuccess); - } - - SECTION("Pass extent fields as max size_t") { - hipExtent extMW = make_hipExtent(std::numeric_limits::max(), - numH, - depth); - hipExtent extMH = make_hipExtent(width, - std::numeric_limits::max(), - depth); - hipExtent extMD = make_hipExtent(width, - numH, - std::numeric_limits::max()); - - ret = hipMemset3D(devPitchedPtr, memsetval, extMW); - REQUIRE(ret != hipSuccess); - - ret = hipMemset3D(devPitchedPtr, memsetval, extMH); - REQUIRE(ret != hipSuccess); - - if ((TestContext::get()).isAmd()) { - ret = hipMemset3D(devPitchedPtr, memsetval, extMD); - REQUIRE(ret != hipSuccess); - } else { - WARN("Test is skipped for max depth." - << "Cuda doesn't check the maximum depth of extent field"); - } - } - - SECTION("Device Ptr out bound and extent(0) passed for memset") { - size_t pitch = devPitchedPtr.pitch; - size_t slicePitch = pitch * extent.height; - constexpr auto advanceOffset = 10; - - // Point devptr to end of allocated memory - char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) - + depth * slicePitch; - - // Advance devptr further to go out of boundary - devPtrMod = devPtrMod + advanceOffset; - hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, - numW * sizeof(char), numH); - hipExtent extent0{}; - ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0); - - // api expected to check extent0 and return success before going for ptr. - REQUIRE(ret == hipSuccess); - } - - SECTION("Device Ptr out bound and valid extent passed for memset") { - size_t pitch = devPitchedPtr.pitch; - size_t slicePitch = pitch * extent.height; - constexpr auto advanceOffset = 10; - - // Point devptr to end of allocated memory - char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) - + depth * slicePitch; - - // Advance devptr further to go out of boundary - devPtrMod = devPtrMod + advanceOffset; - hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, - numW * sizeof(char), numH); - ret = hipMemset3D(modDevPitchedPtr, memsetval, extent); - - REQUIRE(ret != hipSuccess); - } - - HIP_CHECK(hipFree(devPitchedPtr.ptr)); -} - -TEST_CASE("Unit_hipMemset3DAsync_Negative") { - hipError_t ret; - hipPitchedPtr devPitchedPtr; - hipStream_t stream; - constexpr int memsetval = 1; - constexpr size_t numH = 256; - constexpr size_t numW = 256; - constexpr size_t depth = 10; - constexpr size_t width = numW * sizeof(char); - hipExtent extent = make_hipExtent(width, numH, depth); - - HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); - - SECTION("Using uninitialized devpitched ptr") { - hipPitchedPtr devPitchedUnPtr; - - ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream); - REQUIRE(ret != hipSuccess); - } - - SECTION("Reset devPitchedPtr to zero") { - hipPitchedPtr rdevPitchedPtr{}; - - ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream); - REQUIRE(ret != hipSuccess); - } - - SECTION("Pass extent fields as max size_t") { - hipExtent extMW = make_hipExtent(std::numeric_limits::max(), - numH, - depth); - hipExtent extMH = make_hipExtent(width, - std::numeric_limits::max(), - depth); - hipExtent extMD = make_hipExtent(width, - numH, - std::numeric_limits::max()); - - ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream); - REQUIRE(ret != hipSuccess); - - ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream); - REQUIRE(ret != hipSuccess); - - if ((TestContext::get()).isAmd()) { - ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream); - REQUIRE(ret != hipSuccess); - } else { - WARN("Test is skipped for max depth." - << "Cuda doesn't check the maximum depth of extent field"); - } - } - - SECTION("Device Ptr out bound and extent(0) passed for memset") { - size_t pitch = devPitchedPtr.pitch; - size_t slicePitch = pitch * extent.height; - constexpr auto advanceOffset = 10; - - // Point devptr to end of allocated memory - char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) - + depth * slicePitch; - - // Advance devptr further to go out of boundary - devPtrMod = devPtrMod + advanceOffset; - hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, - numW * sizeof(char), numH); - hipExtent extent0{}; - ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream); - - // api expected to check extent0 and return success before going for ptr. - REQUIRE(ret == hipSuccess); - } - - SECTION("Device Ptr out bound and valid extent passed for memset") { - size_t pitch = devPitchedPtr.pitch; - size_t slicePitch = pitch * extent.height; - constexpr auto advanceOffset = 10; - - // Point devptr to end of allocated memory - char *devPtrMod = (reinterpret_cast(devPitchedPtr.ptr)) - + depth * slicePitch; - - // Advance devptr further to go out of boundary - devPtrMod = devPtrMod + advanceOffset; - hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch, - numW * sizeof(char), numH); - ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream); - - REQUIRE(ret != hipSuccess); - } - - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipFree(devPitchedPtr.ptr)); -} diff --git a/catch/unit/memory/hipMemsetInvalidPtr.cc b/catch/unit/memory/hipMemsetInvalidPtr.cc deleted file mode 100644 index 8c87c66ea8..0000000000 --- a/catch/unit/memory/hipMemsetInvalidPtr.cc +++ /dev/null @@ -1,127 +0,0 @@ -/* - * Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. -*/ - -/** -Testcase Scenarios : - 1) Test hipMemset apis with invalid pointer and invalid 2D pitch. - 2) Test hipMemsetAsync apis with invalid pointer and invalid 2D pitch. -*/ - - -#include - -#define N 50 -#define MEMSETVAL 0x42 - -/** - * Testcase validates hipMemset apis behavior with - * invalid pointer and invalid 2D pitch value. - */ -TEST_CASE("Unit_hipMemset_InvalidPtrTests") { - hipError_t ret; - constexpr int Nbytes = N*sizeof(char); - char *A_d; - - SECTION("hipMemset with null") { - ret = hipMemset(NULL, MEMSETVAL , Nbytes); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemset with hostptr") { - char *A_h; - A_h = reinterpret_cast(malloc(Nbytes)); - - ret = hipMemset(A_h, MEMSETVAL, Nbytes); - REQUIRE(ret != hipSuccess); - - free(A_h); - } - - SECTION("hipMemsetD32 with null") { - ret = hipMemsetD32(NULL, MEMSETVAL , Nbytes); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemsetD16 with null") { - ret = hipMemsetD16(NULL, MEMSETVAL , Nbytes); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemsetD8 with null") { - ret = hipMemsetD8(NULL, MEMSETVAL , Nbytes); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemset2D with null") { - constexpr size_t NUM_H = 256, NUM_W = 256; - size_t pitch_A; - size_t width = NUM_W * sizeof(char); - - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, - width , NUM_H)); - ret = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H); - REQUIRE(ret != hipSuccess); - - hipFree(A_d); - } -} - - -/** - * Testcase validates hipMemsetAsync apis behavior with - * invalid pointer and invalid 2D pitch value. - */ -TEST_CASE("Unit_hipMemsetAsync_InvalidPtrTests") { - hipError_t ret; - constexpr int Nbytes = N*sizeof(char); - char *A_d; - - SECTION("hipMemsetAsync with null") { - ret = hipMemsetAsync(NULL, MEMSETVAL, Nbytes , 0); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemsetD32Async with null") { - ret = hipMemsetD32Async(NULL, MEMSETVAL , Nbytes, 0); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemsetD16Async with null") { - ret = hipMemsetD16Async(NULL, MEMSETVAL , Nbytes, 0); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemsetD8Async with null") { - ret = hipMemsetD8Async(NULL, MEMSETVAL , Nbytes, 0); - REQUIRE(ret != hipSuccess); - } - - SECTION("hipMemset2DAsync with null") { - constexpr size_t NUM_H = 256, NUM_W = 256; - size_t pitch_A; - size_t width = NUM_W * sizeof(char); - - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, - width , NUM_H)); - ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0); - REQUIRE(ret != hipSuccess); - - hipFree(A_d); - } -} diff --git a/catch/unit/memory/hipMemsetNegative.cc b/catch/unit/memory/hipMemsetNegative.cc new file mode 100644 index 0000000000..a06638904e --- /dev/null +++ b/catch/unit/memory/hipMemsetNegative.cc @@ -0,0 +1,229 @@ +/* + * 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include +#include +#include "hip/driver_types.h" + +static constexpr size_t memsetVal{0x42}; +static constexpr hipExtent validExtent{184, 57, 16}; +static constexpr size_t height{validExtent.height}; +static constexpr size_t width{validExtent.width}; +static constexpr hipStream_t nullStream{nullptr}; + +inline void testHipMemsetApis(void* dst, int value, size_t sizeBytes, + hipError_t expectedReturn = hipErrorInvalidValue) { + HIP_CHECK_ERROR(hipMemset(dst, value, sizeBytes), expectedReturn); + HIP_CHECK_ERROR(hipMemsetAsync(dst, value, sizeBytes, nullStream), expectedReturn); + + hipDeviceptr_t devicePtrDst = reinterpret_cast(dst); + HIP_CHECK_ERROR(hipMemsetD32(devicePtrDst, value, sizeBytes), expectedReturn); + HIP_CHECK_ERROR(hipMemsetD32Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn); + HIP_CHECK_ERROR(hipMemsetD16(devicePtrDst, value, sizeBytes), expectedReturn); + HIP_CHECK_ERROR(hipMemsetD16Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn); + HIP_CHECK_ERROR(hipMemsetD8(devicePtrDst, value, sizeBytes), expectedReturn); + HIP_CHECK_ERROR(hipMemsetD8Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn); +} + +inline void testHipMemset2DApis(void* dst, size_t pitch, int value, size_t width, size_t height, + hipError_t expectedReturn = hipErrorInvalidValue) { + HIP_CHECK_ERROR(hipMemset2D(dst, pitch, value, width, height), expectedReturn); + HIP_CHECK_ERROR(hipMemset2DAsync(dst, pitch, value, width, height, nullStream), expectedReturn); +} + + +inline void testHipMemset3DApis(hipPitchedPtr& pitchedDevPtr, int value, const hipExtent& extent, + hipError_t expectedReturn = hipErrorInvalidValue) { + HIP_CHECK_ERROR(hipMemset3D(pitchedDevPtr, value, extent), expectedReturn); + HIP_CHECK_ERROR(hipMemset3DAsync(pitchedDevPtr, value, extent, nullStream), expectedReturn); +} + +TEST_CASE("Unit_hipMemset_Negative_InvalidPtr") { + void* dst; + + SECTION("Uninitialized Dst") {} + SECTION("Nullptr as dst") { dst = nullptr; } + + std::unique_ptr hostPtr; + SECTION("Host Pointer as Dst") { + hostPtr.reset(new char[width]); + dst = hostPtr.get(); + } + + testHipMemsetApis(dst, memsetVal, width); +} + + +TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsSize") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-20"); +#endif + +#if !HT_AMD + void* dst; + constexpr size_t outOfBoundsSize{width + 1}; + HIP_CHECK(hipMalloc(&dst, width)); + + testHipMemsetApis(dst, memsetVal, outOfBoundsSize); + HIP_CHECK(hipFree(dst)); +#endif +} + +TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsPtr") { + void* dst; + HIP_CHECK(hipMalloc(&dst, width)); + void* outOfBoundsPtr{reinterpret_cast(dst) + width + 1}; + testHipMemsetApis(outOfBoundsPtr, memsetVal, width); + HIP_CHECK(hipFree(dst)); +} + +TEST_CASE("Unit_hipMemset2D_Negative_InvalidPtr") { + void* dst; + SECTION("Uninitialized Dst") {} + SECTION("Nullptr as Dst") { dst = nullptr; } + + std::unique_ptr hostPtr; + SECTION("Host Pointer as Dst") { + hostPtr.reset(new char[height * width]); + dst = hostPtr.get(); + } + + void* A_d; + size_t pitch_A; + HIP_CHECK(hipMallocPitch(&A_d, &pitch_A, width, height)); + testHipMemset2DApis(dst, pitch_A, memsetVal, width, height); + hipFree(A_d); +} + +TEST_CASE("Unit_hipMemset2D_Negative_InvalidSizes") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52"); +#endif + + void* dst; + size_t realPitch; + HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height)); + + SECTION("Invalid Pitch") { + size_t invalidPitch = 1; + testHipMemset2DApis(dst, invalidPitch, memsetVal, width, height); + } + + SECTION("Invalid Width") { + size_t invalidWidth = realPitch + 1; + testHipMemset2DApis(dst, realPitch, memsetVal, invalidWidth, height); + } + +#if !HT_AMD /* EXSWCPHIPT-52 */ + SECTION("Invalid height") { + size_t invalidHeight = height + 1; + testHipMemset2DApis(dst, realPitch, memsetVal, width, invalidHeight); + } +#endif + HIP_CHECK(hipFree(dst)); +} + +TEST_CASE("Unit_hipMemset2D_Negative_OutOfBoundsPtr") { + void* dst; + size_t realPitch; + + HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height)); + void* outOfBoundsPtr{reinterpret_cast(dst) + realPitch * height + 1}; + testHipMemset2DApis(outOfBoundsPtr, realPitch, memsetVal, width, height); + HIP_CHECK(hipFree(dst)); +} + + +TEST_CASE("Unit_hipMemset3D_Negative_InvalidPtr") { + hipPitchedPtr pitchedDevPtr; + + SECTION("Uninitialized PitchedDevPtr") {} + SECTION("Zero Initialized PitchedDevPtr") { pitchedDevPtr = {}; } + + testHipMemset3DApis(pitchedDevPtr, memsetVal, validExtent); +} + +TEST_CASE("Unit_hipMemset3D_Negative_ModifiedPtr") { + hipPitchedPtr pitchedDevPtr; + + HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent)); + void* allocatedMemory{pitchedDevPtr.ptr}; + + SECTION("Nullptr Dst") { pitchedDevPtr.ptr = nullptr; } + + std::unique_ptr hostPtr; + SECTION("Host Pointer as Dst") { + hostPtr.reset(new char[validExtent.width * validExtent.height * validExtent.depth]); + pitchedDevPtr.ptr = hostPtr.get(); + } + + SECTION("Invalid Pitch") { pitchedDevPtr.pitch = 1; } + + CAPTURE(pitchedDevPtr.ptr, pitchedDevPtr.pitch, pitchedDevPtr.xsize, pitchedDevPtr.ysize); + testHipMemset3DApis(pitchedDevPtr, memsetVal, validExtent); + HIP_CHECK(hipFree(allocatedMemory)); +} + +TEST_CASE("Unit_hipMemset3D_Negative_InvalidSizes") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52"); +#endif + + hipPitchedPtr pitchedDevPtr; + HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent)); + hipExtent invalidExtent{validExtent}; + + SECTION("Max Width") { invalidExtent.width = std::numeric_limits::max(); } + + SECTION("Max Height") { invalidExtent.height = std::numeric_limits::max(); } + +#if !HT_NVIDIA /* This case hangs on Nvidia */ + SECTION("Max Depth") { invalidExtent.depth = std::numeric_limits::max(); } +#endif + + SECTION("Invalid Width") { invalidExtent.width = pitchedDevPtr.pitch + 1; } + +#if !HT_AMD /* EXSWCPHIPT-52 */ + SECTION("Invalid height") { invalidExtent.height += 1; } + + SECTION("Invalid depth") { invalidExtent.depth += 1; } +#endif + + CAPTURE(invalidExtent.width, invalidExtent.height, invalidExtent.depth); + testHipMemset3DApis(pitchedDevPtr, memsetVal, invalidExtent); + HIP_CHECK(hipFree(pitchedDevPtr.ptr)); +} + +TEST_CASE("Unit_hipMemset3D_Negative_OutOfBounds") { + hipPitchedPtr pitchedDevPtr; + + HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent)); + hipPitchedPtr outOfBoundsPtr{pitchedDevPtr}; + outOfBoundsPtr.ptr = reinterpret_cast(pitchedDevPtr.ptr) + + pitchedDevPtr.pitch * validExtent.height * validExtent.depth + 1; + + SECTION("Extent Equal to 0") { + hipExtent zeroExtent{0, 0, 0}; + testHipMemset3DApis(outOfBoundsPtr, memsetVal, zeroExtent, hipSuccess); + } + + SECTION("Valid Extent") { testHipMemset3DApis(outOfBoundsPtr, memsetVal, validExtent); } + + HIP_CHECK(hipFree(pitchedDevPtr.ptr)); +}