From f79b6c5a5510cba5fd5dcdeb5bc3b6b07454ebe9 Mon Sep 17 00:00:00 2001 From: Finlay Date: Thu, 28 Jul 2022 12:32:59 +0100 Subject: [PATCH] EXSWCPHIPT-136 Texture Gather tests for hipMalloc3DArray and hipArray3DCreate (#2818) --- catch/unit/memory/hipArray3DCreate.cc | 62 +++++++++++++++++++++------ catch/unit/memory/hipMalloc3DArray.cc | 43 +++++++++++++++---- 2 files changed, 82 insertions(+), 23 deletions(-) diff --git a/catch/unit/memory/hipArray3DCreate.cc b/catch/unit/memory/hipArray3DCreate.cc index faa3822b9b..9c11152936 100644 --- a/catch/unit/memory/hipArray3DCreate.cc +++ b/catch/unit/memory/hipArray3DCreate.cc @@ -58,15 +58,17 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, #if HT_AMD desc.Flags = 0; #else - desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); + desc.Flags = GENERATE(0, hipArraySurfaceLoadStore, hipArrayTextureGather); #endif constexpr size_t size = 64; - std::vector extents{ - {size, 0, 0}, // 1D array - {size, size, 0}, // 2D array - {size, size, size} // 3D array + std::vector extents; + extents.reserve(3); + extents.push_back({size, size, 0}); // 2D array + if (desc.Flags != hipArrayTextureGather) { + extents.push_back({size, 0, 0}); // 1D array + extents.push_back({size, size, size}); // 3D array }; for (auto& extent : extents) { @@ -100,8 +102,8 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, us #if HT_AMD desc.Flags = 0; #else - desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); - if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) { + desc.Flags = GENERATE(0, hipArraySurfaceLoadStore); + if (desc.Flags == hipArraySurfaceLoadStore) { HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); return; } @@ -200,7 +202,7 @@ constexpr HIP_ARRAY3D_DESCRIPTOR defaultDescriptor(unsigned int flags, size_t si desc.Depth = size; #if HT_NVIDIA - if (flags == CUDA_ARRAY3D_TEXTURE_GATHER) { + if (flags == hipArrayTextureGather) { desc.Depth = 0; } #endif @@ -247,8 +249,8 @@ TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroHeight") { unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); auto desc = defaultDescriptor(flags, 6); #if HT_NVIDIA - std::array exceptions{CUDA_ARRAY3D_LAYERED, - CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST}; + std::array exceptions{hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore}; #else std::array exceptions{}; #endif @@ -290,10 +292,9 @@ TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFlags") { // FIXME: use the same flags for both tests when the values exist for hip #if HT_NVIDIA - unsigned int flags = - GENERATE(0xDEADBEEF, CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST, - CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_CUBEMAP, - CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_CUBEMAP); + unsigned int flags = GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore, + hipArrayTextureGather | hipArrayCubemap, + hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap); #else unsigned int flags = 0xDEADBEEF; #endif @@ -319,3 +320,36 @@ TEST_CASE("Unit_hipArray3DCreate_Negative_NumericLimit") { testInvalidDescription(desc); } + +// texture gather arrays may only be 2D +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_Negative_Non2DTextureGather", "", char, uint2, int4, + float2, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); + return; +#endif + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + desc.Flags = hipArrayTextureGather; + + constexpr size_t size = 64; + + std::array extents{ + make_hipExtent(size, 0, 0), // 1D array + make_hipExtent(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); + + testInvalidDescription(desc); + } +} diff --git a/catch/unit/memory/hipMalloc3DArray.cc b/catch/unit/memory/hipMalloc3DArray.cc index 1a8ec2f0d0..1b7ee27171 100644 --- a/catch/unit/memory/hipMalloc3DArray.cc +++ b/catch/unit/memory/hipMalloc3DArray.cc @@ -136,18 +136,26 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, #if HT_AMD const unsigned int flags = hipArrayDefault; #else - const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); + const unsigned int flags = + GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); #endif constexpr size_t size = 64; - hipExtent extent; - SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); } - SECTION("2D Array") { extent = make_hipExtent(size, size, 0); } - SECTION("3D Array") { extent = make_hipExtent(size, size, size); } + std::vector extents; + extents.reserve(3); + extents.push_back({size, size, 0}); // 2D array + if (flags != hipArrayTextureGather) { + extents.push_back({size, 0, 0}); // 1D array + extents.push_back({size, size, size}); // 3D array + }; - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); - checkArrayIsExpected(array, desc, extent, flags); - HIP_CHECK(hipFreeArray(array)); + for (const auto extent : extents) { + CAPTURE(flags, extent.width, extent.height, extent.depth); + + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + checkArrayIsExpected(array, desc, extent, flags); + HIP_CHECK(hipFreeArray(array)); + } } TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, ushort2, @@ -316,7 +324,7 @@ TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") { HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), hipErrorInvalidValue); } -void testInvalidDescription(hipChannelFormatDesc desc){ +void testInvalidDescription(hipChannelFormatDesc desc) { constexpr size_t s = 6; // 6 to keep cubemap happy hipArray_t array; @@ -420,3 +428,20 @@ TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") { HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag), hipErrorInvalidValue); } + +// texture gather arrays are only allowed to be 2D +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_Negative_Non2DTextureGather", "", char, uchar2, short4, + float2, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); + return; +#endif + hipArray_t array; + const auto desc = hipCreateChannelDesc(); + + constexpr unsigned int flags = hipArrayTextureGather; + constexpr size_t size = 64; + const hipExtent extent = GENERATE(make_hipExtent(size, 0, 0), make_hipExtent(size, size, size)); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flags), hipErrorInvalidValue); +}