From 5fe762ae5deea9a23516c87613c89d8e72d15cec Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 14:11:27 +0100 Subject: [PATCH] EXSWHTEC-319 - Implement tests for 3D texture device functions #368 Change-Id: Ifc7a6feae1a9567df29d35928d8c4328497bfc97 --- catch/hipTestMain/config/config_amd_windows | 9 + catch/unit/texture/CMakeLists.txt | 3 + catch/unit/texture/kernels.hh | 42 +++++ catch/unit/texture/tex3D.cc | 192 ++++++++++++++++++++ catch/unit/texture/tex3DGrad.cc | 191 +++++++++++++++++++ catch/unit/texture/tex3DLod.cc | 191 +++++++++++++++++++ 6 files changed, 628 insertions(+) create mode 100644 catch/unit/texture/tex3D.cc create mode 100644 catch/unit/texture/tex3DGrad.cc create mode 100644 catch/unit/texture/tex3DLod.cc diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index b3b396f6ba..912f7ecef0 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -211,6 +211,15 @@ "Unit_hipHostMalloc_AllocateUseMoreThanAvailGPUMemory", "=== SWDEV-432250:Below tests failed in stress test on 10/11/23 ===", "Unit_hipVectorTypes_test_on_device", +<<<<<<< HEAD +======= + "Unit_Layered1DTexture_Check_DeviceBufferToFromLayered1DArray - ushort4", + "Unit_Layered2DTexture_Check_DeviceBufferToFromLayered2DArray - float4", + "Unit_tex3DLod_Positive_ReadModeElementType", + "Unit_tex3DLod_Positive_ReadModeNormalizedFloat", + "Unit_tex3DGrad_Positive_ReadModeElementType", + "Unit_tex3DGrad_Positive_ReadModeNormalizedFloat", +>>>>>>> 5cdc6efc (Merge branch 'develop' into tex3D_tests) "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", "NOTE: The following test is disabled due to defect - EXSWHTEC-241", diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index 14a74de8f5..85f653573c 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -60,6 +60,9 @@ set(TEST_SRC tex1DLayeredGrad.cc tex1DLayeredLod.cc tex1DLod.cc + tex3D.cc + tex3DLod.cc + tex3DGrad.cc ) if(WIN32) diff --git a/catch/unit/texture/kernels.hh b/catch/unit/texture/kernels.hh index ac5d73ff40..ec4e1449ef 100644 --- a/catch/unit/texture/kernels.hh +++ b/catch/unit/texture/kernels.hh @@ -131,6 +131,48 @@ __global__ void tex3DKernel(TexelType* const out, size_t N_x, size_t N_y, size_t out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3D(tex_obj, x, y, z); } +template +__global__ void tex3DLodKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t depth, size_t num_subdivisions, bool normalized_coords, + float level) { + const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x; + if (tid_x >= N_x) return; + + const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y; + if (tid_y >= N_y) return; + + const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z; + if (tid_z >= N_z) return; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords); + + out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3DLod(tex_obj, x, y, z, level); +} + +template +__global__ void tex3DGradKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t depth, size_t num_subdivisions, bool normalized_coords, + float4 dx, float4 dy) { + const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x; + if (tid_x >= N_x) return; + + const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y; + if (tid_y >= N_y) return; + + const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z; + if (tid_z >= N_z) return; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords); + + out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3DGrad(tex_obj, x, y, z, dx, dy); +} + template __global__ void tex1DLayeredKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, size_t width, size_t num_subdivisions, bool normalized_coords, diff --git a/catch/unit/texture/tex3D.cc b/catch/unit/texture/tex3D.cc new file mode 100644 index 0000000000..ecfdd535bb --- /dev/null +++ b/catch/unit/texture/tex3D.cc @@ -0,0 +1,192 @@ +/* +Copyright (c) 2023 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 + +#include "kernels.hh" +#include "test_fixture.hh" + +/** + * @addtogroup tex3D tex3D + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex3D` and read mode set to `hipReadModeElementType`. The + * test is performed with: + * - normalized coordinates + * - non-normalized coordinates + * - Nearest-point sampling + * - Linear filtering + * - All combinations of different addressing modes. + * Test source + * ------------------------ + * - unit/texture/tex3D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 4, 2); + params.num_subdivisions = 2; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY()); + const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + dim_grid.z = num_blocks_z; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + dim_block.z = num_threads_z; + + tex3DKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(), + fixture.tex.object(), params.Width(), params.Height(), params.Depth(), + params.num_subdivisions, params.tex_desc.normalizedCoords); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumIters(); ++i) { + const auto plane = i % (params.NumItersX() * params.NumItersY()); + float x = plane % params.NumItersX(); + float y = plane / params.NumItersX(); + float z = i / (params.NumItersX() * params.NumItersY()); + + x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1])); + INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + INFO("z: " << std::fixed << std::setprecision(16) << z); + + const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc); + REQUIRE(ref_val.x == fixture.out_alloc_h[i].x); + REQUIRE(ref_val.y == fixture.out_alloc_h[i].y); + REQUIRE(ref_val.z == fixture.out_alloc_h[i].z); + REQUIRE(ref_val.w == fixture.out_alloc_h[i].w); + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex3D` and read mode set to `hipReadModeNormalizedFloat`. The + * test is performed with: + * - normalized coordinates + * - non-normalized coordinates + * - Nearest-point sampling + * - Linear filtering + * - All combinations of different addressing modes. + * Test source + * ------------------------ + * - unit/texture/tex3D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, + unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 2); + params.num_subdivisions = 2; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY()); + const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + dim_grid.z = num_blocks_z; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + dim_block.z = num_threads_z; + + tex3DKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(), + fixture.tex.object(), params.Width(), params.Height(), params.Depth(), + params.num_subdivisions, params.tex_desc.normalizedCoords); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumIters(); ++i) { + const auto plane = i % (params.NumItersX() * params.NumItersY()); + float x = plane % params.NumItersX(); + float y = plane / params.NumItersX(); + float z = i / (params.NumItersX() * params.NumItersY()); + + x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1])); + INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + INFO("z: " << std::fixed << std::setprecision(16) << z); + + + auto ref_val = Vec4Map(fixture.tex_h.Tex3D(x, y, z, params.tex_desc), + NormalizeInteger); + REQUIRE(ref_val.x == fixture.out_alloc_h[i].x); + REQUIRE(ref_val.y == fixture.out_alloc_h[i].y); + REQUIRE(ref_val.z == fixture.out_alloc_h[i].z); + REQUIRE(ref_val.w == fixture.out_alloc_h[i].w); + } +} diff --git a/catch/unit/texture/tex3DGrad.cc b/catch/unit/texture/tex3DGrad.cc new file mode 100644 index 0000000000..a3f3d8ddfd --- /dev/null +++ b/catch/unit/texture/tex3DGrad.cc @@ -0,0 +1,191 @@ +/* +Copyright (c) 2023 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 + +#include "kernels.hh" +#include "test_fixture.hh" + +/** + * @addtogroup tex3DGrad tex3DGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex3DGrad` and read mode set to `hipReadModeElementType`. The + * test is performed with: + * - normalized coordinates + * - non-normalized coordinates + * - Nearest-point sampling + * - Linear filtering + * - All combinations of different addressing modes. + * Test source + * ------------------------ + * - unit/texture/tex3DGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 2); + params.num_subdivisions = 2; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY()); + const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + dim_grid.z = num_blocks_z; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + dim_block.z = num_threads_z; + + tex3DGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(), + fixture.tex.object(), params.Width(), params.Height(), params.Depth(), + params.num_subdivisions, params.tex_desc.normalizedCoords, float4{}, float4{}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumIters(); ++i) { + const auto plane = i % (params.NumItersX() * params.NumItersY()); + float x = plane % params.NumItersX(); + float y = plane / params.NumItersX(); + float z = i / (params.NumItersX() * params.NumItersY()); + + x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1])); + INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + INFO("z: " << std::fixed << std::setprecision(16) << z); + + const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc); + REQUIRE(ref_val.x == fixture.out_alloc_h[i].x); + REQUIRE(ref_val.y == fixture.out_alloc_h[i].y); + REQUIRE(ref_val.z == fixture.out_alloc_h[i].z); + REQUIRE(ref_val.w == fixture.out_alloc_h[i].w); + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex3DGrad` and read mode set to `hipReadModeNormalizedFloat`. + * The test is performed with: + * - normalized coordinates + * - non-normalized coordinates + * - Nearest-point sampling + * - Linear filtering + * - All combinations of different addressing modes. + * Test source + * ------------------------ + * - unit/texture/tex3DGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 2); + params.num_subdivisions = 2; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY()); + const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + dim_grid.z = num_blocks_z; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + dim_block.z = num_threads_z; + + tex3DGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(), + fixture.tex.object(), params.Width(), params.Height(), params.Depth(), + params.num_subdivisions, params.tex_desc.normalizedCoords, float4{}, float4{}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumIters(); ++i) { + const auto plane = i % (params.NumItersX() * params.NumItersY()); + float x = plane % params.NumItersX(); + float y = plane / params.NumItersX(); + float z = i / (params.NumItersX() * params.NumItersY()); + + x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1])); + INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + INFO("z: " << std::fixed << std::setprecision(16) << z); + + auto ref_val = Vec4Map(fixture.tex_h.Tex3D(x, y, z, params.tex_desc), + NormalizeInteger); + REQUIRE(ref_val.x == fixture.out_alloc_h[i].x); + REQUIRE(ref_val.y == fixture.out_alloc_h[i].y); + REQUIRE(ref_val.z == fixture.out_alloc_h[i].z); + REQUIRE(ref_val.w == fixture.out_alloc_h[i].w); + } +} diff --git a/catch/unit/texture/tex3DLod.cc b/catch/unit/texture/tex3DLod.cc new file mode 100644 index 0000000000..bd5ee5f7a5 --- /dev/null +++ b/catch/unit/texture/tex3DLod.cc @@ -0,0 +1,191 @@ +/* +Copyright (c) 2023 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 + +#include "kernels.hh" +#include "test_fixture.hh" + +/** + * @addtogroup tex3DLod tex3DLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex3DLod` and read mode set to `hipReadModeElementType`. The + * test is performed with: + * - normalized coordinates + * - non-normalized coordinates + * - Nearest-point sampling + * - Linear filtering + * - All combinations of different addressing modes. + * Test source + * ------------------------ + * - unit/texture/tex3DLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 2); + params.num_subdivisions = 2; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY()); + const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + dim_grid.z = num_blocks_z; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + dim_block.z = num_threads_z; + + tex3DLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(), + fixture.tex.object(), params.Width(), params.Height(), params.Depth(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0.0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumIters(); ++i) { + const auto plane = i % (params.NumItersX() * params.NumItersY()); + float x = plane % params.NumItersX(); + float y = plane / params.NumItersX(); + float z = i / (params.NumItersX() * params.NumItersY()); + + x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1])); + INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + INFO("z: " << std::fixed << std::setprecision(16) << z); + + const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc); + REQUIRE(ref_val.x == fixture.out_alloc_h[i].x); + REQUIRE(ref_val.y == fixture.out_alloc_h[i].y); + REQUIRE(ref_val.z == fixture.out_alloc_h[i].z); + REQUIRE(ref_val.w == fixture.out_alloc_h[i].w); + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex3DLod` and read mode set to `hipReadModeNormalizedFloat`. The + * test is performed with: + * - normalized coordinates + * - non-normalized coordinates + * - Nearest-point sampling + * - Linear filtering + * - All combinations of different addressing modes. + * Test source + * ------------------------ + * - unit/texture/tex3DLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, + unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 2); + params.num_subdivisions = 2; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY()); + const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + dim_grid.z = num_blocks_z; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + dim_block.z = num_threads_z; + + tex3DLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(), + fixture.tex.object(), params.Width(), params.Height(), params.Depth(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0.0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumIters(); ++i) { + const auto plane = i % (params.NumItersX() * params.NumItersY()); + float x = plane % params.NumItersX(); + float y = plane / params.NumItersX(); + float z = i / (params.NumItersX() * params.NumItersY()); + + x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1])); + INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + INFO("z: " << std::fixed << std::setprecision(16) << z); + + auto ref_val = Vec4Map(fixture.tex_h.Tex3D(x, y, z, params.tex_desc), + NormalizeInteger); + REQUIRE(ref_val.x == fixture.out_alloc_h[i].x); + REQUIRE(ref_val.y == fixture.out_alloc_h[i].y); + REQUIRE(ref_val.z == fixture.out_alloc_h[i].z); + REQUIRE(ref_val.w == fixture.out_alloc_h[i].w); + } +}