AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

hipMemset Negative Testing (#2639)

Browse Source
This commit is contained in:
Fábio
2022-05-25 14:37:20 +01:00
committed by GitHub
parent 205107aa3c
commit b64d0cff29
4 changed files with 231 additions and 367 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/memory/CMakeLists.txt
+2 -4
View File
@@ -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
catch/unit/memory/hipMemset3DNegative.cc
-236
View File
@@ -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 <hip_test_common.hh>
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<std::size_t>::max(),
numH,
depth);
hipExtent extMH = make_hipExtent(width,
std::numeric_limits<std::size_t>::max(),
depth);
hipExtent extMD = make_hipExtent(width,
numH,
std::numeric_limits<std::size_t>::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<char *>(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<char *>(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<std::size_t>::max(),
numH,
depth);
hipExtent extMH = make_hipExtent(width,
std::numeric_limits<std::size_t>::max(),
depth);
hipExtent extMD = make_hipExtent(width,
numH,
std::numeric_limits<std::size_t>::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<char *>(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<char *>(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));
}
catch/unit/memory/hipMemsetInvalidPtr.cc
-127
View File
@@ -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 <hip_test_common.hh>
#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<char *>(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<void**>(&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<void**>(&A_d), &pitch_A,
width , NUM_H));
ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0);
REQUIRE(ret != hipSuccess);
hipFree(A_d);
}
}
catch/unit/memory/hipMemsetNegative.cc
+229
View File
@@ -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 <hip_test_common.hh>
#include <limits>
#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<hipDeviceptr_t>(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<char[]> 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<char*>(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<char[]> 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<char*>(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<char[]> 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<std::size_t>::max(); }
SECTION("Max Height") { invalidExtent.height = std::numeric_limits<std::size_t>::max(); }
#if !HT_NVIDIA /* This case hangs on Nvidia */
SECTION("Max Depth") { invalidExtent.depth = std::numeric_limits<std::size_t>::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<char*>(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));
}