EXSWHTEC-319 - Implement tests for 3D texture device functions #368
Change-Id: Ifc7a6feae1a9567df29d35928d8c4328497bfc97
This commit is contained in:
Mirza Halilčević
کامیت شده توسط
Rakesh Roy
Rakesh Roy
والد
37d8529a9b
کامیت
5fe762ae5d
@@ -211,6 +211,15 @@
|
||||
"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",
|
||||
|
||||
@@ -60,6 +60,9 @@ set(TEST_SRC
|
||||
tex1DLayeredGrad.cc
|
||||
tex1DLayeredLod.cc
|
||||
tex1DLod.cc
|
||||
tex3D.cc
|
||||
tex3DLod.cc
|
||||
tex3DGrad.cc
|
||||
)
|
||||
|
||||
if(WIN32)
|
||||
|
||||
@@ -131,6 +131,48 @@ __global__ void tex3DKernel(TexelType* const out, size_t N_x, size_t N_y, size_t
|
||||
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 tex3DLodKernel(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] = tex3DLod<TexelType>(tex_obj, x, y, z, level);
|
||||
}
|
||||
|
||||
template <typename TexelType>
|
||||
__global__ void tex3DGradKernel(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] = tex3DGrad<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,
|
||||
|
||||
@@ -0,0 +1,192 @@
|
||||
/*
|
||||
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 tex3D tex3D
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `tex3D` 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/tex3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeElementType", "", char, unsigned char, short,
|
||||
unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 4, 2);
|
||||
params.num_subdivisions = 2;
|
||||
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;
|
||||
|
||||
tex3DKernel<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 i = 0u; i < params.NumIters(); ++i) {
|
||||
const auto plane = i % (params.NumItersX() * params.NumItersY());
|
||||
float x = plane % params.NumItersX();
|
||||
float y = plane / params.NumItersX();
|
||||
float z = i / (params.NumItersX() * params.NumItersY());
|
||||
|
||||
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
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);
|
||||
|
||||
const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `tex3D` 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/tex3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short,
|
||||
unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 2);
|
||||
params.num_subdivisions = 2;
|
||||
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;
|
||||
|
||||
tex3DKernel<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 i = 0u; i < params.NumIters(); ++i) {
|
||||
const auto plane = i % (params.NumItersX() * params.NumItersY());
|
||||
float x = plane % params.NumItersX();
|
||||
float y = plane / params.NumItersX();
|
||||
float z = i / (params.NumItersX() * params.NumItersY());
|
||||
|
||||
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
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 ref_val = Vec4Map<TestType>(fixture.tex_h.Tex3D(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,191 @@
|
||||
/*
|
||||
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 tex3DGrad tex3DGrad
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `tex3DGrad` 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/tex3DGrad.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeElementType", "", char, unsigned char, short,
|
||||
unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 2);
|
||||
params.num_subdivisions = 2;
|
||||
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;
|
||||
|
||||
tex3DGradKernel<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 i = 0u; i < params.NumIters(); ++i) {
|
||||
const auto plane = i % (params.NumItersX() * params.NumItersY());
|
||||
float x = plane % params.NumItersX();
|
||||
float y = plane / params.NumItersX();
|
||||
float z = i / (params.NumItersX() * params.NumItersY());
|
||||
|
||||
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
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);
|
||||
|
||||
const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `tex3DGrad` 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/tex3DGrad.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
|
||||
short, unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 2);
|
||||
params.num_subdivisions = 2;
|
||||
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;
|
||||
|
||||
tex3DGradKernel<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 i = 0u; i < params.NumIters(); ++i) {
|
||||
const auto plane = i % (params.NumItersX() * params.NumItersY());
|
||||
float x = plane % params.NumItersX();
|
||||
float y = plane / params.NumItersX();
|
||||
float z = i / (params.NumItersX() * params.NumItersY());
|
||||
|
||||
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
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 ref_val = Vec4Map<TestType>(fixture.tex_h.Tex3D(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,191 @@
|
||||
/*
|
||||
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 tex3DLod tex3DLod
|
||||
* @{
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `tex3DLod` 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/tex3DLod.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeElementType", "", char, unsigned char, short,
|
||||
unsigned short, int, unsigned int, float) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 2);
|
||||
params.num_subdivisions = 2;
|
||||
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;
|
||||
|
||||
tex3DLodKernel<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 i = 0u; i < params.NumIters(); ++i) {
|
||||
const auto plane = i % (params.NumItersX() * params.NumItersY());
|
||||
float x = plane % params.NumItersX();
|
||||
float y = plane / params.NumItersX();
|
||||
float z = i / (params.NumItersX() * params.NumItersY());
|
||||
|
||||
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
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);
|
||||
|
||||
const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test texture fetching with `tex3DLod` 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/tex3DLod.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short,
|
||||
unsigned short) {
|
||||
CHECK_IMAGE_SUPPORT;
|
||||
|
||||
TextureTestParams<TestType> params = {};
|
||||
params.extent = make_hipExtent(2, 2, 2);
|
||||
params.num_subdivisions = 2;
|
||||
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;
|
||||
|
||||
tex3DLodKernel<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 i = 0u; i < params.NumIters(); ++i) {
|
||||
const auto plane = i % (params.NumItersX() * params.NumItersY());
|
||||
float x = plane % params.NumItersX();
|
||||
float y = plane / params.NumItersX();
|
||||
float z = i / (params.NumItersX() * params.NumItersY());
|
||||
|
||||
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
y = GetCoordinate(y, params.NumItersY(), params.Height(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
z = GetCoordinate(z, params.NumItersZ(), params.Depth(), params.num_subdivisions,
|
||||
params.tex_desc.normalizedCoords);
|
||||
|
||||
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 ref_val = Vec4Map<TestType>(fixture.tex_h.Tex3D(x, y, z, params.tex_desc),
|
||||
NormalizeInteger<TestType>);
|
||||
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
|
||||
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
|
||||
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
|
||||
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
|
||||
}
|
||||
}
|
||||
مرجع در شماره جدید
Block a user