From 90a783c959172aa3887876aa6ac99b5744d2c6a2 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:22:50 +0100 Subject: [PATCH] EXSWHTEC-358 - Implement tests for 2D mipmapped texture device functions #432 Change-Id: I8d517ef95eeddfcc442f00e331e3ed0d1123d99c --- catch/hipTestMain/config/config_amd_linux | 11 ++ catch/unit/texture/CMakeLists.txt | 7 + catch/unit/texture/kernels.hh | 83 ++++++++++ catch/unit/texture/tex2D.cc | 87 +++++----- catch/unit/texture/tex2DGrad.cc | 175 +++++++++++++++++++++ catch/unit/texture/tex2DLayered.cc | 90 +++++------ catch/unit/texture/tex2DLayeredGrad.cc | 183 ++++++++++++++++++++++ catch/unit/texture/tex2DLayeredLod.cc | 183 ++++++++++++++++++++++ catch/unit/texture/tex2DLod.cc | 175 +++++++++++++++++++++ catch/unit/texture/tex2Dgather.cc | 104 ++++++++++++ catch/unit/texture/texture_reference.hh | 36 +++++ 11 files changed, 1045 insertions(+), 89 deletions(-) create mode 100644 catch/unit/texture/tex2DGrad.cc create mode 100644 catch/unit/texture/tex2DLayeredGrad.cc create mode 100644 catch/unit/texture/tex2DLayeredLod.cc create mode 100644 catch/unit/texture/tex2DLod.cc create mode 100644 catch/unit/texture/tex2Dgather.cc diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 66107df9e1..99c33f7834 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -128,8 +128,19 @@ "Unit_hipEventIpc", "=== SWDEV-427101:Below test fails randomly in PSDB ===", "Unit_deviceAllocation_InOneThread_AccessInAllThreads", +<<<<<<< HEAD "=== Below test is disabled due to defect EXSWHTEC-347 ===", "Unit_hipPointerSetAttribute_Positive_SyncMemops", +======= + "Unit_tex2DLod_Positive_ReadModeElementType", + "Unit_tex2DLod_Positive_ReadModeNormalizedFloat", + "Unit_tex2DLayeredLod_Positive_ReadModeElementType", + "Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat", + "Unit_tex2DGrad_Positive_ReadModeElementType", + "Unit_tex2DGrad_Positive_ReadModeNormalizedFloat", + "Unit_tex2DLayeredGrad_Positive_ReadModeElementType", + "Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat", +>>>>>>> ab1704d9 (Merge branch 'develop' into tex2D_mipmap_tests) "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index 18881833b9..d1d56d7064 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -69,6 +69,13 @@ set(TEST_SRC texCubemapLayered.cc texCubemapLayeredLod.cc texCubemapLayeredGrad.cc + tex2Dgather.cc + tex2D.cc + tex2DLayered.cc + tex2DGrad.cc + tex2DLayeredGrad.cc + tex2DLod.cc + tex2DLayeredLod.cc ) if(WIN32) diff --git a/catch/unit/texture/kernels.hh b/catch/unit/texture/kernels.hh index 2e0bf256cd..d74caa5b2a 100644 --- a/catch/unit/texture/kernels.hh +++ b/catch/unit/texture/kernels.hh @@ -95,6 +95,22 @@ __global__ void tex1DLayeredGradKernel(TexelType* const out, size_t N, hipTextur out[tid] = tex1DLayeredGrad(tex_obj, x, layer, dx, dy); } +template +__global__ void tex2DgatherKernel(TexelType* const out, int comp, size_t N_x, size_t N_y, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t num_subdivisions, bool normalized_coords) { + 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; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2Dgather(tex_obj, x, y, comp); +} + template __global__ void tex2DKernel(TexelType* const out, size_t N_x, size_t N_y, hipTextureObject_t tex_obj, size_t width, size_t height, @@ -111,6 +127,73 @@ __global__ void tex2DKernel(TexelType* const out, size_t N_x, size_t N_y, out[tid_y * N_x + tid_x] = tex2D(tex_obj, x, y); } +template +__global__ void tex2DGradKernel(TexelType* const out, size_t N_x, size_t N_y, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t num_subdivisions, bool normalized_coords, float2 dx, + float2 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; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2DGrad(tex_obj, x, y, dx, dy); +} + +template +__global__ void tex2DLayeredGradKernel(TexelType* const out, size_t N_x, size_t N_y, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t num_subdivisions, bool normalized_coords, float layer, + float2 dx, float2 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; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2DLayeredGrad(tex_obj, x, y, layer, dx, dy); +} + +template +__global__ void tex2DLodKernel(TexelType* const out, size_t N_x, size_t N_y, + hipTextureObject_t tex_obj, size_t width, size_t height, + 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; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2DLod(tex_obj, x, y, level); +} + +template +__global__ void tex2DLayeredLodKernel(TexelType* const out, size_t N_x, size_t N_y, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t num_subdivisions, bool normalized_coords, int layer, + 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; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2DLayeredLod(tex_obj, x, y, layer, level); +} + template __global__ void tex3DKernel(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, diff --git a/catch/unit/texture/tex2D.cc b/catch/unit/texture/tex2D.cc index 79d6055ede..7b31a03944 100644 --- a/catch/unit/texture/tex2D.cc +++ b/catch/unit/texture/tex2D.cc @@ -50,7 +50,9 @@ THE SOFTWARE. */ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeElementType", "", char, unsigned char, short, unsigned short, int, unsigned int, float) { - TextureTestParams params = {0}; + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; params.extent = make_hipExtent(16, 4, 0); params.num_subdivisions = 4; params.GenerateTextureDesc(); @@ -75,29 +77,27 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeElementType", "", char, unsigned fixture.LoadOutput(); - for (auto j = 0u; j < params.NumItersY(); ++j) { - for (auto i = 0u; i < params.NumItersX(); ++i) { - float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, - params.tex_desc.normalizedCoords); - float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions, - params.tex_desc.normalizedCoords); + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); - INFO("i: " << i); - INFO("j: " << j); - 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("x: " << std::fixed << std::setprecision(16) << x); - INFO("y: " << std::fixed << std::setprecision(16) << y); + 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); - auto index = j * params.NumItersX() + i; + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); - const auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc); - REQUIRE(ref_val.x == fixture.out_alloc_h[index].x); - REQUIRE(ref_val.y == fixture.out_alloc_h[index].y); - REQUIRE(ref_val.z == fixture.out_alloc_h[index].z); - REQUIRE(ref_val.w == fixture.out_alloc_h[index].w); - } + const auto ref_val = fixture.tex_h.Tex2D(x, y, 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); } } @@ -120,7 +120,9 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeElementType", "", char, unsigned */ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, unsigned short) { - TextureTestParams params = {0}; + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; params.extent = make_hipExtent(16, 4, 0); params.num_subdivisions = 4; params.GenerateTextureDesc(hipReadModeNormalizedFloat); @@ -145,30 +147,27 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeNormalizedFloat", "", char, unsi fixture.LoadOutput(); - for (auto j = 0u; j < params.NumItersY(); ++j) { - for (auto i = 0u; i < params.NumItersX(); ++i) { - float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, - params.tex_desc.normalizedCoords); - float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions, - params.tex_desc.normalizedCoords); + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersY(); - INFO("i: " << i); - INFO("j: " << j); - INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); - 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("x: " << std::fixed << std::setprecision(16) << x); - INFO("y: " << std::fixed << std::setprecision(16) << y); + 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); - auto index = j * params.NumItersX() + i; + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); - auto ref_val = - Vec4Map(fixture.tex_h.Tex2D(x, y, params.tex_desc), NormalizeInteger); - REQUIRE(ref_val.x == fixture.out_alloc_h[index].x); - REQUIRE(ref_val.y == fixture.out_alloc_h[index].y); - REQUIRE(ref_val.z == fixture.out_alloc_h[index].z); - REQUIRE(ref_val.w == fixture.out_alloc_h[index].w); - } + auto ref_val = + Vec4Map(fixture.tex_h.Tex2D(x, y, 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); } } \ No newline at end of file diff --git a/catch/unit/texture/tex2DGrad.cc b/catch/unit/texture/tex2DGrad.cc new file mode 100644 index 0000000000..939a6a2543 --- /dev/null +++ b/catch/unit/texture/tex2DGrad.cc @@ -0,0 +1,175 @@ +/* +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 tex2DGrad tex2DGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex2DGrad` 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/tex2DGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DGrad_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + tex2DGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords, + float2{0.5f, 0.5f}, float2{0.5f, 0.5f}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + const auto ref_val = fixture.tex_h.Tex2D(x, y, 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 `tex2DGrad` 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/tex2DGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + tex2DGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords, + float2{0.5f, 0.5f}, float2{0.5f, 0.5f}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex2D(x, y, 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); + } +} \ No newline at end of file diff --git a/catch/unit/texture/tex2DLayered.cc b/catch/unit/texture/tex2DLayered.cc index b05a2e0a32..4929a5b3e0 100644 --- a/catch/unit/texture/tex2DLayered.cc +++ b/catch/unit/texture/tex2DLayered.cc @@ -50,7 +50,9 @@ THE SOFTWARE. */ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeElementType", "", char, unsigned char, short, unsigned short, int, unsigned int, float) { - TextureTestParams params = {0}; + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; params.extent = make_hipExtent(16, 4, 0); params.layers = 2; params.num_subdivisions = 4; @@ -78,30 +80,28 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeElementType", "", char, u fixture.LoadOutput(); - for (auto j = 0u; j < params.NumItersY(); ++j) { - for (auto i = 0u; i < params.NumItersX(); ++i) { - float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, - params.tex_desc.normalizedCoords); - float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions, - params.tex_desc.normalizedCoords); + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); - INFO("Layer: " << layer); - INFO("i: " << i); - INFO("j: " << j); - 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("x: " << std::fixed << std::setprecision(16) << x); - INFO("y: " << std::fixed << std::setprecision(16) << y); + 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); - auto index = j * params.NumItersX() + i; + INFO("Layer: " << layer); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); - const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc); - REQUIRE(ref_val.x == fixture.out_alloc_h[index].x); - REQUIRE(ref_val.y == fixture.out_alloc_h[index].y); - REQUIRE(ref_val.z == fixture.out_alloc_h[index].z); - REQUIRE(ref_val.w == fixture.out_alloc_h[index].w); - } + const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, 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); } } } @@ -125,7 +125,9 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeElementType", "", char, u */ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, unsigned short) { - TextureTestParams params = {0}; + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; params.extent = make_hipExtent(16, 4, 0); params.layers = 2; params.num_subdivisions = 4; @@ -153,31 +155,29 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeNormalizedFloat", "", cha fixture.LoadOutput(); - for (auto j = 0u; j < params.NumItersY(); ++j) { - for (auto i = 0u; i < params.NumItersX(); ++i) { - float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, - params.tex_desc.normalizedCoords); - float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions, - params.tex_desc.normalizedCoords); + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); - INFO("Layer: " << layer); - INFO("i: " << i); - INFO("j: " << j); - 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("x: " << std::fixed << std::setprecision(16) << x); - INFO("y: " << std::fixed << std::setprecision(16) << y); + 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); - auto index = j * params.NumItersX() + i; + INFO("Layer: " << layer); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); - auto ref_val = Vec4Map(fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc), - NormalizeInteger); - REQUIRE(ref_val.x == fixture.out_alloc_h[index].x); - REQUIRE(ref_val.y == fixture.out_alloc_h[index].y); - REQUIRE(ref_val.z == fixture.out_alloc_h[index].z); - REQUIRE(ref_val.w == fixture.out_alloc_h[index].w); - } + auto ref_val = Vec4Map(fixture.tex_h.Tex2DLayered(x, y, layer, 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); } } } \ No newline at end of file diff --git a/catch/unit/texture/tex2DLayeredGrad.cc b/catch/unit/texture/tex2DLayeredGrad.cc new file mode 100644 index 0000000000..d58920bf56 --- /dev/null +++ b/catch/unit/texture/tex2DLayeredGrad.cc @@ -0,0 +1,183 @@ +/* +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 tex2DLayeredGrad tex2DLayeredGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex2DLayeredGrad` 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/tex2DLayeredGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DLayeredGrad_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.layers = 2; + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + for (auto layer = 0u; layer < params.layers; ++layer) { + tex2DLayeredGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords, + layer, float2{0.5f, 0.5f}, float2{0.5f, 0.5f}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Layer: " << layer); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, 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 `tex2DLayeredGrad` 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/tex2DLayeredGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat", "", char, + unsigned char, short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.layers = 2; + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + for (auto layer = 0u; layer < params.layers; ++layer) { + tex2DLayeredGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords, + layer, float2{0.5f, 0.5f}, float2{0.5f, 0.5f}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Layer: " << layer); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + auto ref_val = Vec4Map(fixture.tex_h.Tex2DLayered(x, y, layer, 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); + } + } +} \ No newline at end of file diff --git a/catch/unit/texture/tex2DLayeredLod.cc b/catch/unit/texture/tex2DLayeredLod.cc new file mode 100644 index 0000000000..e51000c204 --- /dev/null +++ b/catch/unit/texture/tex2DLayeredLod.cc @@ -0,0 +1,183 @@ +/* +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 tex2DLayeredLod tex2DLayeredLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex2DLayeredLod` 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/tex2DLayeredLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DLayeredLod_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.layers = 2; + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + for (auto layer = 0u; layer < params.layers; ++layer) { + tex2DLayeredLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords, + layer, 0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Layer: " << layer); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, 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 `tex2DLayeredLod` 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/tex2DLayeredLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.layers = 2; + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + for (auto layer = 0u; layer < params.layers; ++layer) { + tex2DLayeredLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords, + layer, 0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Layer: " << layer); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + auto ref_val = Vec4Map(fixture.tex_h.Tex2DLayered(x, y, layer, 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); + } + } +} \ No newline at end of file diff --git a/catch/unit/texture/tex2DLod.cc b/catch/unit/texture/tex2DLod.cc new file mode 100644 index 0000000000..e875e09133 --- /dev/null +++ b/catch/unit/texture/tex2DLod.cc @@ -0,0 +1,175 @@ +/* +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 tex2DLod tex2DLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex2DLod` 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/tex2DLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DLod_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + tex2DLodKernel> + <<>>(fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), + fixture.tex.object(), params.Width(), params.Height(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + const auto ref_val = fixture.tex_h.Tex2D(x, y, 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 `tex2DLod` 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/tex2DLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ +TEMPLATE_TEST_CASE("Unit_tex2DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, + unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + tex2DLodKernel> + <<>>(fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), + fixture.tex.object(), params.Width(), params.Height(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex2D(x, y, 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); + } +} \ No newline at end of file diff --git a/catch/unit/texture/tex2Dgather.cc b/catch/unit/texture/tex2Dgather.cc new file mode 100644 index 0000000000..6d1c262097 --- /dev/null +++ b/catch/unit/texture/tex2Dgather.cc @@ -0,0 +1,104 @@ +/* +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 tex2Dgather tex2Dgather + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex2Dgather` 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/tex2Dgather.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex2Dgather_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(16, 4, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX()); + const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY()); + + dim3 dim_grid; + dim_grid.x = num_blocks_x; + dim_grid.y = num_blocks_y; + + dim3 dim_block; + dim_block.x = num_threads_x; + dim_block.y = num_threads_y; + + const int comp = GENERATE(0, 1, 2, 3); + + tex2DgatherKernel><<>>( + fixture.out_alloc_d.ptr(), comp, params.NumItersX(), params.NumItersY(), fixture.tex.object(), + params.Width(), params.Height(), params.num_subdivisions, params.tex_desc.normalizedCoords); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) { + float x = i % params.NumItersX(); + float y = i / params.NumItersX(); + + 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); + + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + 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("x: " << std::fixed << std::setprecision(16) << x); + INFO("y: " << std::fixed << std::setprecision(16) << y); + + const auto ref_val = fixture.tex_h.Tex2DGather(x, y, comp, 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); + } +} \ No newline at end of file diff --git a/catch/unit/texture/texture_reference.hh b/catch/unit/texture/texture_reference.hh index 45f7dd8efa..b66fc4a045 100644 --- a/catch/unit/texture/texture_reference.hh +++ b/catch/unit/texture/texture_reference.hh @@ -35,6 +35,42 @@ template class TextureReference { return Tex1DLayered(x, 0, tex_desc); } + TexelType Tex2DGather(float x, float y, int comp, const hipTextureDesc& tex_desc) const { + x = tex_desc.normalizedCoords ? x * extent_.width : x; + y = tex_desc.normalizedCoords ? y * extent_.height : y; + + const auto [i, alpha] = GetLinearFilteringParams(x); + const auto [j, beta] = GetLinearFilteringParams(y); + + const auto T_i0j0 = Sample(i, j, 0, tex_desc.addressMode); + const auto T_i1j0 = Sample(i + 1.0f, j, 0, tex_desc.addressMode); + const auto T_i0j1 = Sample(i, j + 1.0f, 0, tex_desc.addressMode); + const auto T_i1j1 = Sample(i + 1.0f, j + 1.0f, 0, tex_desc.addressMode); + + const auto IndexVec4 = [](auto vec, int comp) { + switch (comp) { + case 0: + return vec.x; + case 1: + return vec.y; + case 2: + return vec.z; + case 3: + return vec.w; + default: + throw std::invalid_argument("Invalid gather comp"); + } + }; + + TexelType texel; + texel.x = IndexVec4(T_i0j1, comp); + texel.y = IndexVec4(T_i1j1, comp); + texel.z = IndexVec4(T_i1j0, comp); + texel.w = IndexVec4(T_i0j0, comp); + + return texel; + } + TexelType Tex2D(float x, float y, const hipTextureDesc& tex_desc) const { return Tex2DLayered(x, y, 0, tex_desc); }