From ca5e7641895844ed3d348f0cbbd0de10eb9496f8 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:05:57 +0100 Subject: [PATCH] EXSWHTEC-320 - Implement tests for cubemap texture device functions #369 Change-Id: I1a247dba4e46ed7e1045dee0467d3fdac1f70cd0 [ROCm/hip-tests commit: 6eec9aa90d9ef7fed025914f8c3834b4175514d1] --- .../hipTestMain/config/config_amd_windows | 7 - .../catch/unit/texture/CMakeLists.txt | 6 + .../hip-tests/catch/unit/texture/kernels.hh | 130 +++++++++++ .../catch/unit/texture/test_fixture.hh | 7 +- .../catch/unit/texture/texCubemap.cc | 201 +++++++++++++++++ .../catch/unit/texture/texCubemapGrad.cc | 201 +++++++++++++++++ .../catch/unit/texture/texCubemapLayered.cc | 209 ++++++++++++++++++ .../unit/texture/texCubemapLayeredGrad.cc | 209 ++++++++++++++++++ .../unit/texture/texCubemapLayeredLod.cc | 209 ++++++++++++++++++ .../catch/unit/texture/texCubemapLod.cc | 201 +++++++++++++++++ .../catch/unit/texture/texture_reference.hh | 69 +++++- 11 files changed, 1439 insertions(+), 10 deletions(-) create mode 100644 projects/hip-tests/catch/unit/texture/texCubemap.cc create mode 100644 projects/hip-tests/catch/unit/texture/texCubemapGrad.cc create mode 100644 projects/hip-tests/catch/unit/texture/texCubemapLayered.cc create mode 100644 projects/hip-tests/catch/unit/texture/texCubemapLayeredGrad.cc create mode 100644 projects/hip-tests/catch/unit/texture/texCubemapLayeredLod.cc create mode 100644 projects/hip-tests/catch/unit/texture/texCubemapLod.cc diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 912f7ecef0..724aef781e 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -211,15 +211,8 @@ "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/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index 85f653573c..18881833b9 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -63,6 +63,12 @@ set(TEST_SRC tex3D.cc tex3DLod.cc tex3DGrad.cc + texCubemap.cc + texCubemapLod.cc + texCubemapGrad.cc + texCubemapLayered.cc + texCubemapLayeredLod.cc + texCubemapLayeredGrad.cc ) if(WIN32) diff --git a/projects/hip-tests/catch/unit/texture/kernels.hh b/projects/hip-tests/catch/unit/texture/kernels.hh index ec4e1449ef..f9f7a6a41e 100644 --- a/projects/hip-tests/catch/unit/texture/kernels.hh +++ b/projects/hip-tests/catch/unit/texture/kernels.hh @@ -173,6 +173,69 @@ __global__ void tex3DGradKernel(TexelType* const out, size_t N_x, size_t N_y, si out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3DGrad(tex_obj, x, y, z, dx, dy); } +template +__global__ void texCubemapKernel(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) { + 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] = texCubemap(tex_obj, x, y, z); +} + +template +__global__ void texCubemapLodKernel(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] = texCubemapLod(tex_obj, x, y, z, level); +} + +template +__global__ void texCubemapGradKernel(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] = + texCubemapGrad(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, @@ -198,4 +261,71 @@ __global__ void tex2DLayeredKernel(TexelType* const out, size_t N_x, size_t N_y, float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); out[tid_y * N_x + tid_x] = tex2DLayered(tex_obj, x, y, layer); +} + +template +__global__ void texCubemapLayeredKernel(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, size_t layer) { + 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] = + texCubemapLayered(tex_obj, x, y, z, layer); +} + +template +__global__ void texCubemapLayeredLodKernel(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, size_t 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; + + 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] = + texCubemapLayeredLod(tex_obj, x, y, z, layer, level); +} + +template +__global__ void texCubemapLayeredGradKernel(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, size_t layer, 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] = + texCubemapLayeredGrad(tex_obj, x, y, z, layer, dx, dy); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/test_fixture.hh b/projects/hip-tests/catch/unit/texture/test_fixture.hh index 0572fd6fae..28ca9ee2df 100644 --- a/projects/hip-tests/catch/unit/texture/test_fixture.hh +++ b/projects/hip-tests/catch/unit/texture/test_fixture.hh @@ -34,6 +34,7 @@ template struct TextureTestParams { size_t layers; size_t num_subdivisions; hipTextureDesc tex_desc; + bool cubemap; size_t Size() const { return extent.width * (extent.height ?: 1) * (extent.depth ?: 1) * (layers ?: 1); @@ -53,6 +54,10 @@ template struct TextureTestParams { size_t Depth() const { return extent.depth; } + unsigned int Flags() const { + return (Layered() ? hipArrayLayered : 0u) | (cubemap ? hipArrayCubemap : 0u); + } + hipExtent LayeredExtent() const { return Layered() ? make_hipExtent(Width(), Height(), layers) : extent; } @@ -115,7 +120,7 @@ template struct TextureTestFix : params{p}, host_alloc{LinearAllocs::hipHostMalloc, sizeof(VecType) * params.Size()}, tex_h{host_alloc.ptr(), params.extent, params.layers}, - tex_alloc_d{params.LayeredExtent(), params.Layered() ? hipArrayLayered : 0u}, + tex_alloc_d{params.LayeredExtent(), params.Flags()}, tex{ResDesc(), ¶ms.tex_desc}, out_alloc_d{LinearAllocs::hipMalloc, sizeof(OutType) * params.NumIters()}, out_alloc_h(params.NumIters()) {} diff --git a/projects/hip-tests/catch/unit/texture/texCubemap.cc b/projects/hip-tests/catch/unit/texture/texCubemap.cc new file mode 100644 index 0000000000..ade3775c52 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/texCubemap.cc @@ -0,0 +1,201 @@ +/* +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 texCubemap texCubemap + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemap` 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/texCubemap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.cubemap = true; + 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; + + texCubemapKernel><<>>( + 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 k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemap` 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/texCubemap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.cubemap = true; + 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; + + texCubemapKernel><<>>( + 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 k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + auto ref_val = Vec4Map(fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/texCubemapGrad.cc b/projects/hip-tests/catch/unit/texture/texCubemapGrad.cc new file mode 100644 index 0000000000..b2024737b3 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/texCubemapGrad.cc @@ -0,0 +1,201 @@ +/* +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 texCubemapGrad texCubemapGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapGrad` 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/texCubemapGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.cubemap = true; + 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; + + texCubemapGradKernel><<>>( + 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 k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapGrad` 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/texCubemapGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.cubemap = true; + 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; + + texCubemapGradKernel><<>>( + 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 k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + auto ref_val = Vec4Map(fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/texCubemapLayered.cc b/projects/hip-tests/catch/unit/texture/texCubemapLayered.cc new file mode 100644 index 0000000000..d7db8d0847 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/texCubemapLayered.cc @@ -0,0 +1,209 @@ +/* +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 texCubemapLayered texCubemapLayered + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLayered` 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/texCubemapLayered.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.layers = 1; + params.cubemap = true; + 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; + + for (auto layer = 0u; layer < params.layers; ++layer) { + texCubemapLayeredKernel><<>>( + 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, layer); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLayered` 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/texCubemapLayered.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat", "", char, + unsigned char, short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.layers = 1; + params.cubemap = true; + 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; + + for (auto layer = 0u; layer < params.layers; ++layer) { + texCubemapLayeredKernel><<>>( + 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, layer); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + auto ref_val = Vec4Map(fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/texCubemapLayeredGrad.cc b/projects/hip-tests/catch/unit/texture/texCubemapLayeredGrad.cc new file mode 100644 index 0000000000..96dd0415b6 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/texCubemapLayeredGrad.cc @@ -0,0 +1,209 @@ +/* +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 texCubemapLayeredGrad texCubemapLayeredGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLayeredGrad` 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/texCubemapLayeredGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeElementType", "", char, + unsigned char, short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.layers = 1; + params.cubemap = true; + 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; + + for (auto layer = 0u; layer < params.layers; ++layer) { + texCubemapLayeredGradKernel><<>>( + 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, layer, float4{}, float4{}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLayeredGrad` 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/texCubemapLayeredGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat", "", char, + unsigned char, short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.layers = 1; + params.cubemap = true; + 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; + + for (auto layer = 0u; layer < params.layers; ++layer) { + texCubemapLayeredGradKernel><<>>( + 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, layer, float4{}, float4{}); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + auto ref_val = Vec4Map(fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/texCubemapLayeredLod.cc b/projects/hip-tests/catch/unit/texture/texCubemapLayeredLod.cc new file mode 100644 index 0000000000..fd9db2c0b7 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/texCubemapLayeredLod.cc @@ -0,0 +1,209 @@ +/* +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 texCubemapLayeredLod texCubemapLayeredLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLayeredLod` 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/texCubemapLayeredLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeElementType", "", char, + unsigned char, short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.layers = 1; + params.cubemap = true; + 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; + + for (auto layer = 0u; layer < params.layers; ++layer) { + texCubemapLayeredLodKernel><<>>( + 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, layer, 0.0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLayeredLod` 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/texCubemapLayeredLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat", "", char, + unsigned char, short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.layers = 1; + params.cubemap = true; + 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; + + for (auto layer = 0u; layer < params.layers; ++layer) { + texCubemapLayeredLodKernel><<>>( + 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, layer, 0.0); + HIP_CHECK(hipGetLastError()); + + fixture.LoadOutput(); + + for (auto k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + auto ref_val = Vec4Map(fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/texCubemapLod.cc b/projects/hip-tests/catch/unit/texture/texCubemapLod.cc new file mode 100644 index 0000000000..0d33048197 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/texCubemapLod.cc @@ -0,0 +1,201 @@ +/* +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 texCubemapLod texCubemapLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLod` 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/texCubemapLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.cubemap = true; + 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; + + texCubemapLodKernel><<>>( + 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 k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `texCubemapLod` 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/texCubemapLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(2, 2, 6); + params.num_subdivisions = 4; + params.cubemap = true; + 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; + + texCubemapLodKernel><<>>( + 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 k = 0u; k < params.NumItersZ(); ++k) { + 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); + float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("j: " << j); + INFO("k: " << k); + 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 index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i; + + auto ref_val = Vec4Map(fixture.tex_h.TexCubemap(x, y, z, 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); + } + } + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/texture_reference.hh b/projects/hip-tests/catch/unit/texture/texture_reference.hh index 21e08abd0e..45f7dd8efa 100644 --- a/projects/hip-tests/catch/unit/texture/texture_reference.hh +++ b/projects/hip-tests/catch/unit/texture/texture_reference.hh @@ -52,6 +52,64 @@ template class TextureReference { } } + TexelType TexCubemap(float x, float y, float z, const hipTextureDesc& tex_desc) const { + x = tex_desc.normalizedCoords ? x * extent_.width : x; + y = tex_desc.normalizedCoords ? y * extent_.height : y; + z = tex_desc.normalizedCoords ? z * extent_.depth : z; + + int face; + float m, s, t; + + if (std::abs(x) > std::abs(y) && std::abs(x) > std::abs(z)) { + if (x >= 0) { + face = 0; + m = x; + s = -z; + t = -y; + } else { + face = 1; + m = -x; + s = z; + t = -y; + } + } else if (std::abs(y) >= std::abs(x) && std::abs(y) > std::abs(z)) { + if (y >= 0) { + face = 2; + m = y; + s = x; + t = z; + } else { + face = 3; + m = -y; + s = x; + t = -z; + } + } else { + if (z >= 0) { + face = 4; + m = z; + s = x; + t = -y; + } else { + face = 5; + m = -z; + s = -x; + t = -y; + } + } + + float coord1 = (s / m + 1) / 2; + float coord2 = (t / m + 1) / 2; + + if (tex_desc.filterMode == hipFilterModePoint) { + return Sample(roundf(coord1), roundf(coord2), face, tex_desc.addressMode); + } else if (tex_desc.filterMode == hipFilterModeLinear) { + return LinearFiltering(coord1, coord2, face, tex_desc.addressMode); + } else { + throw std::invalid_argument("Invalid hipFilterMode value"); + } + } + TexelType Tex1DLayered(float x, int layer, const hipTextureDesc& tex_desc) const { x = tex_desc.normalizedCoords ? x * extent_.width : x; if (tex_desc.filterMode == hipFilterModePoint) { @@ -236,10 +294,17 @@ template class TextureReference { return coord; } + template float FloatToNBitFractional(float x) const { + constexpr size_t mult = 1 << N; + const auto x_trunc = std::trunc(x); + const auto x_frac = std::round((x - x_trunc) * mult) / mult; + return x_trunc + x_frac; + } + std::tuple GetLinearFilteringParams(float coord) const { - const auto coordB = coord - 0.5f; + const auto coordB = FloatToNBitFractional<8>(coord - 0.5f); const auto index = floorf(coordB); - const FixedPoint<8> coeff = coordB - index; + const auto coeff = coordB - index; return {index, coeff}; }