diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index c262620f53..997572ee6d 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -82,6 +82,7 @@ set(TEST_SRC hipMalloc3D.cc hipMalloc3DArray.cc hipArrayCreate.cc + hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc @@ -150,6 +151,7 @@ set(TEST_SRC hipMalloc3D.cc hipMalloc3DArray.cc hipArrayCreate.cc + hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc diff --git a/projects/hip-tests/catch/unit/memory/DriverContext.cc b/projects/hip-tests/catch/unit/memory/DriverContext.cc index 6791650567..dc27a62000 100644 --- a/projects/hip-tests/catch/unit/memory/DriverContext.cc +++ b/projects/hip-tests/catch/unit/memory/DriverContext.cc @@ -24,17 +24,13 @@ THE SOFTWARE. #include DriverContext::DriverContext() { -#if HT_NVIDIA HIP_CHECK(hipInit(0)); HIP_CHECK(hipDeviceGet(&device, 0)); HIP_CHECK(hipDevicePrimaryCtxRetain(&ctx, device)); HIP_CHECK(hipCtxPushCurrent(ctx)); -#endif } DriverContext::~DriverContext() { -#if HT_NVIDIA HIP_CHECK(hipCtxPopCurrent(&ctx)); HIP_CHECK(hipDevicePrimaryCtxRelease(device)); -#endif } diff --git a/projects/hip-tests/catch/unit/memory/DriverContext.hh b/projects/hip-tests/catch/unit/memory/DriverContext.hh index 76afe28f44..5593c512d4 100644 --- a/projects/hip-tests/catch/unit/memory/DriverContext.hh +++ b/projects/hip-tests/catch/unit/memory/DriverContext.hh @@ -26,10 +26,8 @@ THE SOFTWARE. class DriverContext { private: -#if HT_NVIDIA hipCtx_t ctx; hipDevice_t device; -#endif public: DriverContext(); diff --git a/projects/hip-tests/catch/unit/memory/hipArray3DCreate.cc b/projects/hip-tests/catch/unit/memory/hipArray3DCreate.cc new file mode 100644 index 0000000000..47e70e66c5 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipArray3DCreate.cc @@ -0,0 +1,166 @@ +/* +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 "DriverContext.hh" +#include "hipArrayCommon.hh" + +namespace { +void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) { +// hipArray3DGetDescriptor doesn't currently exist (EXSWCPHIPT-87) +#if HT_AMD + std::ignore = array; + std::ignore = expected_desc; +#else + CUDA_ARRAY3D_DESCRIPTOR queried_desc; + cuArray3DGetDescriptor(&queried_desc, array); + + REQUIRE(queried_desc.Width == expected_desc.Width); + REQUIRE(queried_desc.Height == expected_desc.Height); + REQUIRE(queried_desc.Depth == expected_desc.Depth); + REQUIRE(queried_desc.Format == expected_desc.Format); + REQUIRE(queried_desc.NumChannels == expected_desc.NumChannels); + REQUIRE(queried_desc.Flags == expected_desc.Flags); +#endif +} +} // namespace + +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float, + float2, float4) { + using vec_info = vector_info; + DriverContext ctx; + + hiparray array; + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; +#if HT_AMD + desc.Flags = 0; +#else + desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); +#endif + + constexpr size_t size = 64; + + std::vector extents{ + {size, 0, 0}, // 1D array + {size, size, 0}, // 2D array + {size, size, size} // 3D array + }; + + for (auto& extent : extents) { + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + HIP_CHECK(hipArray3DCreate(&array, &desc)); + checkArrayIsExpected(array, desc); + HIP_CHECK(hipArrayDestroy(array)); + } +} + +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, ushort2, + unsigned char, float, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97"); + return; +#endif + + using vec_info = vector_info; + DriverContext ctx; + + hiparray array; + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; +#if HT_AMD + desc.Flags = 0; +#else + desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); + if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); + return; + } +#endif + CAPTURE(desc.Flags); + + const Sizes sizes(desc.Flags); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { + // stored in a vector so some values can be ifdef'd out + std::vector extentsToTest{ + make_hipExtent(sizes.max1D, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0], s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1], 0), // 2D max height + make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0), // 2D max + make_hipExtent(sizes.max3D[0], s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1], s), // 3D max height + make_hipExtent(s, s, sizes.max3D[2]), // 3D max depth + make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]), // 3D max height and depth + make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]), // 3D max width and depth + make_hipExtent(sizes.max3D[0], sizes.max3D[1], s), // 3D max width and height + make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2]) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + auto maxArrayCreateError = hipArray3DCreate(&array, &desc); + // this can try to alloc many GB of memory, so out of memory is acceptable + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + checkArrayIsExpected(array, desc); + HIP_CHECK(hipArrayDestroy(array)); + } + SECTION("Negative") { + std::vector extentsToTest { + make_hipExtent(sizes.max1D + 1, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0] + 1, s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1] + 1, 0), // 2D max height + make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0), // 2D max + make_hipExtent(sizes.max3D[0] + 1, s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1] + 1, s), // 3D max height +#if !HT_NVIDIA // leads to hipSuccess on NVIDIA + make_hipExtent(s, s, sizes.max3D[2] + 1), // 3D max depth +#endif + make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1), // 3D max height and depth + make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1), // 3D max width and depth + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s), // 3D max width and height + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); + } +} diff --git a/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh index c9823806c5..b51022eb56 100644 --- a/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh +++ b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh @@ -131,8 +131,16 @@ struct Sizes { Sizes(unsigned int flag) { int device; HIP_CHECK(hipGetDevice(&device)); + static_assert( + hipArrayDefault == 0, + "hipArrayDefault is assumed to be equivalent to 0 for the following switch statment"); +#if HT_NVIDIA + static_assert(hipArraySurfaceLoadStore == CUDA_ARRAY3D_SURFACE_LDST, + "hipArraySurface is assumed to be equivalent to CUDA_ARRAY3D_SURFACE_LDST for " + "the following switch statment"); +#endif switch (flag) { - case hipArrayDefault: { + case hipArrayDefault: { // 0 hipDeviceProp_t prop; HIP_CHECK(hipGetDeviceProperties(&prop, device)); max1D = prop.maxTexture1D; @@ -140,7 +148,7 @@ struct Sizes { max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]}; return; } - case hipArraySurfaceLoadStore: { + case hipArraySurfaceLoadStore: { // CUDA_ARRAY3D_SURFACE_LDST int value; HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device)); max1D = value; diff --git a/projects/hip-tests/catch/unit/memory/hipMalloc3DArray.cc b/projects/hip-tests/catch/unit/memory/hipMalloc3DArray.cc index 253913345e..be41740242 100644 --- a/projects/hip-tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/projects/hip-tests/catch/unit/memory/hipMalloc3DArray.cc @@ -145,6 +145,7 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { } } +namespace { void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc, const hipExtent& expected_extent, const unsigned int expected_flags) { // hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) @@ -172,6 +173,7 @@ void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected REQUIRE(expected_flags == queried_flags); #endif } +} TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, float2, float4) { @@ -187,19 +189,16 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); } SECTION("2D Array") { extent = make_hipExtent(size, size, 0); - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); } SECTION("3D Array") { extent = make_hipExtent(size, size, size); - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); } + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); checkArrayIsExpected(array, desc, extent, flags); - HIP_CHECK(hipFreeArray(array)); }