diff --git a/catch/unit/texture/fixed_point.hh b/catch/unit/texture/fixed_point.hh new file mode 100644 index 0000000000..ba95c42c44 --- /dev/null +++ b/catch/unit/texture/fixed_point.hh @@ -0,0 +1,73 @@ +/* +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. +*/ + +#pragma once + +#include + +template class FixedPoint { + public: + FixedPoint() = default; + FixedPoint(float x) { fixed_point_ = static_cast(roundf(x * (1 << fractional_bits))); } + + operator float() const { + return (static_cast(fixed_point_) / static_cast(1 << fractional_bits)); + } + + FixedPoint operator+(FixedPoint other) const { + FixedPoint res; + res.fixed_point_ = fixed_point_ + other.fixed_point_; + return res; + } + + FixedPoint operator-(FixedPoint other) const { + FixedPoint res; + res.fixed_point_ = fixed_point_ - other.fixed_point_; + return res; + } + + FixedPoint operator*(FixedPoint other) const { + constexpr auto K = 1 << (fractional_bits - 1); + + FixedPoint res; + int32_t temp; + + temp = static_cast(fixed_point_) * static_cast(other.fixed_point_); + temp += K; + + res.fixed_point_ = Sat16(temp >> fractional_bits); + + return res; + } + + private: + int16_t fixed_point_; + + int16_t Sat16(int32_t x) const { + if (x > 0x7FFF) + return 0x7FFF; + else if (x < -0x8000) + return -0x8000; + else + return (int16_t)x; + } +}; \ No newline at end of file diff --git a/catch/unit/texture/kernels.hh b/catch/unit/texture/kernels.hh new file mode 100644 index 0000000000..056dd3f6e0 --- /dev/null +++ b/catch/unit/texture/kernels.hh @@ -0,0 +1,107 @@ +/* +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. +*/ + +#pragma once + +#include +#include + +namespace cg = cooperative_groups; + +__host__ __device__ inline float GetCoordinate(size_t iteration, size_t N, size_t dim, + size_t num_subdivisions, bool normalized_coords) { + float x = (static_cast(iteration) - N / 2) / num_subdivisions; + return normalized_coords ? x / dim : x; +} + +template +__global__ void tex1DKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, + size_t width, size_t num_subdivisions, bool normalized_coords) { + const auto tid = cg::this_grid().thread_rank(); + if (tid >= N) return; + + float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords); + out[tid] = tex1D(tex_obj, x); +} + +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, + size_t num_subdivisions, bool normalized_coords) { + const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x; + if (tid_x >= N_x) return; + + const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y; + if (tid_y >= N_y) return; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2D(tex_obj, x, y); +} + +template +__global__ void tex3DKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z, + hipTextureObject_t tex_obj, size_t width, size_t height, size_t depth, + 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] = tex3D(tex_obj, x, y, z); +} + +template +__global__ void tex1DLayeredKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj, + size_t width, size_t num_subdivisions, bool normalized_coords, + size_t 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] = tex1DLayered(tex_obj, x, layer); +} + +template +__global__ void tex2DLayeredKernel(TexelType* const out, size_t N_x, size_t N_y, + hipTextureObject_t tex_obj, size_t width, size_t height, + size_t num_subdivisions, bool normalized_coords, 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; + + float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords); + float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords); + + out[tid_y * N_x + tid_x] = tex2DLayered(tex_obj, x, y, layer); +} \ No newline at end of file diff --git a/catch/unit/texture/test_fixture.hh b/catch/unit/texture/test_fixture.hh new file mode 100644 index 0000000000..47ac8b73bc --- /dev/null +++ b/catch/unit/texture/test_fixture.hh @@ -0,0 +1,150 @@ +/* +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. +*/ + +#pragma once + +#include +#include + +#include "texture_reference.hh" +#include "utils.hh" +#include "vec4.hh" + +template struct TextureTestParams { + hipExtent extent; + size_t layers; + size_t num_subdivisions; + hipTextureDesc tex_desc; + + size_t Size() const { + return extent.width * (extent.height ?: 1) * (extent.depth ?: 1) * (layers ?: 1); + } + + size_t NumItersX() const { return 3 * extent.width * num_subdivisions * 2 + 1; } + + size_t NumItersY() const { return 3 * extent.height * num_subdivisions * 2 + 1; } + + size_t NumItersZ() const { return 3 * extent.depth * num_subdivisions * 2 + 1; } + + size_t NumIters() const { return NumItersX() * NumItersY() * NumItersZ(); } + + size_t Width() const { return extent.width; } + + size_t Height() const { return extent.height; } + + size_t Depth() const { return extent.depth; } + + hipExtent LayeredExtent() const { + return Layered() ? make_hipExtent(Width(), Height(), layers) : extent; + } + + bool Layered() const { return layers > 1; } + + void GenerateTextureDesc(decltype(hipReadModeElementType) read_mode = hipReadModeElementType) { + constexpr bool is_floating_point = std::is_floating_point_v; + + memset(&tex_desc, 0, sizeof(tex_desc)); + tex_desc.readMode = read_mode; + + tex_desc.filterMode = hipFilterModePoint; + if (is_floating_point || tex_desc.readMode == hipReadModeNormalizedFloat) { + tex_desc.filterMode = GENERATE(hipFilterModePoint, hipFilterModeLinear); + } + + tex_desc.normalizedCoords = GENERATE(false, true); + + auto address_mode_x = hipAddressModeClamp; + auto address_mode_y = address_mode_x; + auto address_mode_z = address_mode_y; + + if (tex_desc.normalizedCoords) { + address_mode_x = GENERATE(hipAddressModeClamp, hipAddressModeBorder, hipAddressModeWrap, + hipAddressModeMirror); + if (extent.height) + address_mode_y = GENERATE(hipAddressModeClamp, hipAddressModeBorder, hipAddressModeWrap, + hipAddressModeMirror); + if (extent.depth) + address_mode_z = GENERATE(hipAddressModeClamp, hipAddressModeBorder, hipAddressModeWrap, + hipAddressModeMirror); + } else { + address_mode_x = GENERATE(hipAddressModeClamp, hipAddressModeBorder); + if (extent.height) address_mode_y = GENERATE(hipAddressModeClamp, hipAddressModeBorder); + if (extent.depth) address_mode_z = GENERATE(hipAddressModeClamp, hipAddressModeBorder); + } + + tex_desc.addressMode[0] = address_mode_x; + if (extent.height) tex_desc.addressMode[1] = address_mode_y; + if (extent.depth) tex_desc.addressMode[2] = address_mode_z; + } +}; + +template struct TextureTestFixture { + using VecType = vec4; + using OutType = std::conditional_t, VecType>; + + TextureTestParams params; + hipResourceDesc res_desc; + + LinearAllocGuard host_alloc; + TextureReference tex_h; + ArrayAllocGuard tex_alloc_d; + TextureGuard tex; + LinearAllocGuard out_alloc_d; + std::vector out_alloc_h; + + TextureTestFixture(const TextureTestParams& p) + : 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{ResDesc(), ¶ms.tex_desc}, + out_alloc_d{LinearAllocs::hipMalloc, sizeof(OutType) * params.NumIters()}, + out_alloc_h(params.NumIters()) {} + + hipResourceDesc* ResDesc() { + constexpr int test_value_offset = 7; + for (auto i = 0u; i < params.Size(); ++i) { + SetVec4(host_alloc.ptr()[i], i + test_value_offset); + } + + hipMemcpy3DParms memcpy_params = {}; + memcpy_params.dstArray = tex_alloc_d.ptr(); + memcpy_params.extent = params.LayeredExtent(); + memcpy_params.extent.height = memcpy_params.extent.height ?: 1; + memcpy_params.extent.depth = memcpy_params.extent.depth ?: 1; + memcpy_params.srcPtr = make_hipPitchedPtr(tex_h.ptr(0), sizeof(VecType) * params.Width(), + params.Width(), params.Height() ?: 1); + memcpy_params.kind = hipMemcpyHostToDevice; + HIP_CHECK(hipMemcpy3D(&memcpy_params)); + + memset(&res_desc, 0, sizeof(res_desc)); + res_desc.resType = hipResourceTypeArray; + res_desc.res.array.array = tex_alloc_d.ptr(); + return &res_desc; + } + + void LoadOutput() { + HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), sizeof(OutType) * params.NumIters(), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + } +}; \ No newline at end of file diff --git a/catch/unit/texture/texture_reference.hh b/catch/unit/texture/texture_reference.hh new file mode 100644 index 0000000000..d2ee9159e0 --- /dev/null +++ b/catch/unit/texture/texture_reference.hh @@ -0,0 +1,253 @@ +/* +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. +*/ + +#pragma once + +#include + +#include "fixed_point.hh" + +template class TextureReference { + public: + TextureReference(TexelType* alloc, hipExtent extent, size_t layers) + : alloc_{alloc}, extent_{extent}, layers_{layers} {} + + TexelType Tex1D(float x, const hipTextureDesc& tex_desc) const { + return Tex1DLayered(x, 0, tex_desc); + } + + TexelType Tex2D(float x, float y, const hipTextureDesc& tex_desc) const { + return Tex2DLayered(x, y, 0, tex_desc); + } + + TexelType Tex3D(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; + if (tex_desc.filterMode == hipFilterModePoint) { + return Sample(floorf(x), floorf(y), floorf(z), tex_desc.addressMode); + } else if (tex_desc.filterMode == hipFilterModeLinear) { + return LinearFiltering(x, y, z, 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) { + return Sample(floorf(x), layer, tex_desc.addressMode); + } else if (tex_desc.filterMode == hipFilterModeLinear) { + return LinearFiltering(x, layer, tex_desc.addressMode); + } else { + throw std::invalid_argument("Invalid hipFilterMode value"); + } + } + + TexelType Tex2DLayered(float x, float y, int layer, const hipTextureDesc& tex_desc) const { + x = tex_desc.normalizedCoords ? x * extent_.width : x; + y = tex_desc.normalizedCoords ? y * extent_.height : y; + if (tex_desc.filterMode == hipFilterModePoint) { + return Sample(floorf(x), floorf(y), layer, tex_desc.addressMode); + } else if (tex_desc.filterMode == hipFilterModeLinear) { + return LinearFiltering(x, y, layer, tex_desc.addressMode); + } else { + throw std::invalid_argument("Invalid hipFilterMode value"); + } + } + + TexelType* ptr(size_t layer) const { + return alloc_ + layer * extent_.width * (extent_.height ?: 1); + } + + size_t width() const { return extent_.width; } + + size_t height() const { return extent_.height; } + + size_t depth() const { return extent_.depth; } + + private: + TexelType* const alloc_; + const hipExtent extent_; + const size_t layers_; + + template TexelType Vec4Sum(T arg) const { return Vec4Add(arg, Zero()); } + + template TexelType Vec4Sum(T arg, Ts... args) const { + return Vec4Add(arg, Vec4Sum(args...)); + } + + TexelType Zero() const { + TexelType ret; + memset(&ret, 0, sizeof(ret)); + return ret; + } + + float ApplyAddressMode(float coord, size_t dim, hipTextureAddressMode address_mode) const { + switch (address_mode) { + case hipAddressModeClamp: + return ApplyClamp(coord, dim); + case hipAddressModeBorder: + if (CheckBorder(coord, dim)) { + return std::numeric_limits::quiet_NaN(); + } + case hipAddressModeWrap: + return ApplyWrap(coord, dim); + case hipAddressModeMirror: + return ApplyMirror(coord, dim); + default: + throw std::invalid_argument("Invalid hipAddressMode value"); + } + } + + TexelType Sample(float x, int layer, const hipTextureAddressMode* address_mode) const { + x = ApplyAddressMode(x, extent_.width, address_mode[0]); + + if (std::isnan(x)) { + return Zero(); + } + + return ptr(layer)[static_cast(x)]; + } + + TexelType Sample(float x, float y, int layer, const hipTextureAddressMode* address_mode) const { + x = ApplyAddressMode(x, extent_.width, address_mode[0]); + y = ApplyAddressMode(y, extent_.height, address_mode[1]); + + if (std::isnan(x) || std::isnan(y)) { + return Zero(); + } + + return ptr(layer)[static_cast(y) * extent_.width + static_cast(x)]; + } + + TexelType Sample(float x, float y, float z, const hipTextureAddressMode* address_mode) const { + x = ApplyAddressMode(x, extent_.width, address_mode[0]); + y = ApplyAddressMode(y, extent_.height, address_mode[1]); + z = ApplyAddressMode(z, extent_.depth, address_mode[2]); + + if (std::isnan(x) || std::isnan(y) || std::isnan(z)) { + return Zero(); + } + + return ptr(0)[static_cast(z) * extent_.width * extent_.height + + static_cast(y) * extent_.width + static_cast(x)]; + } + + TexelType LinearFiltering(float x, int layer, const hipTextureAddressMode* address_mode) const { + const auto [i, alpha] = GetLinearFilteringParams(x); + + const auto T_i0 = Sample(i, layer, address_mode); + const auto T_i1 = Sample(i + 1.0f, layer, address_mode); + + const auto term_i0 = Vec4Scale((1.0f - alpha), T_i0); + const auto term_i1 = Vec4Scale(alpha, T_i1); + + return Vec4Sum(term_i0, term_i1); + } + + TexelType LinearFiltering(float x, float y, int layer, + const hipTextureAddressMode* address_mode) const { + const auto [i, alpha] = GetLinearFilteringParams(x); + const auto [j, beta] = GetLinearFilteringParams(y); + + const auto T_i0j0 = Sample(i, j, layer, address_mode); + const auto T_i1j0 = Sample(i + 1.0f, j, layer, address_mode); + const auto T_i0j1 = Sample(i, j + 1.0f, layer, address_mode); + const auto T_i1j1 = Sample(i + 1.0f, j + 1.0f, layer, address_mode); + + const auto term_i0j0 = Vec4Scale((1.0f - alpha) * (1.0f - beta), T_i0j0); + const auto term_i1j0 = Vec4Scale(alpha * (1.0f - beta), T_i1j0); + const auto term_i0j1 = Vec4Scale((1.0f - alpha) * beta, T_i0j1); + const auto term_i1j1 = Vec4Scale(alpha * beta, T_i1j1); + + return Vec4Sum(term_i0j0, term_i1j0, term_i0j1, term_i1j1); + } + + TexelType LinearFiltering(float x, float y, float z, + const hipTextureAddressMode* address_mode) const { + const auto [i, alpha] = GetLinearFilteringParams(x); + const auto [j, beta] = GetLinearFilteringParams(y); + const auto [k, gamma] = GetLinearFilteringParams(z); + + const auto T_i0j0k0 = Sample(i, j, k, address_mode); + const auto T_i1j0k0 = Sample(i + 1.0f, j, k, address_mode); + const auto T_i0j1k0 = Sample(i, j + 1.0f, k, address_mode); + const auto T_i1j1k0 = Sample(i + 1.0f, j + 1.0f, k, address_mode); + const auto T_i0j0k1 = Sample(i, j, k + 1.0f, address_mode); + const auto T_i1j0k1 = Sample(i + 1.0f, j, k + 1.0f, address_mode); + const auto T_i0j1k1 = Sample(i, j + 1.0f, k + 1.0f, address_mode); + const auto T_i1j1k1 = Sample(i + 1.0f, j + 1.0f, k + 1.0f, address_mode); + + const auto term_i0j0k0 = Vec4Scale((1.0f - alpha) * (1.0f - beta) * (1.0f - gamma), T_i0j0k0); + const auto term_i1j0k0 = Vec4Scale(alpha * (1.0f - beta) * (1.0f - gamma), T_i1j0k0); + const auto term_i0j1k0 = Vec4Scale((1.0f - alpha) * beta * (1.0f - gamma), T_i0j1k0); + const auto term_i1j1k0 = Vec4Scale(alpha * beta * (1.0f - gamma), T_i1j1k0); + const auto term_i0j0k1 = Vec4Scale((1.0f - alpha) * (1.0f - beta) * gamma, T_i0j0k1); + const auto term_i1j0k1 = Vec4Scale(alpha * (1.0f - beta) * gamma, T_i1j0k1); + const auto term_i0j1k1 = Vec4Scale((1.0f - alpha) * beta * gamma, T_i0j1k1); + const auto term_i1j1k1 = Vec4Scale(alpha * beta * gamma, T_i1j1k1); + + return Vec4Sum(term_i0j0k0, term_i1j0k0, term_i0j1k0, term_i1j1k0, term_i0j0k1, term_i1j0k1, + term_i0j1k1, term_i1j1k1); + } + + float ApplyClamp(float coord, size_t dim) const { + return max(min(coord, static_cast(dim - 1)), 0.0f); + } + + bool CheckBorder(float coord, size_t dim) const { return coord > dim - 1 || coord < 0.0f; } + + float ApplyWrap(float coord, size_t dim) const { + coord /= dim; + coord = coord - floorf(coord); + coord *= dim; + + return coord; + } + + float ApplyMirror(float coord, size_t dim) const { + coord /= dim; + const float frac = coord - floor(coord); + const bool is_reversing = static_cast(floorf(coord)) % 2; + coord = is_reversing ? 1.0f - frac : frac; + coord *= dim; + coord -= (coord == truncf(coord)) * is_reversing; + + 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 = FloatToNBitFractional<8>(coord - 0.5f); + const auto index = floorf(coordB); + const auto coeff = coordB - index; + + return {index, coeff}; + } +}; diff --git a/catch/unit/texture/utils.hh b/catch/unit/texture/utils.hh new file mode 100644 index 0000000000..90e51d48c8 --- /dev/null +++ b/catch/unit/texture/utils.hh @@ -0,0 +1,85 @@ +/* +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. +*/ + +#pragma once + +#include +#include + +#include + +class TextureGuard { + public: + TextureGuard(hipResourceDesc* res_desc, hipTextureDesc* tex_desc) { + HIP_CHECK(hipCreateTextureObject(&tex_obj_, res_desc, tex_desc, nullptr)); + } + + ~TextureGuard() { static_cast(hipDestroyTextureObject(tex_obj_)); } + + TextureGuard(const TextureGuard&) = delete; + TextureGuard& operator=(const TextureGuard&) = delete; + + hipTextureObject_t object() const { return tex_obj_; } + + private: + hipTextureObject_t tex_obj_ = 0; +}; + +template std::enable_if_t, float> NormalizeInteger(const T x) { + // On the GPU, -1.0 will be returned both for the minimum value of a signed type and its + // successor e.g. for char, -1.0 will be returned for both -128 and -127. + auto xf = std::abs(static_cast(x)); + xf = std::min(xf, std::numeric_limits::max()); + return std::copysign(xf / std::numeric_limits::max(), x); +} + +inline std::tuple GetLaunchConfig(size_t max_num_threads, size_t num_iters) { + auto num_threads = std::min(max_num_threads, num_iters); + auto num_blocks = (num_iters + num_threads - 1) / num_threads; + return {num_threads, num_blocks}; +} + +inline std::string AddressModeToString(decltype(hipAddressModeClamp) address_mode) { + switch (address_mode) { + case hipAddressModeClamp: + return "hipAddressModeClamp"; + case hipAddressModeBorder: + return "hipAddressModeBorder"; + case hipAddressModeWrap: + return "hipAddressModeWrap"; + case hipAddressModeMirror: + return "hipAddressModeMirror"; + default: + throw std::invalid_argument("Invalid hipAddressMode value"); + } +} + +inline std::string FilteringModeToString(decltype(hipFilterModePoint) filter_mode) { + switch (filter_mode) { + case hipFilterModePoint: + return "hipFilterModePoint"; + case hipFilterModeLinear: + return "hipFilterModeLinear"; + default: + throw std::invalid_argument("Invalid hipFilterMode value"); + } +} \ No newline at end of file diff --git a/catch/unit/texture/vec4.hh b/catch/unit/texture/vec4.hh new file mode 100644 index 0000000000..db4b073908 --- /dev/null +++ b/catch/unit/texture/vec4.hh @@ -0,0 +1,106 @@ +/* +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. +*/ + +#pragma once + +#include + +template struct vec4_struct { using type = void; }; + +#define DEFINE_VEC4_OVERLOAD(base_type, vec_type) \ + template <> struct vec4_struct { using type = vec_type; } + +DEFINE_VEC4_OVERLOAD(char, char4); +DEFINE_VEC4_OVERLOAD(short, short4); +DEFINE_VEC4_OVERLOAD(int, int4); +DEFINE_VEC4_OVERLOAD(long, long4); +DEFINE_VEC4_OVERLOAD(long long, longlong4); + +DEFINE_VEC4_OVERLOAD(unsigned char, uchar4); +DEFINE_VEC4_OVERLOAD(unsigned short, ushort4); +DEFINE_VEC4_OVERLOAD(unsigned int, uint4); +DEFINE_VEC4_OVERLOAD(unsigned long, ulong4); +DEFINE_VEC4_OVERLOAD(unsigned long long, ulonglong4); + +DEFINE_VEC4_OVERLOAD(float, float4); +DEFINE_VEC4_OVERLOAD(double, float4); + +template using vec4 = typename vec4_struct::type; + +template inline void SetVec4(vec4& vec, const T val) { + vec.x = val; + vec.y = val; + vec.z = val; + vec.w = val; +} + +template +inline void SetVec4(vec4& vec, const T x, const T y, const T z, const T w) { + vec.x = x; + vec.y = y; + vec.z = z; + vec.w = w; +} + +template inline auto MakeVec4(const T val) { + vec4 vec; + SetVec4(vec, val); + + return vec; +} + +template inline void MakeVec4(const T x, const T y, const T z, const T w) { + vec4 vec; + SetVec4(vec, x, y, z, w); + + return vec; +} + +template inline auto Vec4Map(const vec4& vec, F f) { + vec4 ret; + ret.x = f(vec.x); + ret.y = f(vec.y); + ret.z = f(vec.z); + ret.w = f(vec.w); + + return ret; +} + +template inline __host__ __device__ auto Vec4Scale(float s, const T& vec) { + T ret; + ret.x = s * vec.x; + ret.y = s * vec.y; + ret.z = s * vec.z; + ret.w = s * vec.w; + + return ret; +} + +template inline __host__ __device__ auto Vec4Add(const T& vec1, const T& vec2) { + T ret; + ret.x = vec1.x + vec2.x; + ret.y = vec1.y + vec2.y; + ret.z = vec1.z + vec2.z; + ret.w = vec1.w + vec2.w; + + return ret; +} \ No newline at end of file