AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

EXSWHTEC-357 - Implement tests for 1D mipmapped texture device functions #431

Просмотр исходного кода 
Change-Id: I74b4bc2a73ec8bbe21460eda8cdc79210fe2501e
Этот коммит содержится в:
Mirza Halilčević
2023-12-28 14:29:28 +01:00
коммит произвёл Rakesh Roy
родитель 303836a7fd
Коммит 0ecb874b52
8 изменённых файлов: 165 добавлений и 98 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
catch/include/resource_guards.hh
+36
Просмотреть файл
@@ -216,6 +216,42 @@ template <typename T> class ArrayAllocGuard {
const hipExtent extent_;
};
template <typename T> class MipmappedArrayAllocGuard {
public:
// extent should contain logical width
MipmappedArrayAllocGuard(const hipExtent extent, const unsigned int levels,
const unsigned int flags)
: extent_{extent}, levels_{levels} {
hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
HIP_CHECK(hipMallocMipmappedArray(&ptr_, &desc, extent_, levels_, flags));
}
MipmappedArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u)
: MipmappedArrayAllocGuard{extent, 1, flags} {}
~MipmappedArrayAllocGuard() { static_cast<void>(hipFreeMipmappedArray(ptr_)); }
MipmappedArrayAllocGuard(const MipmappedArrayAllocGuard&) = delete;
MipmappedArrayAllocGuard(MipmappedArrayAllocGuard&&) = delete;
hipMipmappedArray_t ptr() const { return ptr_; }
hipArray_t GetLevel(unsigned int level) {
hipArray_t ret;
HIP_CHECK(hipGetMipmappedArrayLevel(&ret, ptr_, level));
return ret;
}
hipExtent extent() const { return extent_; }
unsigned int levels() const { return levels_; }
private:
hipMipmappedArray_t ptr_ = nullptr;
const hipExtent extent_;
const unsigned int levels_;
};
template <typename T> class DrvArrayAllocGuard {
public:
// extent should contain width in bytes
catch/unit/texture/kernels.hh
+1 -1
Просмотреть файл
@@ -87,7 +87,7 @@ __global__ void tex1DGradKernel(TexelType* const out, size_t N, hipTextureObject
template <typename TexelType>
__global__ void tex1DLayeredGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions,
bool normalized_coords, float dx, float dy, int layer) {
bool normalized_coords, int layer, float dx, float dy) {
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
catch/unit/texture/test_fixture.hh
+21 -6
Просмотреть файл
@@ -99,19 +99,25 @@ template <typename TestType> struct TextureTestParams {
tex_desc.addressMode[0] = address_mode_x;
if (extent.height) tex_desc.addressMode[1] = address_mode_y;
if (extent.depth) tex_desc.addressMode[2] = address_mode_z;
tex_desc.mipmapFilterMode = tex_desc.filterMode;
}
};
template <typename TestType, bool normalized_read = false> struct TextureTestFixture {
template <typename TestType, bool normalized_read = false, bool mipmap = false>
struct TextureTestFixture {
using VecType = vec4<TestType>;
using OutType = std::conditional_t<normalized_read, vec4<float>, VecType>;
template <typename T>
using ArrayAllocGuardType =
std::conditional_t<mipmap, MipmappedArrayAllocGuard<T>, ArrayAllocGuard<T>>;
TextureTestParams<TestType> params;
hipResourceDesc res_desc;
LinearAllocGuard<VecType> host_alloc;
TextureReference<VecType> tex_h;
ArrayAllocGuard<VecType> tex_alloc_d;
ArrayAllocGuardType<VecType> tex_alloc_d;
TextureGuard tex;
LinearAllocGuard<OutType> out_alloc_d;
std::vector<OutType> out_alloc_h;
@@ -131,9 +137,13 @@ template <typename TestType, bool normalized_read = false> struct TextureTestFix
SetVec4<TestType>(host_alloc.ptr()[i], i + test_value_offset);
}
hipMemcpy3DParms memcpy_params;
hipMemcpy3DParms memcpy_params = {};
memset(&memcpy_params, 0 sizeof(hipMemcpy3DParms));
memcpy_params.dstArray = tex_alloc_d.ptr();
if constexpr (mipmap) {
memcpy_params.dstArray = tex_alloc_d.GetLevel(0);
} else {
memcpy_params.dstArray = tex_alloc_d.ptr();
}
memcpy_params.extent = params.LayeredExtent();
memcpy_params.extent.height = memcpy_params.extent.height ?: 1;
memcpy_params.extent.depth = memcpy_params.extent.depth ?: 1;
@@ -143,8 +153,13 @@ template <typename TestType, bool normalized_read = false> struct TextureTestFix
HIP_CHECK(hipMemcpy3D(&memcpy_params));
memset(&res_desc, 0, sizeof(res_desc));
res_desc.resType = hipResourceTypeArray;
res_desc.res.array.array = tex_alloc_d.ptr();
if constexpr (mipmap) {
res_desc.resType = hipResourceTypeMipmappedArray;
res_desc.res.mipmap.mipmap = tex_alloc_d.ptr();
} else {
res_desc.resType = hipResourceTypeArray;
res_desc.res.array.array = tex_alloc_d.ptr();
}
return &res_desc;
}
catch/unit/texture/tex1D.cc
-1
Просмотреть файл
@@ -127,7 +127,6 @@ TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeNormalizedFloat", "", char, unsi
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val =
catch/unit/texture/tex1DGrad.cc
+4 -5
Просмотреть файл
@@ -46,7 +46,7 @@ THE SOFTWARE.
* - unit/texture/tex1DGrad.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsigned char, short,
unsigned short, int, unsigned int, float) {
@@ -57,7 +57,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsi
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType> fixture{params};
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
@@ -99,7 +99,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsi
* - unit/texture/tex1DGrad.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
short, unsigned short) {
@@ -110,7 +110,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char,
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true> fixture{params};
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
@@ -127,7 +127,6 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char,
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val =
catch/unit/texture/tex1DLayeredGrad.cc
+50 -40
Просмотреть файл
@@ -46,7 +46,7 @@ THE SOFTWARE.
* - unit/texture/tex1DLayeredGrad.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", char, unsigned char,
short, unsigned short, int, unsigned int, float) {
@@ -54,33 +54,38 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", cha
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(1024, 0, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType> fixture{params};
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DLayeredGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f, 0);
fixture.LoadOutput();
for (auto layer = 0u; layer < params.layers; ++layer) {
tex1DLayeredGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, 0.5f, 0.5f);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
fixture.LoadOutput();
INFO("Index: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
INFO("Layer: " << layer);
INFO("Index: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
}
}
}
@@ -99,7 +104,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", cha
* - unit/texture/tex1DLayeredGrad.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "", char,
unsigned char, short, unsigned short) {
@@ -107,34 +112,39 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "",
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(1024, 0, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true> fixture{params};
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DLayeredGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, 0.5f, 0.5f, 0);
fixture.LoadOutput();
for (auto layer = 0u; layer < params.layers; ++layer) {
tex1DLayeredGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, 0.5f, 0.5f);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
fixture.LoadOutput();
INFO("i: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex1D(x, 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);
INFO("Layer: " << layer);
INFO("i: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex1DLayered(x, layer, 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);
}
}
}
catch/unit/texture/tex1DLayeredLod.cc
+49 -40
Просмотреть файл
@@ -46,7 +46,7 @@ THE SOFTWARE.
* - unit/texture/tex1DLayeredLod.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char, unsigned char,
short, unsigned short, int, unsigned int, float) {
@@ -54,33 +54,38 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(1024, 0, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType> fixture{params};
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DLayeredLodKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, 0, 0);
fixture.LoadOutput();
for (auto layer = 0u; layer < params.layers; ++layer) {
tex1DLayeredLodKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, 0);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
fixture.LoadOutput();
INFO("Index: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
INFO("Layer: " << layer);
INFO("Index: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
}
}
}
@@ -99,7 +104,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char
* - unit/texture/tex1DLayeredLod.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
short, unsigned short) {
@@ -107,34 +112,38 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "",
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(1024, 0, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true> fixture{params};
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DLayeredLodKernel<vec4<float>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, 0, 0);
fixture.LoadOutput();
for (auto layer = 0u; layer < params.layers; ++layer) {
tex1DLayeredLodKernel<vec4<float>><<<num_blocks, num_threads>>>(
fixture.out_alloc_d.ptr(), params.NumItersX(), fixture.tex.object(), params.Width(),
params.num_subdivisions, params.tex_desc.normalizedCoords, layer, 0);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
fixture.LoadOutput();
INFO("i: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
for (auto i = 0u; i < params.NumItersX(); ++i) {
float x = GetCoordinate(i, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex1D(x, 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);
INFO("Layer: " << layer);
INFO("i: " << i);
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex1DLayered(x, layer, 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);
}
}
}
catch/unit/texture/tex1DLod.cc
+4 -5
Просмотреть файл
@@ -46,7 +46,7 @@ THE SOFTWARE.
* - unit/texture/tex1DLod.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsigned char, short,
unsigned short, int, unsigned int, float) {
@@ -57,7 +57,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsig
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType> fixture{params};
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DLodKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
@@ -99,7 +99,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsig
* - unit/texture/tex1DLod.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 5.7
*/
TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char, short,
unsigned short) {
@@ -110,7 +110,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, u
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true> fixture{params};
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DLodKernel<vec4<float>><<<num_blocks, num_threads>>>(
@@ -127,7 +127,6 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, u
INFO("Filtering mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("Normalized coordinates: " << std::boolalpha << params.tex_desc.normalizedCoords);
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val =