diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index fe4c0a8d1d..73394005b6 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -51,6 +51,13 @@ set(TEST_SRC hipMipmappedArrayCreate.cc hipMipmappedArrayDestroy.cc hipMipmappedArrayGetLevel.cc + tex1Dfetch.cc + tex1D.cc + tex1DLayered.cc + tex1DGrad.cc + tex1DLayeredGrad.cc + tex1DLayeredLod.cc + tex1DLod.cc ) if(WIN32) diff --git a/catch/unit/texture/kernels.hh b/catch/unit/texture/kernels.hh index 056dd3f6e0..ac5d73ff40 100644 --- a/catch/unit/texture/kernels.hh +++ b/catch/unit/texture/kernels.hh @@ -33,6 +33,14 @@ __host__ __device__ inline float GetCoordinate(size_t iteration, size_t N, size_ return normalized_coords ? x / dim : x; } +template +__global__ void tex1DfetchKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj) { + const auto tid = cg::this_grid().thread_rank(); + if (tid >= N) return; + + out[tid] = tex1D(tex_obj, tid); +} + template __global__ void tex1DKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, size_t width, size_t num_subdivisions, bool normalized_coords) { @@ -43,6 +51,50 @@ __global__ void tex1DKernel(TexelType* const out, size_t N, hipTextureObject_t t out[tid] = tex1D(tex_obj, x); } +template +__global__ void tex1DLodKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, + size_t width, size_t num_subdivisions, bool normalized_coords, + float level_of_detail) { + const auto tid = cg::this_grid().thread_rank(); + if (tid >= N) return; + + float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords); + out[tid] = tex1DLod(tex_obj, x, level_of_detail); +} + +template +__global__ void tex1DLayeredLodKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, + size_t width, size_t num_subdivisions, bool normalized_coords, + int layer, float level_of_detail) { + const auto tid = cg::this_grid().thread_rank(); + if (tid >= N) return; + + float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords); + out[tid] = tex1DLayeredLod(tex_obj, x, layer, level_of_detail); +} + +template +__global__ void tex1DGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, + size_t width, size_t num_subdivisions, bool normalized_coords, + float dx, float dy) { + const auto tid = cg::this_grid().thread_rank(); + if (tid >= N) return; + + float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords); + out[tid] = tex1DGrad(tex_obj, x, dx, dy); +} + +template +__global__ void tex1DLayeredGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, + size_t width, size_t num_subdivisions, + bool normalized_coords, float dx, float dy, int layer) { + const auto tid = cg::this_grid().thread_rank(); + if (tid >= N) return; + + float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords); + out[tid] = tex1DLayeredGrad(tex_obj, x, layer, dx, dy); +} + 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, diff --git a/catch/unit/texture/tex1D.cc b/catch/unit/texture/tex1D.cc new file mode 100644 index 0000000000..de92fca2ee --- /dev/null +++ b/catch/unit/texture/tex1D.cc @@ -0,0 +1,140 @@ +/* +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 tex1D tex1D + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1D` 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/tex1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Index: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = fixture.tex_h.Tex1D(x, 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 `tex1D` 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/tex1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, + unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex1D(x, 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/tex1DGrad.cc b/catch/unit/texture/tex1DGrad.cc new file mode 100644 index 0000000000..5b893d0954 --- /dev/null +++ b/catch/unit/texture/tex1DGrad.cc @@ -0,0 +1,140 @@ +/* +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 tex1DGrad tex1DGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1DGrad` 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/tex1DGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Index: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = fixture.tex_h.Tex1D(x, 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 `tex1DGrad` 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/tex1DGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex1D(x, 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/tex1DLayered.cc b/catch/unit/texture/tex1DLayered.cc new file mode 100644 index 0000000000..e0ad4f707c --- /dev/null +++ b/catch/unit/texture/tex1DLayered.cc @@ -0,0 +1,149 @@ +/* +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 tex1DLayered tex1DLayered + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1DLayered` 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/tex1DLayered.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLayered_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.layers = 2; + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + + for (auto layer = 0u; layer < params.layers; ++layer) { + tex1DLayeredKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, layer); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("i: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + const auto ref_val = fixture.tex_h.Tex1DLayered(x, 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 `tex1DLayered` 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/tex1DLayered.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLayered_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.layers = 2; + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + + for (auto layer = 0u; layer < params.layers; ++layer) { + tex1DLayeredKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, layer); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Layer: " << layer); + INFO("Index: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = Vec4Map(fixture.tex_h.Tex1DLayered(x, 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); + } + } +} diff --git a/catch/unit/texture/tex1DLayeredGrad.cc b/catch/unit/texture/tex1DLayeredGrad.cc new file mode 100644 index 0000000000..07c9734619 --- /dev/null +++ b/catch/unit/texture/tex1DLayeredGrad.cc @@ -0,0 +1,140 @@ +/* +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 tex1DLayeredGrad tex1DLayeredGrad + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1DLayeredGrad` 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/tex1DLayeredGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DLayeredGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f, 0); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Index: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = fixture.tex_h.Tex1D(x, 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 `tex1DLayeredGrad` 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/tex1DLayeredGrad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "", char, + unsigned char, short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DLayeredGradKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f, 0); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex1D(x, 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/tex1DLayeredLod.cc b/catch/unit/texture/tex1DLayeredLod.cc new file mode 100644 index 0000000000..874c99c85a --- /dev/null +++ b/catch/unit/texture/tex1DLayeredLod.cc @@ -0,0 +1,140 @@ +/* +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 tex1DLayeredLod tex1DLayeredLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1DLayeredLod` 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/tex1DLayeredLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char, unsigned char, + short, unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DLayeredLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0, 0); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Index: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = fixture.tex_h.Tex1D(x, 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 `tex1DLayeredLod` 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/tex1DLayeredLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DLayeredLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0, 0); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex1D(x, 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/tex1DLod.cc b/catch/unit/texture/tex1DLod.cc new file mode 100644 index 0000000000..ceee1211b3 --- /dev/null +++ b/catch/unit/texture/tex1DLod.cc @@ -0,0 +1,140 @@ +/* +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 tex1DLod tex1DLod + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1DLod` 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/tex1DLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("Index: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = fixture.tex_h.Tex1D(x, 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 `tex1DLod` 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/tex1DLod.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short, + unsigned short) { + CHECK_IMAGE_SUPPORT; + + TextureTestParams params = {}; + params.extent = make_hipExtent(1024, 0, 0); + params.num_subdivisions = 4; + params.GenerateTextureDesc(hipReadModeNormalizedFloat); + + TextureTestFixture fixture{params}; + + const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX()); + tex1DLodKernel><<>>( + fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(), + params.num_subdivisions, params.tex_desc.normalizedCoords, 0); + + fixture.LoadOutput(); + + for (auto i = 0u; i < params.NumItersX(); ++i) { + float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions, + params.tex_desc.normalizedCoords); + + INFO("i: " << i); + INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords); + INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0])); + INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode)); + INFO("x: " << std::fixed << std::setprecision(16) << x); + + auto ref_val = + Vec4Map(fixture.tex_h.Tex1D(x, 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/tex1Dfetch.cc b/catch/unit/texture/tex1Dfetch.cc new file mode 100644 index 0000000000..9354d92a46 --- /dev/null +++ b/catch/unit/texture/tex1Dfetch.cc @@ -0,0 +1,162 @@ +/* +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 +#include + +#include "kernels.hh" +#include "utils.hh" +#include "vec4.hh" + +/** + * @addtogroup tex1D tex1D + * @{ + * @ingroup TextureTest + */ + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1Dfetch` and read mode set to `hipReadModeElementType`. + * Test source + * ------------------------ + * - unit/texture/tex1Dfetch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeElementType", "", char, unsigned char, short, + unsigned short, int, unsigned int, float) { + CHECK_IMAGE_SUPPORT; + + std::vector> tex_h(1024); + for (auto i = 0u; i < tex_h.size(); ++i) { + tex_h[i].x = i + 7; + tex_h[i].y = i + 7; + tex_h[i].z = i + 7; + tex_h[i].w = i + 7; + } + + const auto alloc_size = tex_h.size() * sizeof(vec4); + LinearAllocGuard> tex_alloc_d(LinearAllocs::hipMalloc, alloc_size); + HIP_CHECK(hipMemcpy(tex_alloc_d.ptr(), tex_h.data(), alloc_size, hipMemcpyHostToDevice)); + + hipResourceDesc res_desc; + memset(&res_desc, 0, sizeof(res_desc)); + res_desc.resType = hipResourceTypeLinear; + res_desc.res.linear.devPtr = tex_alloc_d.ptr(); + res_desc.res.linear.desc = hipCreateChannelDesc>(); + res_desc.res.linear.sizeInBytes = alloc_size; + + hipTextureDesc tex_desc; + memset(&tex_desc, 0, sizeof(tex_desc)); + tex_desc.filterMode = hipFilterModePoint; + tex_desc.readMode = hipReadModeElementType; + tex_desc.normalizedCoords = false; + tex_desc.addressMode[0] = hipAddressModeClamp; + + LinearAllocGuard> out_alloc_d(LinearAllocs::hipMalloc, alloc_size); + TextureGuard tex(&res_desc, &tex_desc); + + const auto num_threads = std::min(1024, tex_h.size()); + const auto num_blocks = (tex_h.size() + num_threads - 1) / num_threads; + tex1DfetchKernel> + <<>>(out_alloc_d.ptr(), tex_h.size(), tex.object()); + + std::vector> out_alloc_h(tex_h.size()); + HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), alloc_size, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + for (auto i = 0u; i < out_alloc_h.size(); ++i) { + INFO("Index: " << i); + const auto ref_val = tex_h[i]; + REQUIRE(ref_val.x == out_alloc_h[i].x); + REQUIRE(ref_val.y == out_alloc_h[i].y); + REQUIRE(ref_val.z == out_alloc_h[i].z); + REQUIRE(ref_val.w == out_alloc_h[i].w); + } +} + +/** + * Test Description + * ------------------------ + * - Test texture fetching with `tex1Dfetch` and read mode set to `hipReadModeNormalizedFloat`. + * Test source + * ------------------------ + * - unit/texture/tex1Dfetch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat", "", char, unsigned char, + short, unsigned short) { + CHECK_IMAGE_SUPPORT; + + std::vector> tex_h(1024); + for (auto i = 0u; i < tex_h.size(); ++i) { + tex_h[i].x = i + 7; + tex_h[i].y = i + 7; + tex_h[i].z = i + 7; + tex_h[i].w = i + 7; + } + + const auto alloc_size = tex_h.size() * sizeof(vec4); + LinearAllocGuard> tex_alloc_d(LinearAllocs::hipMalloc, alloc_size); + HIP_CHECK(hipMemcpy(tex_alloc_d.ptr(), tex_h.data(), alloc_size, hipMemcpyHostToDevice)); + + hipResourceDesc res_desc; + memset(&res_desc, 0, sizeof(res_desc)); + res_desc.resType = hipResourceTypeLinear; + res_desc.res.linear.devPtr = tex_alloc_d.ptr(); + res_desc.res.linear.desc = hipCreateChannelDesc>(); + res_desc.res.linear.sizeInBytes = alloc_size; + + hipTextureDesc tex_desc; + memset(&tex_desc, 0, sizeof(tex_desc)); + tex_desc.filterMode = hipFilterModePoint; + tex_desc.readMode = hipReadModeElementType; + tex_desc.normalizedCoords = false; + tex_desc.addressMode[0] = hipAddressModeClamp; + + LinearAllocGuard> out_alloc_d(LinearAllocs::hipMalloc, alloc_size); + TextureGuard tex(&res_desc, &tex_desc); + + const auto num_threads = std::min(1024, tex_h.size()); + const auto num_blocks = (tex_h.size() + num_threads - 1) / num_threads; + tex1DfetchKernel> + <<>>(out_alloc_d.ptr(), tex_h.size(), tex.object()); + + std::vector> out_alloc_h(tex_h.size()); + HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), alloc_size, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + for (auto i = 0u; i < out_alloc_h.size(); ++i) { + INFO("Index: " << i); + const auto ref_val = Vec4Map(tex_h[i], NormalizeInteger); + REQUIRE(ref_val.x == out_alloc_h[i].x); + REQUIRE(ref_val.y == out_alloc_h[i].y); + REQUIRE(ref_val.z == out_alloc_h[i].z); + REQUIRE(ref_val.w == out_alloc_h[i].w); + } +} \ No newline at end of file