EXSWHTEC-320 - Implement tests for cubemap texture device functions #369
Change-Id: I1a247dba4e46ed7e1045dee0467d3fdac1f70cd0
[ROCm/hip-tests commit: 6eec9aa90d]
This commit is contained in:
Mirza Halilčević
committato da
Rakesh Roy
Rakesh Roy
parent
8d0525beab
commit
ca5e764189
@@ -211,15 +211,8 @@
|
||||
"Unit_hipHostMalloc_AllocateUseMoreThanAvailGPUMemory",
|
||||
"=== SWDEV-432250:Below tests failed in stress test on 10/11/23 ===",
|
||||
"Unit_hipVectorTypes_test_on_device",
|
||||
<<<<<<< HEAD
|
||||
=======
|
||||
"Unit_Layered1DTexture_Check_DeviceBufferToFromLayered1DArray - ushort4",
|
||||
"Unit_Layered2DTexture_Check_DeviceBufferToFromLayered2DArray - float4",
|
||||
"Unit_tex3DLod_Positive_ReadModeElementType",
|
||||
"Unit_tex3DLod_Positive_ReadModeNormalizedFloat",
|
||||
"Unit_tex3DGrad_Positive_ReadModeElementType",
|
||||
"Unit_tex3DGrad_Positive_ReadModeNormalizedFloat",
|
||||
>>>>>>> 5cdc6efc (Merge branch 'develop' into tex3D_tests)
|
||||
"=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
|
||||
"Unit_hiprtc_stdheaders",
|
||||
"NOTE: The following test is disabled due to defect - EXSWHTEC-241",
|
||||
|
||||
@@ -63,6 +63,12 @@ set(TEST_SRC
|
||||
tex3D.cc
|
||||
tex3DLod.cc
|
||||
tex3DGrad.cc
|
||||
texCubemap.cc
|
||||
texCubemapLod.cc
|
||||
texCubemapGrad.cc
|
||||
texCubemapLayered.cc
|
||||
texCubemapLayeredLod.cc
|
||||
texCubemapLayeredGrad.cc
|
||||
)
|
||||
|
||||
if(WIN32)
|
||||
|
||||
@@ -173,6 +173,69 @@ __global__ void tex3DGradKernel(TexelType* const out, size_t N_x, size_t N_y, si
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3DGrad<TexelType>(tex_obj, x, y, z, dx, dy);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void texCubemapKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
|
||||
hipTextureObject_t tex_obj, size_t width, size_t height,
|
||||
size_t depth, size_t num_subdivisions, bool normalized_coords) {
|
||||
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid_x >= N_x) return;
|
||||
|
||||
const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (tid_y >= N_y) return;
|
||||
|
||||
const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (tid_z >= N_z) return;
|
||||
|
||||
float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords);
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = texCubemap<TexelType>(tex_obj, x, y, z);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void texCubemapLodKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
|
||||
hipTextureObject_t tex_obj, size_t width, size_t height,
|
||||
size_t depth, size_t num_subdivisions, bool normalized_coords,
|
||||
float level) {
|
||||
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid_x >= N_x) return;
|
||||
|
||||
const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (tid_y >= N_y) return;
|
||||
|
||||
const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (tid_z >= N_z) return;
|
||||
|
||||
float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords);
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = texCubemapLod<TexelType>(tex_obj, x, y, z, level);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void texCubemapGradKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
|
||||
hipTextureObject_t tex_obj, size_t width, size_t height,
|
||||
size_t depth, size_t num_subdivisions, bool normalized_coords,
|
||||
float4 dx, float4 dy) {
|
||||
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid_x >= N_x) return;
|
||||
|
||||
const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (tid_y >= N_y) return;
|
||||
|
||||
const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (tid_z >= N_z) return;
|
||||
|
||||
float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords);
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
|
||||
texCubemapGrad<TexelType>(tex_obj, x, y, z, dx, dy);
|
||||
}
|
||||
|
||||
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,
|
||||
@@ -198,4 +261,71 @@ __global__ void tex2DLayeredKernel(TexelType* const out, size_t N_x, size_t N_y,
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_y * N_x + tid_x] = tex2DLayered<TexelType>(tex_obj, x, y, layer);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void texCubemapLayeredKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
|
||||
hipTextureObject_t tex_obj, size_t width, size_t height,
|
||||
size_t depth, size_t num_subdivisions,
|
||||
bool normalized_coords, size_t layer) {
|
||||
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid_x >= N_x) return;
|
||||
|
||||
const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (tid_y >= N_y) return;
|
||||
|
||||
const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (tid_z >= N_z) return;
|
||||
|
||||
float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords);
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
|
||||
texCubemapLayered<TexelType>(tex_obj, x, y, z, layer);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void texCubemapLayeredLodKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
|
||||
hipTextureObject_t tex_obj, size_t width, size_t height,
|
||||
size_t depth, size_t num_subdivisions,
|
||||
bool normalized_coords, size_t layer, float level) {
|
||||
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid_x >= N_x) return;
|
||||
|
||||
const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (tid_y >= N_y) return;
|
||||
|
||||
const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (tid_z >= N_z) return;
|
||||
|
||||
float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords);
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
|
||||
texCubemapLayeredLod<TexelType>(tex_obj, x, y, z, layer, level);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void texCubemapLayeredGradKernel(TexelType* const out, size_t N_x, size_t N_y,
|
||||
size_t N_z, hipTextureObject_t tex_obj, size_t width,
|
||||
size_t height, size_t depth, size_t num_subdivisions,
|
||||
bool normalized_coords, size_t layer, float4 dx,
|
||||
float4 dy) {
|
||||
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid_x >= N_x) return;
|
||||
|
||||
const auto tid_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (tid_y >= N_y) return;
|
||||
|
||||
const auto tid_z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (tid_z >= N_z) return;
|
||||
|
||||
float x = GetCoordinate(tid_x, N_x, width, num_subdivisions, normalized_coords);
|
||||
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
|
||||
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
|
||||
|
||||
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
|
||||
texCubemapLayeredGrad<TexelType>(tex_obj, x, y, z, layer, dx, dy);
|
||||
}
|
||||
@@ -34,6 +34,7 @@ template <typename TestType> struct TextureTestParams {
|
||||
size_t layers;
|
||||
size_t num_subdivisions;
|
||||
hipTextureDesc tex_desc;
|
||||
bool cubemap;
|
||||
|
||||
size_t Size() const {
|
||||
return extent.width * (extent.height ?: 1) * (extent.depth ?: 1) * (layers ?: 1);
|
||||
@@ -53,6 +54,10 @@ template <typename TestType> struct TextureTestParams {
|
||||
|
||||
size_t Depth() const { return extent.depth; }
|
||||
|
||||
unsigned int Flags() const {
|
||||
return (Layered() ? hipArrayLayered : 0u) | (cubemap ? hipArrayCubemap : 0u);
|
||||
}
|
||||
|
||||
hipExtent LayeredExtent() const {
|
||||
return Layered() ? make_hipExtent(Width(), Height(), layers) : extent;
|
||||
}
|
||||
@@ -115,7 +120,7 @@ template <typename TestType, bool normalized_read = false> struct TextureTestFix
|
||||
: params{p},
|
||||
host_alloc{LinearAllocs::hipHostMalloc, sizeof(VecType) * params.Size()},
|
||||
tex_h{host_alloc.ptr(), params.extent, params.layers},
|
||||
tex_alloc_d{params.LayeredExtent(), params.Layered() ? hipArrayLayered : 0u},
|
||||
tex_alloc_d{params.LayeredExtent(), params.Flags()},
|
||||
tex{ResDesc(), ¶ms.tex_desc},
|
||||
out_alloc_d{LinearAllocs::hipMalloc, sizeof(OutType) * params.NumIters()},
|
||||
out_alloc_h(params.NumIters()) {}
|
||||
|
||||
@@ -0,0 +1,201 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include "kernels.hh"
|
||||
#include "test_fixture.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup texCubemap texCubemap
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemap` and read mode set to `hipReadModeElementType`. The
|
||||
* test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemap.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeElementType", "", char, unsigned char, short,
|
||||
unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc();
|
||||
|
||||
TextureTestFixture<TestType> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
texCubemapKernel<vec4<TestType>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemap` and read mode set to `hipReadModeNormalizedFloat`.
|
||||
* The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemap.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
|
||||
short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
|
||||
|
||||
TextureTestFixture<TestType, true> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
texCubemapKernel<vec4<float>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,201 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include "kernels.hh"
|
||||
#include "test_fixture.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup texCubemapGrad texCubemapGrad
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapGrad` and read mode set to `hipReadModeElementType`.
|
||||
* The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapGrad.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeElementType", "", char, unsigned char,
|
||||
short, unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc();
|
||||
|
||||
TextureTestFixture<TestType> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
texCubemapGradKernel<vec4<TestType>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, float4{}, float4{});
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapGrad` and read mode set to
|
||||
* `hipReadModeNormalizedFloat`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapGrad.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
|
||||
short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
|
||||
|
||||
TextureTestFixture<TestType, true> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
texCubemapGradKernel<vec4<float>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, float4{}, float4{});
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,209 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include "kernels.hh"
|
||||
#include "test_fixture.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup texCubemapLayered texCubemapLayered
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLayered` and read mode set to
|
||||
* `hipReadModeElementType`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLayered.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeElementType", "", char, unsigned char,
|
||||
short, unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.layers = 1;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc();
|
||||
|
||||
TextureTestFixture<TestType> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
for (auto layer = 0u; layer < params.layers; ++layer) {
|
||||
texCubemapLayeredKernel<vec4<TestType>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, layer);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("Layer: " << layer);
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLayered` and read mode set to
|
||||
* `hipReadModeNormalizedFloat`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLayered.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat", "", char,
|
||||
unsigned char, short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.layers = 1;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
|
||||
|
||||
TextureTestFixture<TestType, true> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
for (auto layer = 0u; layer < params.layers; ++layer) {
|
||||
texCubemapLayeredKernel<vec4<float>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, layer);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("Layer: " << layer);
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,209 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include "kernels.hh"
|
||||
#include "test_fixture.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup texCubemapLayeredGrad texCubemapLayeredGrad
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLayeredGrad` and read mode set to
|
||||
* `hipReadModeElementType`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLayeredGrad.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeElementType", "", char,
|
||||
unsigned char, short, unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.layers = 1;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc();
|
||||
|
||||
TextureTestFixture<TestType> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
for (auto layer = 0u; layer < params.layers; ++layer) {
|
||||
texCubemapLayeredGradKernel<vec4<TestType>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, float4{}, float4{});
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("Layer: " << layer);
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLayeredGrad` and read mode set to
|
||||
* `hipReadModeNormalizedFloat`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLayeredGrad.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat", "", char,
|
||||
unsigned char, short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.layers = 1;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
|
||||
|
||||
TextureTestFixture<TestType, true> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
for (auto layer = 0u; layer < params.layers; ++layer) {
|
||||
texCubemapLayeredGradKernel<vec4<float>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, float4{}, float4{});
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("Layer: " << layer);
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,209 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include "kernels.hh"
|
||||
#include "test_fixture.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup texCubemapLayeredLod texCubemapLayeredLod
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLayeredLod` and read mode set to
|
||||
* `hipReadModeElementType`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLayeredLod.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeElementType", "", char,
|
||||
unsigned char, short, unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.layers = 1;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc();
|
||||
|
||||
TextureTestFixture<TestType> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
for (auto layer = 0u; layer < params.layers; ++layer) {
|
||||
texCubemapLayeredLodKernel<vec4<TestType>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, 0.0);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("Layer: " << layer);
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLayeredLod` and read mode set to
|
||||
* `hipReadModeNormalizedFloat`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLayeredLod.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat", "", char,
|
||||
unsigned char, short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.layers = 1;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
|
||||
|
||||
TextureTestFixture<TestType, true> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
for (auto layer = 0u; layer < params.layers; ++layer) {
|
||||
texCubemapLayeredLodKernel<vec4<float>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, 0.0);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("Layer: " << layer);
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,201 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include "kernels.hh"
|
||||
#include "test_fixture.hh"
|
||||
|
||||
/**
|
||||
* @addtogroup texCubemapLod texCubemapLod
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLod` and read mode set to `hipReadModeElementType`.
|
||||
* The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLod.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeElementType", "", char, unsigned char,
|
||||
short, unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc();
|
||||
|
||||
TextureTestFixture<TestType> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
texCubemapLodKernel<vec4<TestType>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, 0.0);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `texCubemapLod` and read mode set to
|
||||
* `hipReadModeNormalizedFloat`. The test is performed with:
|
||||
* - normalized coordinates
|
||||
* - non-normalized coordinates
|
||||
* - Nearest-point sampling
|
||||
* - Linear filtering
|
||||
* - All combinations of different addressing modes.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/texCubemapLod.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
|
||||
short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 6);
|
||||
params.num_subdivisions = 4;
|
||||
params.cubemap = true;
|
||||
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
|
||||
|
||||
TextureTestFixture<TestType, true> fixture{params};
|
||||
|
||||
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
|
||||
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
|
||||
const auto [num_threads_z, num_blocks_z] = GetLaunchConfig(10, params.NumItersZ());
|
||||
|
||||
dim3 dim_grid;
|
||||
dim_grid.x = num_blocks_x;
|
||||
dim_grid.y = num_blocks_y;
|
||||
dim_grid.z = num_blocks_z;
|
||||
|
||||
dim3 dim_block;
|
||||
dim_block.x = num_threads_x;
|
||||
dim_block.y = num_threads_y;
|
||||
dim_block.z = num_threads_z;
|
||||
|
||||
texCubemapLodKernel<vec4<float>><<<dim_grid, dim_block>>>(
|
||||
fixture.out_alloc_d.ptr(), params.NumItersX(), params.NumItersY(), params.NumItersZ(),
|
||||
fixture.tex.object(), params.Width(), params.Height(), params.Depth(),
|
||||
params.num_subdivisions, params.tex_desc.normalizedCoords, 0.0);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
fixture.LoadOutput();
|
||||
|
||||
for (auto k = 0u; k < params.NumItersZ(); ++k) {
|
||||
for (auto j = 0u; j < params.NumItersY(); ++j) {
|
||||
for (auto i = 0u; i < params.NumItersX(); ++i) {
|
||||
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float y = GetCoordinate(j, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
float z = GetCoordinate(k, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
INFO("i: " << i);
|
||||
INFO("j: " << j);
|
||||
INFO("k: " << k);
|
||||
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
|
||||
INFO("Address mode X: " << AddressModeToString(params.tex_desc.addressMode[0]));
|
||||
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
|
||||
INFO("Address mode Z: " << AddressModeToString(params.tex_desc.addressMode[2]));
|
||||
INFO("x: " << std::fixed << std::setprecision(16) << x);
|
||||
INFO("y: " << std::fixed << std::setprecision(16) << y);
|
||||
INFO("z: " << std::fixed << std::setprecision(16) << z);
|
||||
|
||||
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
|
||||
|
||||
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -52,6 +52,64 @@ template <typename TexelType> class TextureReference {
|
||||
}
|
||||
}
|
||||
|
||||
TexelType TexCubemap(float x, float y, float z, const hipTextureDesc& tex_desc) const {
|
||||
x = tex_desc.normalizedCoords ? x * extent_.width : x;
|
||||
y = tex_desc.normalizedCoords ? y * extent_.height : y;
|
||||
z = tex_desc.normalizedCoords ? z * extent_.depth : z;
|
||||
|
||||
int face;
|
||||
float m, s, t;
|
||||
|
||||
if (std::abs(x) > std::abs(y) && std::abs(x) > std::abs(z)) {
|
||||
if (x >= 0) {
|
||||
face = 0;
|
||||
m = x;
|
||||
s = -z;
|
||||
t = -y;
|
||||
} else {
|
||||
face = 1;
|
||||
m = -x;
|
||||
s = z;
|
||||
t = -y;
|
||||
}
|
||||
} else if (std::abs(y) >= std::abs(x) && std::abs(y) > std::abs(z)) {
|
||||
if (y >= 0) {
|
||||
face = 2;
|
||||
m = y;
|
||||
s = x;
|
||||
t = z;
|
||||
} else {
|
||||
face = 3;
|
||||
m = -y;
|
||||
s = x;
|
||||
t = -z;
|
||||
}
|
||||
} else {
|
||||
if (z >= 0) {
|
||||
face = 4;
|
||||
m = z;
|
||||
s = x;
|
||||
t = -y;
|
||||
} else {
|
||||
face = 5;
|
||||
m = -z;
|
||||
s = -x;
|
||||
t = -y;
|
||||
}
|
||||
}
|
||||
|
||||
float coord1 = (s / m + 1) / 2;
|
||||
float coord2 = (t / m + 1) / 2;
|
||||
|
||||
if (tex_desc.filterMode == hipFilterModePoint) {
|
||||
return Sample(roundf(coord1), roundf(coord2), face, tex_desc.addressMode);
|
||||
} else if (tex_desc.filterMode == hipFilterModeLinear) {
|
||||
return LinearFiltering(coord1, coord2, face, tex_desc.addressMode);
|
||||
} else {
|
||||
throw std::invalid_argument("Invalid hipFilterMode value");
|
||||
}
|
||||
}
|
||||
|
||||
TexelType Tex1DLayered(float x, int layer, const hipTextureDesc& tex_desc) const {
|
||||
x = tex_desc.normalizedCoords ? x * extent_.width : x;
|
||||
if (tex_desc.filterMode == hipFilterModePoint) {
|
||||
@@ -236,10 +294,17 @@ template <typename TexelType> class TextureReference {
|
||||
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 = coord - 0.5f;
|
||||
const auto coordB = FloatToNBitFractional<8>(coord - 0.5f);
|
||||
const auto index = floorf(coordB);
|
||||
const FixedPoint<8> coeff = coordB - index;
|
||||
const auto coeff = coordB - index;
|
||||
|
||||
return {index, coeff};
|
||||
}
|
||||
|
||||
Fai riferimento in un nuovo problema
Block a user