EXSWHTEC-322 - Implement common utilities for texture fetching tests (#365)
Change-Id: I7cdf56008dce2551706bc5c0ab1a49133a6337e9
This commit is contained in:
Mirza Halilcevic
committato da
Rakesh Roy
Rakesh Roy
parent
2f783afe8c
commit
16128f659a
@@ -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 <cmath>
|
||||
|
||||
template <size_t fractional_bits> class FixedPoint {
|
||||
public:
|
||||
FixedPoint() = default;
|
||||
FixedPoint(float x) { fixed_point_ = static_cast<int16_t>(roundf(x * (1 << fractional_bits))); }
|
||||
|
||||
operator float() const {
|
||||
return (static_cast<float>(fixed_point_) / static_cast<float>(1 << fractional_bits));
|
||||
}
|
||||
|
||||
FixedPoint operator+(FixedPoint other) const {
|
||||
FixedPoint<fractional_bits> res;
|
||||
res.fixed_point_ = fixed_point_ + other.fixed_point_;
|
||||
return res;
|
||||
}
|
||||
|
||||
FixedPoint operator-(FixedPoint other) const {
|
||||
FixedPoint<fractional_bits> res;
|
||||
res.fixed_point_ = fixed_point_ - other.fixed_point_;
|
||||
return res;
|
||||
}
|
||||
|
||||
FixedPoint operator*(FixedPoint other) const {
|
||||
constexpr auto K = 1 << (fractional_bits - 1);
|
||||
|
||||
FixedPoint<fractional_bits> res;
|
||||
int32_t temp;
|
||||
|
||||
temp = static_cast<int32_t>(fixed_point_) * static_cast<int32_t>(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;
|
||||
}
|
||||
};
|
||||
@@ -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 <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
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<float>(iteration) - N / 2) / num_subdivisions;
|
||||
return normalized_coords ? x / dim : x;
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__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<TexelType>(tex_obj, x);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__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<TexelType>(tex_obj, x, y);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__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<TexelType>(tex_obj, x, y, z);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__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<TexelType>(tex_obj, x, layer);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__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<TexelType>(tex_obj, x, y, layer);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
#include "texture_reference.hh"
|
||||
#include "utils.hh"
|
||||
#include "vec4.hh"
|
||||
|
||||
template <typename TestType> 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<TestType>;
|
||||
|
||||
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 <typename TestType, bool normalized_read = false> struct TextureTestFixture {
|
||||
using VecType = vec4<TestType>;
|
||||
using OutType = std::conditional_t<normalized_read, vec4<float>, VecType>;
|
||||
|
||||
TextureTestParams<TestType> params;
|
||||
hipResourceDesc res_desc;
|
||||
|
||||
LinearAllocGuard<VecType> host_alloc;
|
||||
TextureReference<VecType> tex_h;
|
||||
ArrayAllocGuard<VecType> tex_alloc_d;
|
||||
TextureGuard tex;
|
||||
LinearAllocGuard<OutType> out_alloc_d;
|
||||
std::vector<OutType> out_alloc_h;
|
||||
|
||||
TextureTestFixture(const TextureTestParams<TestType>& 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<TestType>(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());
|
||||
}
|
||||
};
|
||||
@@ -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 <cmath>
|
||||
|
||||
#include "fixed_point.hh"
|
||||
|
||||
template <typename TexelType> 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 <typename T> TexelType Vec4Sum(T arg) const { return Vec4Add(arg, Zero()); }
|
||||
|
||||
template <typename T, typename... Ts> 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<float>::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<size_t>(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<size_t>(y) * extent_.width + static_cast<size_t>(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<size_t>(z) * extent_.width * extent_.height +
|
||||
static_cast<size_t>(y) * extent_.width + static_cast<size_t>(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<float>(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<ssize_t>(floorf(coord)) % 2;
|
||||
coord = is_reversing ? 1.0f - frac : frac;
|
||||
coord *= dim;
|
||||
coord -= (coord == truncf(coord)) * is_reversing;
|
||||
|
||||
return coord;
|
||||
}
|
||||
|
||||
template <size_t N> 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<float, float> 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};
|
||||
}
|
||||
};
|
||||
@@ -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 <algorithm>
|
||||
#include <cmath>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
class TextureGuard {
|
||||
public:
|
||||
TextureGuard(hipResourceDesc* res_desc, hipTextureDesc* tex_desc) {
|
||||
HIP_CHECK(hipCreateTextureObject(&tex_obj_, res_desc, tex_desc, nullptr));
|
||||
}
|
||||
|
||||
~TextureGuard() { static_cast<void>(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 <typename T> std::enable_if_t<std::is_integral_v<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<float>(x));
|
||||
xf = std::min<float>(xf, std::numeric_limits<T>::max());
|
||||
return std::copysign(xf / std::numeric_limits<T>::max(), x);
|
||||
}
|
||||
|
||||
inline std::tuple<size_t, size_t> GetLaunchConfig(size_t max_num_threads, size_t num_iters) {
|
||||
auto num_threads = std::min<size_t>(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");
|
||||
}
|
||||
}
|
||||
@@ -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 <hip/hip_runtime_api.h>
|
||||
|
||||
template <typename T> struct vec4_struct { using type = void; };
|
||||
|
||||
#define DEFINE_VEC4_OVERLOAD(base_type, vec_type) \
|
||||
template <> struct vec4_struct<base_type> { 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 <typename T> using vec4 = typename vec4_struct<T>::type;
|
||||
|
||||
template <typename T> inline void SetVec4(vec4<T>& vec, const T val) {
|
||||
vec.x = val;
|
||||
vec.y = val;
|
||||
vec.z = val;
|
||||
vec.w = val;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline void SetVec4(vec4<T>& 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 <typename T> inline auto MakeVec4(const T val) {
|
||||
vec4<T> vec;
|
||||
SetVec4(vec, val);
|
||||
|
||||
return vec;
|
||||
}
|
||||
|
||||
template <typename T> inline void MakeVec4(const T x, const T y, const T z, const T w) {
|
||||
vec4<T> vec;
|
||||
SetVec4(vec, x, y, z, w);
|
||||
|
||||
return vec;
|
||||
}
|
||||
|
||||
template <typename T, typename F> inline auto Vec4Map(const vec4<T>& vec, F f) {
|
||||
vec4<decltype(f(vec.x))> ret;
|
||||
ret.x = f(vec.x);
|
||||
ret.y = f(vec.y);
|
||||
ret.z = f(vec.z);
|
||||
ret.w = f(vec.w);
|
||||
|
||||
return ret;
|
||||
}
|
||||
|
||||
template <typename T> 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 <typename T> 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;
|
||||
}
|
||||
Fai riferimento in un nuovo problema
Block a user