From a2ae1bcb87c9aca4254794a3a55684b4729fc58c Mon Sep 17 00:00:00 2001 From: milos-mozetic <118800401+milos-mozetic@users.noreply.github.com> Date: Mon, 14 Aug 2023 17:15:35 +0200 Subject: [PATCH] EXSWHTEC-224 - Test cases ID clean up and documentation for Texture Management (#92) - Update hipBindTextureToMipmappedArray.cc - Fix build error for undeclared identifier 'ret' - Disable tests which fail in external CI [ROCm/hip-tests commit: 7c7884a2d9e88d43f1137a3b48990c111cd5e965] --- .../config/config_amd_linux_common.json | 3 + .../config/config_amd_windows_common.json | 24 +- .../config/config_nvidia_linux_common.json | 8 +- .../catch/include/hip_test_defgroups.hh | 7 + .../catch/unit/texture/CMakeLists.txt | 3 + .../texture/hipBindTextureToMipmappedArray.cc | 208 ++++++++++ .../hipCreateTextureObject_ArgValidation.cc | 34 +- .../texture/hipCreateTextureObject_Array.cc | 35 +- .../texture/hipCreateTextureObject_Linear.cc | 36 +- .../texture/hipCreateTextureObject_Pitch2D.cc | 46 ++- .../catch/unit/texture/hipGetChanDesc.cc | 74 +++- .../catch/unit/texture/hipTexObjPitch.cc | 22 +- .../catch/unit/texture/hipTexObjectTests.cc | 359 +++++++++++++++++ .../unit/texture/hipTextureMipmapObj2D.cc | 19 + .../unit/texture/hipTextureObj1DCheckModes.cc | 18 + .../catch/unit/texture/hipTextureObj2D.cc | 18 + .../unit/texture/hipTextureObj2DCheckModes.cc | 18 + .../unit/texture/hipTextureObj3DCheckModes.cc | 126 +++--- .../unit/texture/hipTextureObjectTests.cc | 369 ++++++++++++++++++ 19 files changed, 1321 insertions(+), 106 deletions(-) create mode 100644 projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTexObjectTests.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureObjectTests.cc 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 940eed28fc..2e7d092244 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 @@ -111,6 +111,9 @@ "Unit_deviceAllocation_Malloc_ComplexDataType", "Unit_deviceAllocation_New_ComplexDataType", "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1", + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===", + "Unit_hipGetChannelDesc_Negative_Parameters", + "Unit_hipGraphAddChildGraphNode_CmplxNstGrph_UpdKerFun_Clone", "=== Below tests fail in stress test on 24/07/23 ===", "Unit_hipStreamCreateWithPriority_ValidateWithEvents", "Unit_hipEventIpc", 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 141d5a35e0..405556ecc1 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 @@ -212,15 +212,19 @@ "Unit_hipMemcpyDtoHAsync_Negative_Parameters", "Unit_hipMemcpyHtoDAsync_Negative_Parameters", "Unit_hipMemcpyDtoDAsync_Negative_Parameters", - "=== Below hiprtc tests are disabled temporarily, will be renabled once patches for SWDEV-395996 are merged ===", - "Unit_hiprtc_saxpy.Unit_hiprtc_saxpy", - "Unit_hiprtc_warpsize.Unit_hiprtc_warpsize", - "Unit_hiprtc_functional.Unit_hiprtc_functional", - "Unit_hipStreamCaptureRtc.Unit_hipStreamCaptureRtc", - "Unit_hiprtc_cpp17.Unit_hiprtc_cpp17", - "Unit_hiprtc_namehandling.Unit_hiprtc_namehandling", - "Unit_hiprtc_getloweredname.Unit_hiprtc_getloweredname", - "Unit_hiprtc_test_hip_bfloat16.Unit_hiprtc_test_hip_bfloat16", - "Unit_RTC_LinkerAPI.Unit_RTC_LinkerAPI" + "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===", + "Unit_hipGetChannelDesc_Negative_Parameters", + "Unit_hipTextureMipmapRef2D_Positive_Check", + "Unit_hipTextureMipmapRef2D_Negative_Parameters", + "=== Below hiprtc tests are disabled temporarily, will be renabled once patches for SWDEV-395996 are merged ===", + "Unit_hiprtc_saxpy.Unit_hiprtc_saxpy", + "Unit_hiprtc_warpsize.Unit_hiprtc_warpsize", + "Unit_hiprtc_functional.Unit_hiprtc_functional", + "Unit_hipStreamCaptureRtc.Unit_hipStreamCaptureRtc", + "Unit_hiprtc_cpp17.Unit_hiprtc_cpp17", + "Unit_hiprtc_namehandling.Unit_hiprtc_namehandling", + "Unit_hiprtc_getloweredname.Unit_hiprtc_getloweredname", + "Unit_hiprtc_test_hip_bfloat16.Unit_hiprtc_test_hip_bfloat16", + "Unit_RTC_LinkerAPI.Unit_RTC_LinkerAPI" ] } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json index baa25636a0..691386f1a1 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux_common.json @@ -20,6 +20,12 @@ "Unit_ChannelDescriptor_Positive_Basic_4D - ulong4", "Unit_ChannelDescriptor_Positive_Basic_4D - long4", "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/38 ===", - "Unit_hipFreeAsync_negative" + "Unit_hipFreeAsync_negative", + "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===", + "Unit_hipGetTexObjectResourceDesc_positive", + "Unit_hipGetTexObjectResourceDesc_Negative_Parameters", + "Unit_hipGetTexObjectTextureDesc_positive", + "Unit_hipGetTexObjectTextureDesc_Negative_Parameters", + "Unit_hipTexObjectDestroy_positive" ] } diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 92857a12bc..119b095dcc 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -94,6 +94,13 @@ THE SOFTWARE. */ /** + * @defgroup TextureTest Texture Management + * @{ + * This section describes tests for the texture management functions of HIP runtime API. + * @} + */ + + /** * @defgroup ShflTest warp shuffle function Management * @{ * This section describes the warp shuffle types & functions of HIP runtime API. diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index 93a188683c..9c7bbfa4a6 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -42,6 +42,9 @@ set(TEST_SRC hipTextureObj3DCheckModes.cc hipTextureObj1DCheckSRGBModes.cc hipTextureObj2DCheckSRGBModes.cc + hipTexObjectTests.cc + hipTextureObjectTests.cc + hipBindTextureToMipmappedArray.cc ) hip_add_exe_to_target(NAME TextureTest diff --git a/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc b/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc new file mode 100644 index 0000000000..670e4f75eb --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc @@ -0,0 +1,208 @@ +/* +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 + +/** + * @addtogroup hipBindTextureToMipmappedArray hipBindTextureToMipmappedArray + * @{ + * @ingroup TextureTest + * `hipBindTextureToMipmappedArray(const textureReference* tex, + * hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc)` - + * Binds a mipmapped array to a texture. + */ + +texture texRef; + +// MipMap is currently supported only on windows +#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) +__global__ void tex2DKernel(float* outputData, int width, float level) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[y * width + x] = tex2DLod(texRef, x, y, level); +} + +static void runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) { + INFO("Width: " << width << "Height: " << height << "mip: " << mipmap_level); + + // Create new width & height to be tested + unsigned int orig_width = width; + unsigned int orig_height = height; + unsigned int i, j; + width /= pow(2, mipmap_level); + height /= pow(2, mipmap_level); + unsigned int size = width * height * sizeof(float); + + float* hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (i = 0; i < height; i++) { + for (j = 0; j < width; j++) { + hData[i * width + j] = i * width + j; + } + } + + // Allocate memory for Mipmapped array and set data to mipmap_level + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + + hipMipmappedArray* mip_array_ptr; + HIP_CHECK(hipMallocMipmappedArray(&mip_array_ptr, &channelDesc, + make_hipExtent(orig_width, orig_height, 0), 2 * mipmap_level, + hipArrayDefault)); + + hipArray* hipArray = nullptr; + HIP_CHECK(hipGetMipmappedArrayLevel(&hipArray, mip_array_ptr, mipmap_level)); + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), + height, hipMemcpyHostToDevice)); + + // Set texture parameters + texRef.addressMode[0] = hipAddressModeWrap; + texRef.addressMode[1] = hipAddressModeWrap; + texRef.filterMode = hipFilterModePoint; + texRef.normalized = 0; + + // Bind the array to the texture + HIP_CHECK(hipBindTextureToMipmappedArray(&texRef, mip_array_ptr, &channelDesc)); + + // Allocate device memory for result + float* dData = nullptr; + HIP_CHECK(hipMalloc(&dData, size)); + REQUIRE(dData != nullptr); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, mipmap_level); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + // Allocate memory on host and copy result from device to host + float* hOutputData = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + + for (i = 0; i < height; i++) { + for (j = 0; j < width; j++) { + if (hData[i * width + j] != hOutputData[i * width + j]) { + INFO("Difference found at [ " << i << j << " ]: " << hData[i * width + j] + << hOutputData[i * width + j]); + REQUIRE(false); + } + } + } + HIP_CHECK(hipUnbindTexture(texRef)); + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipFreeMipmappedArray(mip_array_ptr)); + free(hData); +} +#endif + +/** + * Test Description + * ------------------------ + * - Maps texture to the mipmapped array for different mipmapped array + * sizes and number of levels. + * Test source + * ------------------------ + * - unit/texture/hipBindTextureToMipmappedArray.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Host specific (WINDOWS) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") { + CHECK_IMAGE_SUPPORT + // Height Width Vector + std::vector hw_vec = {2048, 1024, 512, 256, 64}; + std::vector mip_vec = {8, 4, 2, 1}; +#ifdef _WIN32 + for (auto& hw : hw_vec) { + for (auto& mip : mip_vec) { + if ((hw / static_cast(pow(2, (mip * 2)))) > 0) { + runMipMapTest(hw, hw, mip); + } + } + } +#else + SUCCEED("Mipmaps are Supported only on windows, skipping the test."); +#endif +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When texture reference is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When mipmapped array handle is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When channel descriptor is `nullptr` + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/texture/hipBindTextureToMipmappedArray.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Host specific (WINDOWS) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipTextureMipmapRef2D_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + +#ifdef _WIN32 + unsigned int width = 64; + unsigned int height = 64; + unsigned int mipmap_level = 1; + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + + hipMipmappedArray* mip_array_ptr; + HIP_CHECK(hipMallocMipmappedArray(&mip_array_ptr, &channelDesc, make_hipExtent(width, height, 0), + mipmap_level, hipArrayDefault)); + + texRef.addressMode[0] = hipAddressModeWrap; + texRef.addressMode[1] = hipAddressModeWrap; + texRef.filterMode = hipFilterModePoint; + texRef.normalized = 0; + hipError_t ret; + + SECTION("textureReference is nullptr") { + ret = hipBindTextureToMipmappedArray(nullptr, mip_array_ptr, &channelDesc); + REQUIRE(ret != hipSuccess); + } + + SECTION("MipmappedArray is nullptr") { + hipError_t ret = hipBindTextureToMipmappedArray(&texRef, nullptr, &channelDesc); + REQUIRE(ret != hipSuccess); + } + + SECTION("Channel descriptor is nullptr") { + ret = hipBindTextureToMipmappedArray(&texRef, mip_array_ptr, nullptr); + REQUIRE(ret != hipSuccess); + } + + HIP_CHECK(hipFreeMipmappedArray(mip_array_ptr)); +#else + SUCCEED("Mipmaps are Supported only on windows, skipping the test."); +#endif +} diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc index 0ce2d51f2d..d0d3d965f7 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc @@ -19,10 +19,34 @@ THE SOFTWARE. #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + #define N 512 -/* - * Validate argument list of texture object api. +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments for [hipCreateTextureObject](@ref hipCreateTextureObject): + * -# When output pointer to the texture object is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When resource descriptor is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When texture descriptor is `nullptr` + * - Expected output: do not return `hipSuccess` + * - Validates handling of invalid arguments for [hipDestroyTextureObject](@ref hipDestroyTextureObject): + * -# When texture object handle is `nullptr` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/texture/hipCreateTextureObject_ArgValidation.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { CHECK_IMAGE_SUPPORT @@ -34,7 +58,7 @@ TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { hipTextureDesc texDesc; hipTextureObject_t texObj; - /** Initialization */ + // Initialization HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float))); // Populate resource descriptor memset(&resDesc, 0, sizeof(resDesc)); @@ -49,7 +73,7 @@ TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { texDesc.readMode = hipReadModeElementType; - /** Sections */ + // Sections SECTION("TextureObject as nullptr") { ret = hipCreateTextureObject(nullptr, &resDesc, &texDesc, nullptr); REQUIRE(ret != hipSuccess); @@ -78,6 +102,6 @@ TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { REQUIRE(ret == hipSuccess); } - /** De-Initialization */ + // De-Initialization HIP_CHECK(hipFree(texBuf)); } diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc index 366547d655..5395502f84 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc @@ -19,8 +19,24 @@ THE SOFTWARE. #include -/* - * Validates Array Resource texture object with negative/functional tests. +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of a regular `nullptr` array + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/texture/hipCreateTextureObject_Array.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { CHECK_IMAGE_SUPPORT @@ -44,9 +60,18 @@ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { REQUIRE(ret != hipSuccess); } -/* - * Validates MipMappedArray Resource texture object - * with negative/functional tests. +/** + * Test Description + * ------------------------ + * - Validates handling of a regular `nullptr` mipmapped array + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/texture/hipCreateTextureObject_Array.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") { CHECK_IMAGE_SUPPORT diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc index fb12586d37..286c12bfc6 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc @@ -19,11 +19,37 @@ THE SOFTWARE. #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + #define UNALIGN_OFFSET 1 #define N 512 -/* - * Validates Linear Resource texture object with negative/functional tests. +/** + * Test Description + * ------------------------ + * - Validates handling of invalid linear resource: + * -# When device pointer is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When size in bytes is 0 + * - Expected output: do not return `hipSuccess` + * -# When size in bytes is `size_t` maximum + * - Expected output: do not return `hipSuccess` + * -# When a valid resource view descriptor is provided + * - Platform specific (AMD) + * - Expected output: do not return `hipSuccess` + * -# When device pointer is not aligned appropriately + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/texture/hipCreateTextureObject_Linear.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipCreateTextureObject_LinearResource") { CHECK_IMAGE_SUPPORT @@ -37,14 +63,14 @@ TEST_CASE("Unit_hipCreateTextureObject_LinearResource") { hipTextureObject_t texObj; hipDeviceProp_t devProp; - /** Initialization */ + // Initialization HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float))); HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); memset(&resDesc, 0, sizeof(resDesc)); memset(&texDesc, 0, sizeof(texDesc)); resDesc.resType = hipResourceTypeLinear; - /** Sections */ + // Sections SECTION("hipResourceTypeLinear and devPtr(nullptr)") { // Populate resource descriptor resDesc.res.linear.devPtr = nullptr; @@ -130,6 +156,6 @@ TEST_CASE("Unit_hipCreateTextureObject_LinearResource") { } } - /** De-Initialization */ + // De-Initialization HIP_CHECK(hipFree(texBuf)); } diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc index fff56d7dd7..f81d03d021 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc @@ -19,14 +19,48 @@ THE SOFTWARE. #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + #define UNALIGN_OFFSET 1 #define SIZE_H 20 #define SIZE_W 30 #define N 512 - -/* - * Validates Pitch2D Resource texture object with negative and functional tests +/** + * Test Description + * ------------------------ + * - Validates handling of invalid 2D pitch resource: + * -# When device pointer is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When device pointer is not aligned appropriately + * - Expected output: do not return `hipSuccess` + * -# When pitch is not aligned appropriately + * - Expected output: do not return `hipSuccess` + * -# When height is 0 + * - Expected output: do not return `hipSuccess` + * -# When height is 0 and device pointer is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When height is `size_t` maximum + * - Expected output: do not return `hipSuccess` + * -# When width is 0 + * - Expected output: do not return `hipSuccess` + * -# When width is 0 and device pointer is `nullptr` + * - Expected output: do not return `hipSuccess` + * -# When width is `size_t` maximum + * - Expected output: do not return `hipSuccess` + * -# When pitch is `size_t` maximum + * - Expected output: do not return `hipSuccess` + * Test source + * ------------------------ + * - unit/texture/hipCreateTextureObject_Pitch2D.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { CHECK_IMAGE_SUPPORT @@ -39,7 +73,7 @@ TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { size_t devPitchA; float *devPtrA; - /** Initialization */ + // Initialization HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtrA), &devPitchA, SIZE_W*sizeof(float), SIZE_H)); HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); @@ -47,7 +81,7 @@ TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { memset(&texDesc, 0, sizeof(texDesc)); resDesc.resType = hipResourceTypePitch2D; - /** Sections */ + // Sections SECTION("hipResourceTypePitch2D and devPtr(nullptr)") { // Populate resource descriptor resDesc.res.pitch2D.devPtr = nullptr; @@ -212,7 +246,7 @@ TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { REQUIRE(ret != hipSuccess); } - /** De-Initialization */ + // De-Initialization HIP_CHECK(hipFree(devPtrA)); } diff --git a/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc b/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc index b7585059c8..2f781254a7 100644 --- a/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc +++ b/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc @@ -19,27 +19,79 @@ THE SOFTWARE. #include +/** + * @addtogroup hipGetChannelDesc hipGetChannelDesc + * @{ + * @ingroup TextureTest + * `hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array)` - + * Gets the channel descriptor in an array. + */ + #define R 8 // rows, height #define C 8 // columns, width - +/** + * Test Description + * ------------------------ + * - Creates a regular channel description. + * - Creates array using previously created description. + * - Checks that valid description is returned. + * Test source + * ------------------------ + * - unit/texture/hipGetChanDesc.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGetChannelDesc_CreateAndGet") { - CHECK_IMAGE_SUPPORT + CHECK_IMAGE_SUPPORT; hipChannelFormatDesc chan_test, chan_desc; - hipArray *hipArray; + hipArray* hip_array; chan_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned); - HIP_CHECK(hipMallocArray(&hipArray, &chan_desc, C, R, 0)); - HIP_CHECK(hipGetChannelDesc(&chan_test, hipArray)); + HIP_CHECK(hipMallocArray(&hip_array, &chan_desc, C, R, 0)); + HIP_CHECK(hipGetChannelDesc(&chan_test, hip_array)); - if ((chan_test.x != 32) || (chan_test.y != 0) - || (chan_test.z != 0) || (chan_test.f != 0)) { - INFO("Mismatch observed : " << chan_test.x << chan_test.y - << chan_test.z << chan_test.f); + if ((chan_test.x != 32) || (chan_test.y != 0) || (chan_test.z != 0) || (chan_test.f != 0)) { + INFO("Mismatch observed : " << chan_test.x << chan_test.y << chan_test.z << chan_test.f); REQUIRE(false); } - - HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipFreeArray(hip_array)); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the description is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When array handle is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipGetChanDesc.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetChannelDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + + hipChannelFormatDesc chan_test, chan_desc; + hipArray* hip_array; + + chan_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned); + HIP_CHECK(hipMallocArray(&hip_array, &chan_desc, C, R, 0)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipGetChannelDesc(nullptr, hip_array), hipErrorInvalidValue); + } + + SECTION("array is nullptr") { + HIP_CHECK_ERROR(hipGetChannelDesc(&chan_test, nullptr), hipErrorInvalidHandle); + } + + HIP_CHECK(hipFreeArray(hip_array)); } diff --git a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc index 364bc79a21..24f9c21aeb 100644 --- a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc +++ b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc @@ -20,6 +20,15 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + * `hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, + * const hipTextureDesc* pTexDesc, const struct hipResourceViewDesc* pResViewDesc)` - + * Creates a texture object. + */ + #define SIZE_H 20 #define SIZE_W 179 @@ -35,7 +44,18 @@ static __global__ void texture2dCopyKernel(hipTextureObject_t texObj, #endif } - +/** + * Test Description + * ------------------------ + * - Basic test where resource type is 2D pitch. + * Test source + * ------------------------ + * - unit/texture/hipTexObjPitch.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int, unsigned char, int16_t, char, unsigned int) { CHECK_IMAGE_SUPPORT diff --git a/projects/hip-tests/catch/unit/texture/hipTexObjectTests.cc b/projects/hip-tests/catch/unit/texture/hipTexObjectTests.cc new file mode 100644 index 0000000000..f7285f93f8 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTexObjectTests.cc @@ -0,0 +1,359 @@ +/* +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 + + +class TexObjectTestWrapper { + private: + float* host_data_; + bool ommit_destroy_; + + public: + hipTextureObject_t texture_object = 0; + HIP_RESOURCE_DESC res_desc; + HIP_TEXTURE_DESC tex_desc; + HIP_RESOURCE_VIEW_DESC res_view_desc; + HIP_ARRAY_DESCRIPTOR array_desc; + hiparray array_member; + size_t size; /* size in bytes*/ + int width; /* width in elements */ + + TexObjectTestWrapper(bool useResourceViewDescriptor, bool ommitDestroy = false) + : ommit_destroy_(ommitDestroy), width(128) { + int i; + size = width * sizeof(float); + + host_data_ = (float*)malloc(size); + memset(host_data_, 0, size); + + for (i = 0; i < width; i++) { + host_data_[i] = i; + } + + memset(&array_desc, 0, sizeof(array_desc)); + array_desc.Format = HIP_AD_FORMAT_FLOAT; + array_desc.NumChannels = 1; + array_desc.Width = width; + array_desc.Height = 0; + + HIP_CHECK(hipArrayCreate(&array_member, &array_desc)); + HIP_CHECK(hipMemcpyHtoA(reinterpret_cast(array_member), 0, host_data_, size)); + + memset(&res_desc, 0, sizeof(res_desc)); + res_desc.resType = HIP_RESOURCE_TYPE_ARRAY; + res_desc.res.array.hArray = array_member; + res_desc.flags = 0; + + memset(&tex_desc, 0, sizeof(tex_desc)); + tex_desc.filterMode = HIP_TR_FILTER_MODE_POINT; + tex_desc.flags = 0; + + memset(&res_view_desc, 0, sizeof(res_view_desc)); + + + if (useResourceViewDescriptor) { +#if HT_AMD + res_view_desc.format = HIP_RES_VIEW_FORMAT_FLOAT_1X32; + res_view_desc.width = size; +#else + /* Resource View Descriptors are not supported on NVIDIA currently */ + useResourceViewDescriptor = false; +#endif + } + + + HIP_CHECK(hipTexObjectCreate(&texture_object, &res_desc, &tex_desc, + useResourceViewDescriptor ? &res_view_desc : nullptr)); + } + + ~TexObjectTestWrapper() { + if (!ommit_destroy_) { + HIP_CHECK(hipTexObjectDestroy(texture_object)); + } + HIP_CHECK(hipArrayDestroy(array_member)); + free(host_data_); + } +}; + +/** + * @addtogroup hipTexObjectGetResourceDesc hipTexObjectGetResourceDesc + * @{ + * @ingroup TextureTest + * `hipTexObjectGetResourceDesc(HIP_RESOURCE_DESC* pResDesc, + * hipTextureObject_t texObject)` - + * Gets resource descriptor of a texture object. + */ + +/** + * Test Description + * ------------------------ + * - Creates new texture object and an empty resource descriptor. + * - Gets resource descriptor from the texture. + * - Compares it to the empty created one. + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetTexObjectResourceDesc_positive") { + CHECK_IMAGE_SUPPORT; + + TexObjectTestWrapper tex_obj_wrapper(false); + + HIP_RESOURCE_DESC check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + HIP_CHECK(hipTexObjectGetResourceDesc(&check_desc, tex_obj_wrapper.texture_object)); + + REQUIRE(check_desc.resType == tex_obj_wrapper.res_desc.resType); + REQUIRE(check_desc.res.array.hArray == tex_obj_wrapper.res_desc.res.array.hArray); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the resource descriptor is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When the texture is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetTexObjectResourceDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + + TexObjectTestWrapper tex_obj_wrapper(false); + + HIP_RESOURCE_DESC check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipTexObjectGetResourceDesc(nullptr, tex_obj_wrapper.texture_object), + hipErrorInvalidValue); + } + + SECTION("texture is invalid") { + HIP_CHECK_ERROR(hipTexObjectGetResourceDesc(&check_desc, static_cast(0)), + hipErrorInvalidValue); + } +} + +/** + * End doxygen group hipTexObjectGetResourceDesc. + * @} + */ + +/** + * @addtogroup hipTexObjectGetResourceViewDesc hipTexObjectGetResourceViewDesc + * @{ + * @ingroup TextureTest + * `hipTexObjectGetResourceViewDesc(HIP_RESOURCE_VIEW_DESC* pResViewDesc, + * hipTextureObject_t texObject)` - + * Gets resource view descriptor of a texture object. + */ + +/** + * Test Description + * ------------------------ + * - Creates new texture object and an empty resource view descriptor. + * - Gets resource view descriptor from the texture. + * - Compares it to the empty created one. + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipGetTexObjectResourceViewDesc_positive") { + CHECK_IMAGE_SUPPORT; + + TexObjectTestWrapper tex_obj_wrapper(true); + + HIP_RESOURCE_VIEW_DESC check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + HIP_CHECK(hipTexObjectGetResourceViewDesc(&check_desc, tex_obj_wrapper.texture_object)); + + REQUIRE(check_desc.format == tex_obj_wrapper.res_view_desc.format); + REQUIRE(check_desc.width == tex_obj_wrapper.res_view_desc.width); +} +#endif + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the resource view descriptor is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When the texture is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipGetTexObjectResourceViewDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + TexObjectTestWrapper tex_obj_wrapper(true); + + HIP_RESOURCE_VIEW_DESC check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipTexObjectGetResourceViewDesc(nullptr, tex_obj_wrapper.texture_object), + hipErrorInvalidValue); + } + + SECTION("texture is invalid") { + HIP_CHECK_ERROR( + hipTexObjectGetResourceViewDesc(&check_desc, static_cast(0)), + hipErrorInvalidValue); + } +} +#endif + +/** + * End doxygen group hipTexObjectGetResourceViewDesc. + * @} + */ + +/** + * @addtogroup hipTexObjectGetTextureDesc hipTexObjectGetTextureDesc + * @{ + * @ingroup TextureTest + * `hipTexObjectGetTextureDesc(HIP_TEXTURE_DESC* pTexDesc, + * hipTextureObject_t texObject)` - + * Gets texture descriptor of a texture object. + */ + +/** + * Test Description + * ------------------------ + * - Creates new texture object and an empty texture descriptor. + * - Gets texture descriptor from the texture. + * - Compares it to the empty created one. + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetTexObjectTextureDesc_positive") { + CHECK_IMAGE_SUPPORT; + + TexObjectTestWrapper tex_obj_wrapper(false); + + HIP_TEXTURE_DESC check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + HIP_CHECK(hipTexObjectGetTextureDesc(&check_desc, tex_obj_wrapper.texture_object)); + + REQUIRE(check_desc.filterMode == tex_obj_wrapper.tex_desc.filterMode); + REQUIRE(check_desc.flags == tex_obj_wrapper.tex_desc.flags); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the texture descriptor is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When the texture is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetTexObjectTextureDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + + TexObjectTestWrapper tex_obj_wrapper(false); + + HIP_TEXTURE_DESC check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipTexObjectGetTextureDesc(nullptr, tex_obj_wrapper.texture_object), + hipErrorInvalidValue); + } + + SECTION("texture is invalid") { + HIP_CHECK_ERROR(hipTexObjectGetTextureDesc(&check_desc, static_cast(0)), + hipErrorInvalidValue); + } +} + +/** + * End doxygen group hipTexObjectGetTextureDesc. + * @} + */ + +/** + * @addtogroup hipTexObjectDestroy hipTexObjectDestroy + * @{ + * @ingroup TextureTest + * `hipTexObjectDestroy(hipTextureObject_t texObject)` - + * Destroys a texture object. + */ + +/** + * Test Description + * ------------------------ + * - Successfully destroys regular texture object. + * Test source + * ------------------------ + * - unit/texture/hipTexObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipTexObjectDestroy_positive") { + CHECK_IMAGE_SUPPORT; + + TexObjectTestWrapper tex_obj_wrapper(false, true); + REQUIRE(hipTexObjectDestroy(tex_obj_wrapper.texture_object) == hipSuccess); +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc index 2b70c13c9f..72b6d196a0 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc @@ -19,6 +19,12 @@ THE SOFTWARE. #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + // Height Width Vector std::vector hw_vector = {2048, 1024, 512, 256, 64}; std::vector mip_vector = {8, 4, 2, 1}; @@ -121,6 +127,19 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int } #endif +/** + * Test Description + * ------------------------ + * - Basic test where resource type is a mipmapped array. + * Test source + * ------------------------ + * - unit/texture/hipTextureMipmapObj2D.cc + * Test requirements + * ------------------------ + * - Host specific (WINDOWS) + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipTextureMipmapObj2D_Check") { CHECK_IMAGE_SUPPORT diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc index 390ef5abc0..a3e8e10f5e 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc @@ -21,6 +21,12 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + template __global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject, int width, float offsetX) { @@ -100,6 +106,18 @@ static void runTest(const int width, const float offsetX) { REQUIRE(result); } +/** + * Test Description + * ------------------------ + * - Uses different addressing and filtering modes for 1D array. + * Test source + * ------------------------ + * - unit/texture/hipTextureObj1DCheckModes.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipTextureObj1DCheckModes") { CHECK_IMAGE_SUPPORT diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc index d545962165..ee4e895b32 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc @@ -19,6 +19,12 @@ THE SOFTWARE. #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT @@ -28,6 +34,18 @@ __global__ void tex2DKernel(float* outputData, #endif } +/** + * Test Description + * ------------------------ + * - Basic test where the resource type is a 2D array. + * Test source + * ------------------------ + * - unit/texture/hipTextureObj2D.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipTextureObj2D_Check") { CHECK_IMAGE_SUPPORT diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc index 84f1081a2a..0e95023dc0 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc @@ -21,6 +21,12 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + template __global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject, int width, int height, float offsetX, @@ -112,6 +118,18 @@ line1: REQUIRE(result); } +/** + * Test Description + * ------------------------ + * - Uses different addressing and filtering modes for 2D arrays. + * Test source + * ------------------------ + * - unit/texture/hipTextureObj2DCheckModes.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipTextureObj2DCheckModes") { CHECK_IMAGE_SUPPORT diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc index 538237b4ad..708caf1201 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc @@ -22,32 +22,38 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipCreateTextureObject hipCreateTextureObject + * @{ + * @ingroup TextureTest + */ + bool LinearFilter3D = false; -template -__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject, - int width, int height, int depth, float offsetX, - float offsetY, float offsetZ) { +template +__global__ void tex3DKernel(float* outputData, hipTextureObject_t textureObject, int width, + int height, int depth, float offsetX, float offsetY, float offsetZ) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; - outputData[z * width * height + y * width + x] = tex3D(textureObject, - normalizedCoords ? (x + offsetX) / width : x + offsetX, - normalizedCoords ? (y + offsetY) / height : y + offsetY, - normalizedCoords ? (z + offsetZ) / depth : z + offsetZ); + outputData[z * width * height + y * width + x] = + tex3D(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX, + normalizedCoords ? (y + offsetY) / height : y + offsetY, + normalizedCoords ? (z + offsetZ) / depth : z + offsetZ); #endif } -template -static void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY, - const float offsetZ) { - //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, offsetX=%f, offsetY=%f, offsetZ=%f)\n", - // __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, - // depth, offsetX, offsetY, offsetZ); +template +static void runTest(const int width, const int height, const int depth, const float offsetX, + const float offsetY, const float offsetZ) { + // printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, + // offsetX=%f, offsetY=%f, offsetZ=%f)\n", + // __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, + // depth, offsetX, offsetY, offsetZ); bool result = true; unsigned int size = width * height * depth * sizeof(float); - float *hData = (float*) malloc(size); + float* hData = (float*)malloc(size); memset(hData, 0, size); for (int i = 0; i < depth; i++) { @@ -61,13 +67,14 @@ static void runTest(const int width, const int height, const int depth, const fl // Allocate array and copy image data hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); - hipArray *arr; + hipArray* arr; - HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + HIP_CHECK( + hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); hipMemcpy3DParms myparms; memset(&myparms, 0, sizeof(myparms)); - myparms.srcPos = make_hipPos(0,0,0); - myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(float), width, height); myparms.dstArray = arr; myparms.extent = make_hipExtent(width, height, depth); @@ -105,20 +112,20 @@ static void runTest(const int width, const int height, const int depth, const fl return; } - float *dData = nullptr; - HIP_CHECK(hipMalloc((void**) &dData, size)); + float* dData = nullptr; + HIP_CHECK(hipMalloc((void**)&dData, size)); HIP_CHECK(hipMemset(dData, 0, size)); - dim3 dimBlock(8, 8, 8); // 512 threads - dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, + dim3 dimBlock(8, 8, 8); // 512 threads + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y, (depth + dimBlock.z - 1) / dimBlock.z); - hipLaunchKernelGGL(tex3DKernel, dimGrid, dimBlock, 0, 0, dData, - textureObject, width, height, depth, offsetX, offsetY, offsetZ); - HIP_CHECK(hipGetLastError()); + hipLaunchKernelGGL(tex3DKernel, dimGrid, dimBlock, 0, 0, dData, textureObject, + width, height, depth, offsetX, offsetY, offsetZ); + HIP_CHECK(hipGetLastError()); HIP_CHECK(hipDeviceSynchronize()); - float *hOutputData = (float*) malloc(size); + float* hOutputData = (float*)malloc(size); memset(hOutputData, 0, size); HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); @@ -130,8 +137,8 @@ static void runTest(const int width, const int height, const int depth, const fl width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData); if (!hipTextureSamplingVerify(hOutputData[index], expectedValue)) { - INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i << "):" << - hOutputData[index] << " expected:" << expectedValue); + INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i + << "):" << hOutputData[index] << " expected:" << expectedValue); result = false; goto line1; } @@ -145,9 +152,20 @@ line1: HIP_CHECK(hipFreeArray(arr)); free(hData); REQUIRE(result); - } +/** + * Test Description + * ------------------------ + * - Uses different addressing and filtering modes for 3D arrays. + * Test source + * ------------------------ + * - unit/texture/hipTextureObj3DCheckModes.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipTextureObj3DCheckModes") { CHECK_IMAGE_SUPPORT @@ -159,58 +177,42 @@ TEST_CASE("Unit_hipTextureObj3DCheckModes") { } SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { - runTest - (256, 256, 256, -3.9, 6.1, 9.5); - runTest - (256, 256, 256, 4.4, -7.0, 5.3); + runTest(256, 256, 256, -3.9, 6.1, 9.5); + runTest(256, 256, 256, 4.4, -7.0, 5.3); } SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { - runTest - (256, 256, 256, -8.5, 2.9, 5.8); - runTest - (256, 256, 256, 12.5, 6.7, 11.4); + runTest(256, 256, 256, -8.5, 2.9, 5.8); + runTest(256, 256, 256, 12.5, 6.7, 11.4); } SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") { - runTest - (256, 256, 256, -0.4, -0.4, -0.4); - runTest - (256, 256, 256, 4, 14.6, -0.3); + runTest(256, 256, 256, -0.4, -0.4, -0.4); + runTest(256, 256, 256, 4, 14.6, -0.3); } SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { - runTest - (256, 256, 256, 6.9, 7.4, 0.4); - runTest - (256, 256, 256, 12.5, 23.7, 0.34); + runTest(256, 256, 256, 6.9, 7.4, 0.4); + runTest(256, 256, 256, 12.5, 23.7, 0.34); } SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { - runTest - (256, 256, 256, -3, 8.9, -4); - runTest - (256, 256, 256, 4, -0.1, 8.2); + runTest(256, 256, 256, -3, 8.9, -4); + runTest(256, 256, 256, 4, -0.1, 8.2); } SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { - runTest - (256, 256, 256, -8.5, 15.9, 0.1); - runTest - (256, 256, 256, 12.5, -17.9, -0.35); + runTest(256, 256, 256, -8.5, 15.9, 0.1); + runTest(256, 256, 256, 12.5, -17.9, -0.35); } SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { - runTest - (256, 256, 256, -3, 5.8, 0.89); - runTest - (256, 256, 256, 4, 9.1, 2.08); + runTest(256, 256, 256, -3, 5.8, 0.89); + runTest(256, 256, 256, 4, 9.1, 2.08); } SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { - runTest - (256, 256, 256, -8.5, 6.6, 3.67); - runTest - (256, 256, 256, 12.5, 0.01, -9.9); + runTest(256, 256, 256, -8.5, 6.6, 3.67); + runTest(256, 256, 256, 12.5, 0.01, -9.9); } } diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObjectTests.cc b/projects/hip-tests/catch/unit/texture/hipTextureObjectTests.cc new file mode 100644 index 0000000000..aa19fe2520 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureObjectTests.cc @@ -0,0 +1,369 @@ +/* +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 + + +class TextureObjectTestWrapper { + private: + float* host_data_; + bool ommit_destroy_; + + public: + hipTextureObject_t texture_object = 0; + hipResourceDesc res_desc; + hipTextureDesc tex_desc; + hipChannelFormatDesc channel_desc; + hipResourceViewDesc res_vew_desc; + hipArray* array_member; + size_t size; /* size in bytes*/ + int width; /* width in elements */ + + TextureObjectTestWrapper(bool useResourceViewDescriptor, bool ommitDestroy = false) + : ommit_destroy_(ommitDestroy), width(128) { + int i; + size = width * sizeof(float); + + host_data_ = (float*)malloc(size); + memset(host_data_, 0, size); + + for (i = 0; i < width; i++) { + host_data_[i] = i; + } + + channel_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipMallocArray(&array_member, &channel_desc, width); + + HIP_CHECK( + hipMemcpy2DToArray(array_member, 0, 0, host_data_, size, size, 1, hipMemcpyHostToDevice)); + + memset(&res_desc, 0, sizeof(res_desc)); + res_desc.resType = hipResourceTypeArray; + res_desc.res.array.array = array_member; + + memset(&tex_desc, 0, sizeof(tex_desc)); + tex_desc.addressMode[0] = hipAddressModeClamp; + tex_desc.filterMode = hipFilterModePoint; + tex_desc.readMode = hipReadModeElementType; + tex_desc.normalizedCoords = false; + + memset(&res_vew_desc, 0, sizeof(res_vew_desc)); + + if (useResourceViewDescriptor) { +#if HT_AMD + res_vew_desc.format = hipResViewFormatFloat1; + res_vew_desc.width = size; +#else + std::cout << "Resource View Descriptors are not supported on NVIDIA currently" << std::endl; + useResourceViewDescriptor = false; +#endif + } + + + HIP_CHECK(hipCreateTextureObject(&texture_object, &res_desc, &tex_desc, + useResourceViewDescriptor ? &res_vew_desc : nullptr)); + } + + ~TextureObjectTestWrapper() { + if (!ommit_destroy_) { + HIP_CHECK(hipDestroyTextureObject(texture_object)); + } + HIP_CHECK(hipFreeArray(array_member)); + free(host_data_); + } +}; + +/** + * @addtogroup hipGetTextureObjectResourceDesc hipGetTextureObjectResourceDesc + * @{ + * @ingroup TextureTest + * `hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, + * hipTextureObject_t textureObject)` - + * Gets resource descriptor for the texture object. + */ + +/** + * Test Description + * ------------------------ + * - Creates new texture object and an empty resource descriptor. + * - Gets resource descriptor from the texture. + * - Compares it to the empty created one. + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetTextureObjectResourceDesc_positive") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(false); + + hipResourceDesc check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + HIP_CHECK(hipGetTextureObjectResourceDesc(&check_desc, tex_obj_wrapper.texture_object)); + + REQUIRE(check_desc.resType == tex_obj_wrapper.res_desc.resType); + REQUIRE(check_desc.res.array.array == tex_obj_wrapper.res_desc.res.array.array); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the resource descriptor is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When the texture is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGetTextureObjectResourceDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(false); + + hipResourceDesc check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipGetTextureObjectResourceDesc(nullptr, tex_obj_wrapper.texture_object), + hipErrorInvalidValue); + } + + SECTION("texture is invalid") { + HIP_CHECK_ERROR( + hipGetTextureObjectResourceDesc(&check_desc, static_cast(0)), + hipErrorInvalidValue); + } +} + +/** + * End doxygen group hipGetTextureObjectResourceDesc. + * @} + */ + +/** + * @addtogroup hipGetTextureObjectResourceViewDesc hipGetTextureObjectResourceViewDesc + * @{ + * @ingroup TextureTest + * `hipGetTextureObjectResourceViewDesc(struct hipResourceViewDesc* pResViewDesc, + * hipTextureObject_t textureObject)` - + * Gets resource view descriptor for the texture object. + */ + +/** + * Test Description + * ------------------------ + * - Creates new texture object and an empty resource view descriptor. + * - Gets resource view descriptor from the texture. + * - Compares it to the empty created one. + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipGetTextureObjectResourceViewDesc_positive") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(true); + + hipResourceViewDesc check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + HIP_CHECK(hipGetTextureObjectResourceViewDesc(&check_desc, tex_obj_wrapper.texture_object)); + + REQUIRE(check_desc.format == tex_obj_wrapper.res_vew_desc.format); + REQUIRE(check_desc.width == tex_obj_wrapper.res_vew_desc.width); +} +#endif + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the resource view descriptor is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When the texture is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipGetTextureObjectResourceViewDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(true); + + hipResourceViewDesc check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipGetTextureObjectResourceViewDesc(nullptr, tex_obj_wrapper.texture_object), + hipErrorInvalidValue); + } + + SECTION("texture is invalid") { + HIP_CHECK_ERROR( + hipGetTextureObjectResourceViewDesc(&check_desc, static_cast(0)), + hipErrorInvalidValue); + } + + HipTest::HIP_SKIP_TEST("Skipping on NVIDIA platform"); +} +#endif + +/** + * End doxygen group hipGetTextureObjectResourceViewDesc. + * @} + */ + +/** + * @addtogroup hipGetTextureObjectTextureDesc hipGetTextureObjectTextureDesc + * @{ + * @ingroup TextureTest + * `hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, + * hipTextureObject_t textureObject)` - + * Gets texture descriptor for the texture object. + */ + +/** + * Test Description + * ------------------------ + * - Creates new texture object and an empty texture descriptor. + * - Gets texture descriptor from the texture. + * - Compares it to the empty created one. + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipGetTextureObjectTextureDesc_positive") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(false); + + hipTextureDesc check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + HIP_CHECK(hipGetTextureObjectTextureDesc(&check_desc, tex_obj_wrapper.texture_object)); + + REQUIRE(check_desc.addressMode[0] == tex_obj_wrapper.tex_desc.addressMode[0]); + REQUIRE(check_desc.filterMode == tex_obj_wrapper.tex_desc.filterMode); + REQUIRE(check_desc.readMode == tex_obj_wrapper.tex_desc.readMode); + REQUIRE(check_desc.normalizedCoords == tex_obj_wrapper.tex_desc.normalizedCoords); +} +#endif + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the texture descriptor is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When the texture is not valid (0) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - Platform specific (AMD) + * - HIP_VERSION >= 5.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipGetTextureObjectTextureDesc_Negative_Parameters") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(false); + + hipTextureDesc check_desc; + memset(&check_desc, 0, sizeof(check_desc)); + + SECTION("desc is nullptr") { + HIP_CHECK_ERROR(hipGetTextureObjectTextureDesc(nullptr, tex_obj_wrapper.texture_object), + hipErrorInvalidValue); + } + + SECTION("texture is invalid") { + HIP_CHECK_ERROR(hipGetTextureObjectTextureDesc(&check_desc, static_cast(0)), + hipErrorInvalidValue); + } +} +#endif + +/** + * End doxygen group hipGetTextureObjectTextureDesc. + * @} + */ + +/** + * @addtogroup hipDestroyTextureObject hipDestroyTextureObject + * @{ + * @ingroup TextureTest + * `hipDestroyTextureObject(hipTextureObject_t textureObject)` - + * Destroys a texture object. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipCreateTextureObject_ArgValidation + */ + +/** + * Test Description + * ------------------------ + * - Successfully destroys regular texture object. + * Test source + * ------------------------ + * - unit/texture/hipTextureObjectTests.cc + * Test requirements + * ------------------------ + * - Textures supported on device + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipDestroyTextureObject_positive") { + CHECK_IMAGE_SUPPORT; + + TextureObjectTestWrapper tex_obj_wrapper(false, true); + REQUIRE(hipDestroyTextureObject(tex_obj_wrapper.texture_object) == hipSuccess); +}