SWDEV-393139 - Add mipmap tests
Add 1D/2D/3D mipmap tests covering necessary data types and
config parameters.
Change-Id: I6079d9f3a2ddf4e42b8a6f7f3902322cfca02cff
[ROCm/hip-tests commit: 07ec2d5a0c]
This commit is contained in:
taosang2
Rakesh Roy
@@ -128,7 +128,12 @@ template<
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ")";
|
||||
if constexpr (std::is_same<decltype(T::x), char>::value ||
|
||||
std::is_same<decltype(T::x), unsigned char>::value) {
|
||||
os << "(" << static_cast<int>(t.x) << ")";
|
||||
} else {
|
||||
os << "(" << t.x << ")";
|
||||
}
|
||||
return os.str();
|
||||
}
|
||||
|
||||
@@ -138,7 +143,12 @@ template<
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ", " << t.y << ")";
|
||||
if constexpr (std::is_same<decltype(T::x), char>::value ||
|
||||
std::is_same<decltype(T::x), unsigned char>::value) {
|
||||
os << "(" << static_cast<int>(t.x) << ", " << static_cast<int>(t.y) << ")";
|
||||
} else {
|
||||
os << "(" << t.x << ", " << t.y << ")";
|
||||
}
|
||||
return os.str();
|
||||
}
|
||||
|
||||
@@ -148,7 +158,13 @@ template<
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ", " << t.y << ", " << t.z << ")";
|
||||
if constexpr (std::is_same<decltype(T::x), char>::value ||
|
||||
std::is_same<decltype(T::x), unsigned char>::value) {
|
||||
os << "(" << static_cast<int>(t.x) << ", " << static_cast<int>(t.y) << ", " <<
|
||||
static_cast<int>(t.z) << ")";
|
||||
} else {
|
||||
os << "(" << t.x << ", " << t.y << ", " << t.z << ")";
|
||||
}
|
||||
return os.str();
|
||||
}
|
||||
|
||||
@@ -158,7 +174,13 @@ template<
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")";
|
||||
if constexpr (std::is_same<decltype(T::x), char>::value ||
|
||||
std::is_same<decltype(T::x), unsigned char>::value) {
|
||||
os << "(" << static_cast<int>(t.x) << ", " << static_cast<int>(t.y) << ", " <<
|
||||
static_cast<int>(t.z) << ", " << static_cast<int>(t.w) << ")";
|
||||
} else {
|
||||
os << "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")";
|
||||
}
|
||||
return os.str();
|
||||
}
|
||||
|
||||
@@ -168,7 +190,12 @@ template<
|
||||
std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os << t;
|
||||
if constexpr (std::is_same<T, char>::value ||
|
||||
std::is_same<T, unsigned char>::value) {
|
||||
os << static_cast<int>(t);
|
||||
} else {
|
||||
os << t;
|
||||
}
|
||||
return os.str();
|
||||
}
|
||||
|
||||
@@ -180,7 +207,13 @@ static inline T getRandom() {
|
||||
} else {
|
||||
r = std::rand() / (RAND_MAX + 1.);
|
||||
}
|
||||
return static_cast<T>(std::numeric_limits < T > ::max() * r);
|
||||
if constexpr (std::is_floating_point<T>::value) {
|
||||
// Restrict any float within (-1000, 1000)
|
||||
// to prevent too big float value that would make caculation sick
|
||||
return static_cast<T>(r * 1000.);
|
||||
} else {
|
||||
return static_cast<T>(std::numeric_limits<T>::max() * r);
|
||||
}
|
||||
}
|
||||
|
||||
template<
|
||||
@@ -214,3 +247,62 @@ template<
|
||||
static inline void initVal(T &val) {
|
||||
val = getRandom<T>();
|
||||
}
|
||||
|
||||
/*Convert normalized floatx to typex*/
|
||||
template <typename T, typename F> inline __device__ T getTypeFromNormalizedFloat(const F &f) {
|
||||
T t;
|
||||
if constexpr (std::is_scalar<T>::value)
|
||||
t = static_cast<T>(f.x * std::numeric_limits<T>::max());
|
||||
else {
|
||||
if constexpr (rank<T>() > 0)
|
||||
t.x = static_cast<decltype(T::x)>(f.x * std::numeric_limits<decltype(T::x)>::max());
|
||||
if constexpr (rank<T>() > 1)
|
||||
t.y = static_cast<decltype(T::y)>(f.y * std::numeric_limits<decltype(T::y)>::max());
|
||||
if constexpr (rank<T>() > 2)
|
||||
t.z = static_cast<decltype(T::z)>(f.z * std::numeric_limits<decltype(T::z)>::max());
|
||||
if constexpr (rank<T>() > 3)
|
||||
t.w = static_cast<decltype(T::w)>(f.w * std::numeric_limits<decltype(T::w)>::max());
|
||||
}
|
||||
return t;
|
||||
}
|
||||
|
||||
/*Convert typex to normalized floatx*/
|
||||
template <class T>
|
||||
inline auto getNormalizedFloatType(const T &t) {
|
||||
if constexpr (std::is_scalar<T>::value)
|
||||
return static_cast<float>(t) / std::numeric_limits<T>::max();
|
||||
else {
|
||||
if constexpr (rank<T>() == 1) {
|
||||
float1 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max()};
|
||||
return f;
|
||||
}
|
||||
if constexpr (rank<T>() == 2) {
|
||||
float2 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max(),
|
||||
static_cast<float>(t.y) / std::numeric_limits<decltype(T::y)>::max()};
|
||||
return f;
|
||||
}
|
||||
if constexpr (rank<T>() == 3) {
|
||||
float3 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max(),
|
||||
static_cast<float>(t.y) / std::numeric_limits<decltype(T::y)>::max(),
|
||||
static_cast<float>(t.z) / std::numeric_limits<decltype(T::z)>::max()};
|
||||
return f;
|
||||
}
|
||||
if constexpr (rank<T>() == 4) {
|
||||
float4 f{static_cast<float>(t.x) / std::numeric_limits<decltype(T::x)>::max(),
|
||||
static_cast<float>(t.y) / std::numeric_limits<decltype(T::y)>::max(),
|
||||
static_cast<float>(t.z) / std::numeric_limits<decltype(T::z)>::max(),
|
||||
static_cast<float>(t.w) / std::numeric_limits<decltype(T::w)>::max()};
|
||||
return f;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/*Check if T is floatx*/
|
||||
template <typename T> inline bool constexpr isFloat() {
|
||||
if constexpr (std::is_scalar<T>::value)
|
||||
return std::is_floating_point<T>::value;
|
||||
else {
|
||||
return std::is_floating_point<decltype(T::x)>::value;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -49,6 +49,11 @@ inline __host__ __device__ operator*=(T &a, const decltype(T::x) &b)
|
||||
}
|
||||
#endif // HT_NVIDIA
|
||||
|
||||
template <typename T> struct mipmapLevelArray {
|
||||
T* data; // level array data
|
||||
hipExtent e; // level array size
|
||||
};
|
||||
|
||||
// From CIE 1931 color space to sRGB
|
||||
inline float hipSRGBMap(float fc) {
|
||||
double c = static_cast<double>(fc);
|
||||
@@ -216,7 +221,7 @@ T hipTextureGetValue(const T *data, const int x, const int width,
|
||||
default:
|
||||
break;
|
||||
}
|
||||
if (sRGB && std::is_same<T, float4>::value) {
|
||||
if constexpr (sRGB && std::is_same<T, float4>::value) {
|
||||
result = hipSRGBUnmap(result);
|
||||
}
|
||||
return result;
|
||||
|
||||
@@ -30,7 +30,6 @@ set(TEST_SRC
|
||||
hipSimpleTexture3D.cc
|
||||
hipTextureRef2D.cc
|
||||
hipSimpleTexture2DLayered.cc
|
||||
hipTextureMipmapObj2D.cc
|
||||
hipBindTex2DPitch.cc
|
||||
hipBindTexRef1DFetch.cc
|
||||
hipTex1DFetchCheckModes.cc
|
||||
@@ -47,6 +46,15 @@ set(TEST_SRC
|
||||
hipBindTextureToMipmappedArray.cc
|
||||
)
|
||||
|
||||
if(WIN32)
|
||||
set(TEST_SRC
|
||||
${TEST_SRC}
|
||||
hipTextureMipmapObj1D.cc
|
||||
hipTextureMipmapObj2D.cc
|
||||
hipTextureMipmapObj3D.cc
|
||||
)
|
||||
endif()
|
||||
|
||||
hip_add_exe_to_target(NAME TextureTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
|
||||
@@ -0,0 +1,471 @@
|
||||
/*
|
||||
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 <hip_array_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_texture_helper.hh>
|
||||
#include <algorithm>
|
||||
|
||||
#pragma clang diagnostic ignored "-Wunused-variable"
|
||||
#pragma clang diagnostic ignored "-Wunused-parameter"
|
||||
|
||||
//#define TEST_TEXTURE // Only for float2
|
||||
static constexpr bool printLog = false; // Print log for debugging
|
||||
|
||||
// Populate mipmap next level array
|
||||
template <typename T, hipTextureReadMode readMode>
|
||||
static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut,
|
||||
hipTextureObject_t texIn, unsigned int width, T* data = nullptr) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
float px = 1.0 / float(width);
|
||||
|
||||
if (x < width) {
|
||||
if constexpr (readMode == hipReadModeElementType) {
|
||||
T t = tex1D<T>(texIn, x * px);
|
||||
#ifdef TEST_TEXTURE
|
||||
printf("(%d/%u):(%f)\n", x, width, t.x;
|
||||
#endif
|
||||
surf1Dwrite<T>(t, surfOut, x * sizeof(T));
|
||||
if (data) data[x] = t; // record it for later verification
|
||||
}
|
||||
if constexpr (readMode == hipReadModeNormalizedFloat) {
|
||||
float4 t = tex1D<float4>(texIn, x * px);
|
||||
T tc = getTypeFromNormalizedFloat<T, float4>(t);
|
||||
surf1Dwrite<T>(tc, surfOut, x * sizeof(T));
|
||||
if (data) data[x] = tc;
|
||||
}
|
||||
// Users have freedom to use other methods to init level array
|
||||
// for example, use averge of surrounding pixes
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width,
|
||||
float offsetX, float lod, T* data = nullptr) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
float px = 1.0 / float(width);
|
||||
|
||||
if (x < width) {
|
||||
T t = tex1DLod<T>(texMipmap, (x + offsetX) * px, lod);
|
||||
if (data) data[x] = t;
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T, hipTextureReadMode readMode = hipReadModeElementType,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void populateMipmaps(hipMipmappedArray_t mipmapArray, hipExtent size,
|
||||
std::vector<mipmapLevelArray<T>>& mipmapData) {
|
||||
size_t width = size.width;
|
||||
unsigned int level = 0;
|
||||
|
||||
while (width != 1) {
|
||||
hipArray_t levelArray = nullptr, nextLevelArray = nullptr;
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&levelArray, mipmapArray, level));
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&nextLevelArray, mipmapArray, level + 1));
|
||||
|
||||
hipExtent levelArraySize {0, 0, 0};
|
||||
HIP_CHECK(hipArrayGetInfo(nullptr, &levelArraySize, nullptr, levelArray));
|
||||
if(levelArraySize.width != width) {
|
||||
fprintf(stderr, "Level %u: size (%zu, %zu, %zu) != Expected size (%zu, 0, 0)\n", level,
|
||||
levelArraySize.width, levelArraySize.height, levelArraySize.depth, width);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
width = width >> 1 ? width >> 1: 1;
|
||||
|
||||
hipExtent nextLevelArraySize {0, 0, 0};
|
||||
HIP_CHECK(hipArrayGetInfo(nullptr, &nextLevelArraySize, nullptr, nextLevelArray));
|
||||
if(nextLevelArraySize.width != width) {
|
||||
fprintf(stderr, "Next level %u: size (%zu, %zu, %zu) != Expected size (%zu, 0, 0)\n",
|
||||
level + 1, nextLevelArraySize.width, nextLevelArraySize.height,
|
||||
nextLevelArraySize.depth, width);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
hipTextureObject_t texIn;
|
||||
hipResourceDesc texRes;
|
||||
memset(&texRes, 0, sizeof(hipResourceDesc));
|
||||
texRes.resType = hipResourceTypeArray;
|
||||
texRes.res.array.array = levelArray;
|
||||
|
||||
hipTextureDesc texDescr;
|
||||
memset(&texDescr, 0, sizeof(hipTextureDesc));
|
||||
texDescr.normalizedCoords = 1; // To populate next level array smoothly
|
||||
texDescr.filterMode = filterMode;
|
||||
texDescr.addressMode[0] = addressMode;
|
||||
texDescr.addressMode[1] = addressMode;
|
||||
texDescr.addressMode[2] = addressMode;
|
||||
texDescr.readMode = readMode;
|
||||
HIP_CHECK(hipCreateTextureObject(&texIn, &texRes, &texDescr, NULL));
|
||||
|
||||
hipSurfaceObject_t surfOut;
|
||||
hipResourceDesc surfRes;
|
||||
memset(&surfRes, 0, sizeof(hipResourceDesc));
|
||||
surfRes.resType = hipResourceTypeArray;
|
||||
surfRes.res.array.array = nextLevelArray;
|
||||
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfOut, &surfRes));
|
||||
size_t size = width * sizeof(T);
|
||||
mipmapLevelArray<T> data{nullptr, {width, 0, 0}};
|
||||
HIP_CHECK(hipHostMalloc((void**)&data.data, size));
|
||||
memset(data.data, 0, size);
|
||||
|
||||
dim3 blockSize(16, 1, 1);
|
||||
dim3 gridSize((width + blockSize.x - 1) / blockSize.x, 1, 1);
|
||||
|
||||
populateMipmapNextLevelArray<T, readMode>
|
||||
<<<gridSize, blockSize>>>(surfOut, texIn, width, data.data);
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfOut));
|
||||
HIP_CHECK(hipDestroyTextureObject(texIn));
|
||||
HIP_CHECK(hipFreeArray(levelArray));
|
||||
HIP_CHECK(hipFreeArray(nextLevelArray));
|
||||
mipmapData.push_back(data); // For later verification
|
||||
level++;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void verifyMipmapLevel(hipTextureObject_t texMipmap, T* data, size_t width,
|
||||
float level, float offsetX) {
|
||||
T* hOutput = nullptr;
|
||||
size_t size = width * sizeof(T);
|
||||
HIP_CHECK(hipHostMalloc((void**)&hOutput, size));
|
||||
memset(hOutput, 0, size);
|
||||
|
||||
dim3 blockSize(16, 1, 1);
|
||||
dim3 gridSize((width + blockSize.x - 1) / blockSize.x, 1, 1);
|
||||
|
||||
getMipmap<T><<<gridSize, blockSize>>>(texMipmap, width, offsetX, level, hOutput);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
for (unsigned int i = 0; i < width; i++) {
|
||||
T cpuExpected = getExpectedValue<T, addressMode, filterMode, false>(width, i + offsetX, data);
|
||||
T gpuOutput = hOutput[i];
|
||||
bool match = hipTextureSamplingVerify<T, filterMode, false>(gpuOutput, cpuExpected);
|
||||
if (!match) {
|
||||
WARN("Mismatch at (level "
|
||||
<< level << ": " << i << " -> " << (i + offsetX)
|
||||
<< ") GPU output : " << getString(gpuOutput)
|
||||
<< " CPU expected: " << getString(cpuExpected) << ", data[" << i
|
||||
<< "]:" << getString(data[i]) << "\n");
|
||||
REQUIRE(false);
|
||||
} else if (printLog) {
|
||||
WARN("Matching at (level "
|
||||
<< level << ": " << i << " -> " << (i + offsetX)
|
||||
<< ") GPU output : " << getString(gpuOutput)
|
||||
<< " CPU expected: " << getString(cpuExpected) << ", data[" << i
|
||||
<< "]:" << getString(data[i]) << "\n");
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipHostFree(hOutput));
|
||||
}
|
||||
|
||||
template <typename T, hipTextureReadMode readMode = hipReadModeElementType,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void testMipmapTextureObj(size_t width, float offsetX = 0.) {
|
||||
std::vector<mipmapLevelArray<T>> mipmapData;
|
||||
size_t size = width * sizeof(T);
|
||||
mipmapLevelArray<T> data{nullptr, {width, 0, 0}};
|
||||
HIP_CHECK(hipHostMalloc((void**)&data.data, size));
|
||||
memset(data.data, 0, size);
|
||||
|
||||
for (int i = 0; i < width; i++) {
|
||||
if (isFloat<T>() && filterMode == hipFilterModeLinear) {
|
||||
/*
|
||||
* For linear sampling of images, the GPU does not use IEEE floating point types, it uses
|
||||
* lower-precision sampling optimized formats. Also, those formats often change between GPU
|
||||
* generations. So counting on IEEE precision and accuracy when doing linear sampling
|
||||
* is mistaken. To workaround this issue, we can initialize float pixels on a retively
|
||||
* smoothy surface.
|
||||
*/
|
||||
data.data[i] = T( float(i) * (float(i) - width + 1));
|
||||
} else {
|
||||
initVal(data.data[i]); // Randomize initial values
|
||||
}
|
||||
}
|
||||
mipmapData.push_back(data); // record level 0 data for later verification
|
||||
|
||||
// Get the max mipmap levels in terms of image size
|
||||
const unsigned int maxLevels = 1 + std::log2(width);
|
||||
|
||||
// create mipmap array
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
|
||||
hipMipmappedArray_t mipmapArray = nullptr;
|
||||
hipExtent extent { width, 0, 0 };
|
||||
HIP_CHECK(hipMallocMipmappedArray(&mipmapArray, &desc, extent, maxLevels));
|
||||
|
||||
// Initialize level 0
|
||||
hipArray_t levelArray;
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&levelArray, mipmapArray, 0));
|
||||
hipMemcpy3DParms copyParams{};
|
||||
copyParams.srcPtr = make_hipPitchedPtr(data.data, width * sizeof(T),
|
||||
width, 1);
|
||||
copyParams.dstArray = levelArray;
|
||||
copyParams.extent.width = width;
|
||||
copyParams.extent.height = 1;
|
||||
copyParams.extent.depth = 1;
|
||||
copyParams.kind = hipMemcpyHostToDevice;
|
||||
HIP_CHECK(hipMemcpy3D(©Params));
|
||||
|
||||
// Populate other mipmap levels based on level 0
|
||||
populateMipmaps<T, readMode, filterMode, addressMode>(mipmapArray, extent, mipmapData);
|
||||
|
||||
if(maxLevels != mipmapData.size()) {
|
||||
fprintf(stderr, "maxLevels %u != mipmapData.size() %zu\n", maxLevels, mipmapData.size());
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
hipResourceDesc resDescr;
|
||||
memset(&resDescr, 0, sizeof(hipResourceDesc));
|
||||
resDescr.resType = hipResourceTypeMipmappedArray; // For mipmap texture
|
||||
resDescr.res.mipmap.mipmap = mipmapArray;
|
||||
|
||||
hipTextureDesc texDescr;
|
||||
memset(&texDescr, 0, sizeof(hipTextureDesc));
|
||||
|
||||
texDescr.normalizedCoords = 1; // normalizedCoords must be 1 for mipmap array
|
||||
texDescr.filterMode = filterMode;
|
||||
texDescr.mipmapFilterMode = filterMode;
|
||||
texDescr.addressMode[0] = addressMode;
|
||||
texDescr.addressMode[1] = addressMode;
|
||||
texDescr.addressMode[2] = addressMode;
|
||||
// Ignored in AMD Hw sampler SRD, but cuda need it
|
||||
texDescr.maxMipmapLevelClamp = float(maxLevels - 1);
|
||||
texDescr.readMode = readMode;
|
||||
|
||||
hipTextureObject_t texMipmap = nullptr;
|
||||
HIP_CHECK(hipCreateTextureObject(&texMipmap, &resDescr, &texDescr, NULL));
|
||||
|
||||
for(unsigned int level = 0; level < mipmapData.size(); level++){
|
||||
mipmapLevelArray<T> &data = mipmapData.at(level);
|
||||
|
||||
if constexpr (hipReadModeNormalizedFloat == readMode) {
|
||||
typedef decltype(getNormalizedFloatType(*data.data)) TargetType;
|
||||
std::vector<TargetType> fData;
|
||||
fData.resize(data.e.width);
|
||||
for (unsigned int i = 0; i < data.e.width; i++) {
|
||||
fData[i] = getNormalizedFloatType<T>(data.data[i]);
|
||||
}
|
||||
verifyMipmapLevel<TargetType, filterMode, addressMode>(
|
||||
texMipmap, fData.data(), data.e.width, level, offsetX);
|
||||
} else { // hipReadModeElementType == readMode
|
||||
verifyMipmapLevel<T, filterMode, addressMode>(
|
||||
texMipmap, data.data, data.e.width, level, offsetX);
|
||||
}
|
||||
HIP_CHECK(hipHostFree(data.data));
|
||||
memset(&data, 0, sizeof(data));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroyTextureObject(texMipmap));
|
||||
HIP_CHECK(hipFreeMipmappedArray(mipmapArray));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeElementType and hipFilterModePoint
|
||||
in 1D,
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj1D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 23") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeClamp>(23, 0.49);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeClamp>(67, -0.3);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeBorder>(131, 0.15);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeBorder>(263, 0.96);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeNormalizedFloat on integer types in 1D
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj1D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat", "",
|
||||
char, uchar, short, ushort,
|
||||
char1, uchar1, short1, ushort1,
|
||||
char2, uchar2, short2, ushort2,
|
||||
char4, uchar4, short4, ushort4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 23") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeClamp>(23, -0.9);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeClamp>(131, 0.15);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(67, -0.3);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(263, 0.13);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeBorder>(131, -0.34);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 23") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeBorder>(23, 0.4);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(263, 0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(67, -0.67);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeElementType and hipFilterModeLinear
|
||||
on float types in 1D,
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj1D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType float only", "",
|
||||
float, float1, float2, float4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 23, 0.") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 0.7);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 23") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, -0.67);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(263, 0.13);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(131, 0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(67, -0.97);
|
||||
}
|
||||
}
|
||||
@@ -19,6 +19,16 @@ THE SOFTWARE.
|
||||
|
||||
#pragma clang diagnostic ignored "-Wunused-parameter"
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_array_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_texture_helper.hh>
|
||||
#include <algorithm>
|
||||
|
||||
#pragma clang diagnostic ignored "-Wunused-variable"
|
||||
#pragma clang diagnostic ignored "-Wunused-parameter"
|
||||
|
||||
//#define TEST_TEXTURE // Only for float2
|
||||
static constexpr bool printLog = false; // Print log for debugging
|
||||
|
||||
/**
|
||||
* @addtogroup hipCreateTextureObject hipCreateTextureObject
|
||||
@@ -26,133 +36,469 @@ THE SOFTWARE.
|
||||
* @ingroup TextureTest
|
||||
*/
|
||||
|
||||
// Height Width Vector
|
||||
std::vector<unsigned int> hw_vector = {2048, 1024, 512, 256, 64};
|
||||
std::vector<unsigned int> mip_vector = {8, 4, 2, 1};
|
||||
|
||||
// MipMap is currently supported only on windows
|
||||
#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT))
|
||||
__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width,
|
||||
float level) {
|
||||
// Populate mipmap next level array
|
||||
template <typename T, hipTextureReadMode readMode>
|
||||
static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut,
|
||||
hipTextureObject_t texIn, unsigned int width, unsigned int height, T* data = nullptr) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2DLod<float>(textureObject, x, y, level);
|
||||
|
||||
float px = 1.0 / float(width);
|
||||
float py = 1.0 / float(height);
|
||||
|
||||
if (x < width && y < height) {
|
||||
if constexpr (readMode == hipReadModeElementType) {
|
||||
T t = tex2D<T>(texIn, x * px, y * py);
|
||||
#ifdef TEST_TEXTURE
|
||||
printf("(%d/%u, %d/%u):(%f, %f)\n", x, width, y, height, t.x, t.y);
|
||||
#endif
|
||||
surf2Dwrite<T>(t, surfOut, x * sizeof(T), y);
|
||||
if (data) data[y * width + x] = t; // record it for later verification
|
||||
}
|
||||
if constexpr (readMode == hipReadModeNormalizedFloat) {
|
||||
float4 t = tex2D<float4>(texIn, x * px, y * py);
|
||||
T tc = getTypeFromNormalizedFloat<T, float4>(t);
|
||||
surf2Dwrite<T>(tc, surfOut, x * sizeof(T), y);
|
||||
if (data) data[y * width + x] = tc;
|
||||
}
|
||||
// Users have freedom to use other methods to init level array
|
||||
// for example, use averge of surrounding pixes
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
static void runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) {
|
||||
INFO("Width: " << width << "Height: " << height << "mip: " << mipmap_level);
|
||||
template <typename T>
|
||||
static __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width,
|
||||
unsigned int height, float offsetX, float offsetY, float lod, T* data = nullptr) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
|
||||
// Create new width & height to be tested
|
||||
unsigned int orig_width = width;
|
||||
unsigned int orig_height = height;
|
||||
unsigned int i, j;
|
||||
width /= pow(2, mipmap_level);
|
||||
height /= pow(2, mipmap_level);
|
||||
unsigned int size = width * height * sizeof(float);
|
||||
float px = 1.0 / float(width);
|
||||
float py = 1.0 / float(height);
|
||||
|
||||
float* hData = reinterpret_cast<float*>(malloc(size));
|
||||
REQUIRE(hData != nullptr);
|
||||
memset(hData, 0, size);
|
||||
for (i = 0; i < height; i++) {
|
||||
for (j = 0; j < width; j++) {
|
||||
hData[i * width + j] = i * width + j;
|
||||
}
|
||||
if (x < width && y < height) {
|
||||
T t = tex2DLod<T>(texMipmap, (x + offsetX) * px, (y + offsetY) * py, lod);
|
||||
if (data) data[y * width + x] = t;
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
|
||||
HIP_ARRAY3D_DESCRIPTOR mipmapped_array_desc;
|
||||
memset(&mipmapped_array_desc, 0x00, sizeof(HIP_ARRAY3D_DESCRIPTOR));
|
||||
mipmapped_array_desc.Width = orig_width;
|
||||
mipmapped_array_desc.Height = orig_height;
|
||||
mipmapped_array_desc.Depth = 0;
|
||||
mipmapped_array_desc.Format = HIP_AD_FORMAT_FLOAT;
|
||||
mipmapped_array_desc.NumChannels =
|
||||
((channelDesc.x != 0) + (channelDesc.y != 0) + (channelDesc.z != 0) + (channelDesc.w != 0));
|
||||
mipmapped_array_desc.Flags = 0;
|
||||
template <typename T, hipTextureReadMode readMode = hipReadModeElementType,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void populateMipmaps(hipMipmappedArray_t mipmapArray, hipExtent size,
|
||||
std::vector<mipmapLevelArray<T>>& mipmapData) {
|
||||
size_t width = size.width;
|
||||
size_t height = size.height;
|
||||
|
||||
hipMipmappedArray* mip_array_ptr;
|
||||
HIP_CHECK(hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc, 2 * mipmap_level));
|
||||
unsigned int level = 0;
|
||||
|
||||
hipArray* hipArray = nullptr;
|
||||
HIP_CHECK(hipMipmappedArrayGetLevel(&hipArray, mip_array_ptr, mipmap_level));
|
||||
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
|
||||
while (width != 1 || height != 1) {
|
||||
hipArray_t levelArray = nullptr, nextLevelArray = nullptr;
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&levelArray, mipmapArray, level));
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&nextLevelArray, mipmapArray, level + 1));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
hipExtent levelArraySize {0, 0, 0};
|
||||
HIP_CHECK(hipArrayGetInfo(nullptr, &levelArraySize, nullptr, levelArray));
|
||||
if(levelArraySize.width != width || levelArraySize.height != height) {
|
||||
fprintf(stderr, "Level %u: size (%zu, %zu, %zu) != Expected size (%zu, %zu, 0)\n", level,
|
||||
levelArraySize.width, levelArraySize.height, levelArraySize.depth,
|
||||
width, height);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
// Specify texture object parameters
|
||||
hipTextureDesc texDesc;
|
||||
memset(&texDesc, 0, sizeof(texDesc));
|
||||
texDesc.addressMode[0] = hipAddressModeWrap;
|
||||
texDesc.addressMode[1] = hipAddressModeWrap;
|
||||
texDesc.filterMode = hipFilterModePoint;
|
||||
texDesc.readMode = hipReadModeElementType;
|
||||
texDesc.normalizedCoords = 0;
|
||||
width = width >> 1 ? width >> 1: 1;
|
||||
height = height >> 1 ? height >> 1 : 1;
|
||||
|
||||
// Create texture object
|
||||
hipTextureObject_t textureObject = 0;
|
||||
HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, nullptr));
|
||||
hipExtent nextLevelArraySize {0, 0, 0};
|
||||
HIP_CHECK(hipArrayGetInfo(nullptr, &nextLevelArraySize, nullptr, nextLevelArray));
|
||||
if(nextLevelArraySize.width != width || nextLevelArraySize.height != height) {
|
||||
fprintf(stderr, "Next level %u: size (%zu, %zu, %zu) != Expected size (%zu, %zu, 0)\n",
|
||||
level + 1, nextLevelArraySize.width, nextLevelArraySize.height,
|
||||
nextLevelArraySize.depth, width, height);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
float* dData = nullptr;
|
||||
HIP_CHECK(hipMalloc(&dData, size));
|
||||
REQUIRE(dData != nullptr);
|
||||
hipTextureObject_t texIn;
|
||||
hipResourceDesc texRes;
|
||||
memset(&texRes, 0, sizeof(hipResourceDesc));
|
||||
texRes.resType = hipResourceTypeArray;
|
||||
texRes.res.array.array = levelArray;
|
||||
|
||||
dim3 dimBlock(16, 16, 1);
|
||||
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
|
||||
hipTextureDesc texDescr;
|
||||
memset(&texDescr, 0, sizeof(hipTextureDesc));
|
||||
texDescr.normalizedCoords = 1; // To populate next level array smoothly
|
||||
texDescr.filterMode = filterMode;
|
||||
texDescr.addressMode[0] = addressMode;
|
||||
texDescr.addressMode[1] = addressMode;
|
||||
texDescr.addressMode[2] = addressMode;
|
||||
texDescr.readMode = readMode;
|
||||
HIP_CHECK(hipCreateTextureObject(&texIn, &texRes, &texDescr, NULL));
|
||||
|
||||
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, width,
|
||||
(2 * mipmap_level));
|
||||
HIP_CHECK(hipGetLastError());
|
||||
hipSurfaceObject_t surfOut;
|
||||
hipResourceDesc surfRes;
|
||||
memset(&surfRes, 0, sizeof(hipResourceDesc));
|
||||
surfRes.resType = hipResourceTypeArray;
|
||||
surfRes.res.array.array = nextLevelArray;
|
||||
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfOut, &surfRes));
|
||||
size_t size = width * height * sizeof(T);
|
||||
mipmapLevelArray<T> data{nullptr, {width, height, 0}};
|
||||
HIP_CHECK(hipHostMalloc((void**)&data.data, size));
|
||||
memset(data.data, 0, size);
|
||||
|
||||
dim3 blockSize(16, 16, 1);
|
||||
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
|
||||
(height + blockSize.y - 1) / blockSize.y, 1);
|
||||
|
||||
populateMipmapNextLevelArray<T, readMode>
|
||||
<<<gridSize, blockSize>>>(surfOut, texIn, width, height, data.data);
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfOut));
|
||||
HIP_CHECK(hipDestroyTextureObject(texIn));
|
||||
HIP_CHECK(hipFreeArray(levelArray));
|
||||
HIP_CHECK(hipFreeArray(nextLevelArray));
|
||||
mipmapData.push_back(data); // For later verification
|
||||
level++;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void verifyMipmapLevel(hipTextureObject_t texMipmap, T* data, size_t width, size_t height,
|
||||
float level, float offsetX, float offsetY) {
|
||||
T* hOutput = nullptr;
|
||||
size_t size = width * height * sizeof(T);
|
||||
HIP_CHECK(hipHostMalloc((void**)&hOutput, size));
|
||||
memset(hOutput, 0, size);
|
||||
|
||||
dim3 blockSize(16, 16, 1);
|
||||
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
|
||||
(height + blockSize.y - 1) / blockSize.y, 1);
|
||||
|
||||
getMipmap<T><<<gridSize, blockSize>>>(texMipmap, width, height, offsetX,
|
||||
offsetY, level, hOutput);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
float* hOutputData = reinterpret_cast<float*>(malloc(size));
|
||||
REQUIRE(hOutputData != nullptr);
|
||||
memset(hOutputData, 0, size);
|
||||
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
|
||||
|
||||
for (i = 0; i < height; i++) {
|
||||
for (j = 0; j < width; j++) {
|
||||
if (hData[i * width + j] != hOutputData[i * width + j]) {
|
||||
INFO("Difference found at [ " << i << j << " ]: " << hData[i * width + j]
|
||||
<< hOutputData[i * width + j]);
|
||||
for (unsigned int j = 0; j < height; j++) {
|
||||
for (unsigned int i = 0; i < width; i++) {
|
||||
unsigned int index = j * width + i;
|
||||
T cpuExpected = getExpectedValue<T, addressMode, filterMode, false>(
|
||||
width, height, i + offsetX, j + offsetY, data);
|
||||
T gpuOutput = hOutput[index];
|
||||
bool match = hipTextureSamplingVerify<T, filterMode, false>(gpuOutput, cpuExpected);
|
||||
if (!match) {
|
||||
WARN("Mismatch at (level "
|
||||
<< level << ": " << i << ", " << j << " -> " << (i + offsetX) << ", " << (j + offsetY)
|
||||
<< ") GPU output : " << getString(gpuOutput)
|
||||
<< " CPU expected: " << getString(cpuExpected) << ", data[" << index
|
||||
<< "]:" << getString(data[index]) << "\n");
|
||||
REQUIRE(false);
|
||||
} else if (printLog) {
|
||||
WARN("Matching at (level "
|
||||
<< level << ": " << i << ", " << j << " -> " << (i + offsetX) << ", " << (j + offsetY)
|
||||
<< ") GPU output : " << getString(gpuOutput)
|
||||
<< " CPU expected: " << getString(cpuExpected) << ", data[" << index
|
||||
<< "]:" << getString(data[index]) << "\n");
|
||||
}
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipDestroyTextureObject(textureObject));
|
||||
HIP_CHECK(hipFree(dData));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
free(hData);
|
||||
HIP_CHECK(hipHostFree(hOutput));
|
||||
}
|
||||
|
||||
template <typename T, hipTextureReadMode readMode = hipReadModeElementType,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void testMipmapTextureObj(size_t width, size_t height, float offsetX = 0.,
|
||||
float offsetY = 0.) {
|
||||
std::vector<mipmapLevelArray<T>> mipmapData;
|
||||
size_t size = width * height * sizeof(T);
|
||||
mipmapLevelArray<T> data{nullptr, {width, height, 0}};
|
||||
HIP_CHECK(hipHostMalloc((void**)&data.data, size));
|
||||
memset(data.data, 0, size);
|
||||
for (int j = 0; j < height; j++)
|
||||
{
|
||||
for (int i = 0; i < width; i++)
|
||||
{
|
||||
if (isFloat<T>() && filterMode == hipFilterModeLinear) {
|
||||
/*
|
||||
* For linear sampling of images, the GPU does not use IEEE floating point types, it uses
|
||||
* lower-precision sampling optimized formats. Also, those formats often change between GPU
|
||||
* generations. So counting on IEEE precision and accuracy when doing linear sampling
|
||||
* is mistaken. To workaround this issue, we can initialize float pixels on a retively
|
||||
* smoothy surface.
|
||||
*/
|
||||
data.data[j * width + i] =
|
||||
T( float(i) * (float(i) - width + 1) * float(j) * (float(j) - height + 1) );
|
||||
} else {
|
||||
initVal(data.data[j * width + i]); // Randomize initial values
|
||||
}
|
||||
}
|
||||
}
|
||||
mipmapData.push_back(data); // record level 0 data for later verification
|
||||
|
||||
// Get the max mipmap levels in terms of image size
|
||||
const unsigned int maxLevels = 1 + std::log2(std::max(width, height));
|
||||
|
||||
// create mipmap array
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
|
||||
hipMipmappedArray_t mipmapArray = nullptr;
|
||||
hipExtent extent { width, height, 0 };
|
||||
HIP_CHECK(hipMallocMipmappedArray(&mipmapArray, &desc, extent, maxLevels));
|
||||
|
||||
// Initialize level 0
|
||||
hipArray_t levelArray;
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&levelArray, mipmapArray, 0));
|
||||
hipMemcpy3DParms copyParams{};
|
||||
copyParams.srcPtr = make_hipPitchedPtr(data.data, width * sizeof(T),
|
||||
width, height);
|
||||
copyParams.dstArray = levelArray;
|
||||
copyParams.extent.width = width;
|
||||
copyParams.extent.height = height;
|
||||
copyParams.extent.depth = 1;
|
||||
copyParams.kind = hipMemcpyHostToDevice;
|
||||
HIP_CHECK(hipMemcpy3D(©Params));
|
||||
|
||||
// Populate other mipmap levels based on level 0
|
||||
populateMipmaps<T, readMode, filterMode, addressMode>(mipmapArray, extent, mipmapData);
|
||||
|
||||
if(maxLevels != mipmapData.size()) {
|
||||
fprintf(stderr, "maxLevels %u != mipmapData.size() %zu\n", maxLevels, mipmapData.size());
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
hipResourceDesc resDescr;
|
||||
memset(&resDescr, 0, sizeof(hipResourceDesc));
|
||||
resDescr.resType = hipResourceTypeMipmappedArray; // For mipmap texture
|
||||
resDescr.res.mipmap.mipmap = mipmapArray;
|
||||
|
||||
hipTextureDesc texDescr;
|
||||
memset(&texDescr, 0, sizeof(hipTextureDesc));
|
||||
|
||||
texDescr.normalizedCoords = 1; // normalizedCoords must be 1 for mipmap array
|
||||
texDescr.filterMode = filterMode;
|
||||
texDescr.mipmapFilterMode = filterMode;
|
||||
texDescr.addressMode[0] = addressMode;
|
||||
texDescr.addressMode[1] = addressMode;
|
||||
texDescr.addressMode[2] = addressMode;
|
||||
// Ignored in AMD Hw sampler SRD, but cuda need it
|
||||
texDescr.maxMipmapLevelClamp = float(maxLevels - 1);
|
||||
texDescr.readMode = readMode;
|
||||
|
||||
hipTextureObject_t texMipmap = nullptr;
|
||||
HIP_CHECK(hipCreateTextureObject(&texMipmap, &resDescr, &texDescr, NULL));
|
||||
|
||||
for(unsigned int level = 0; level < mipmapData.size(); level++){
|
||||
mipmapLevelArray<T> &data = mipmapData.at(level);
|
||||
|
||||
if constexpr (hipReadModeNormalizedFloat == readMode) {
|
||||
typedef decltype(getNormalizedFloatType(*data.data)) TargetType;
|
||||
std::vector<TargetType> fData;
|
||||
fData.resize(data.e.width * data.e.height);
|
||||
for (unsigned int j = 0; j < data.e.height; j++) {
|
||||
for (unsigned int i = 0; i < data.e.width; i++) {
|
||||
unsigned int index = j * data.e.width + i;
|
||||
fData[index] = getNormalizedFloatType<T>(data.data[index]);
|
||||
}
|
||||
}
|
||||
verifyMipmapLevel<TargetType, filterMode, addressMode>(
|
||||
texMipmap, fData.data(), data.e.width, data.e.height, level, offsetX, offsetY);
|
||||
|
||||
} else { // hipReadModeElementType == readMode
|
||||
verifyMipmapLevel<T, filterMode, addressMode>(
|
||||
texMipmap, data.data, data.e.width, data.e.height, level, offsetX, offsetY);
|
||||
}
|
||||
HIP_CHECK(hipHostFree(data.data));
|
||||
memset(&data, 0, sizeof(data));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroyTextureObject(texMipmap));
|
||||
HIP_CHECK(hipFreeMipmappedArray(mipmapArray));
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test where resource type is a mipmapped array.
|
||||
* - The suite will test following functions with hipReadModeElementType and hipFilterModePoint
|
||||
in 2D
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - unit/texture/hipTextureMipmapObj2D.cc
|
||||
* - catch\unit\texture\hipTextureMipmapObj2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.2
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEST_CASE("Unit_hipTextureMipmapObj2D_Check") {
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
|
||||
#ifdef _WIN32
|
||||
for (auto& hw : hw_vector) {
|
||||
for (auto& mip : mip_vector) {
|
||||
if ((hw / static_cast<int>(pow(2, (mip * 2)))) > 0) {
|
||||
runMipMapTest(hw, hw, mip);
|
||||
}
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 23, 21") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeClamp>(23, 21, 0.4, -0.9);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 67, 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeClamp>(67, 131, -0.3, -0.67);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 131, 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeBorder>(131, 263, 0.15, -0.34);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 263, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeBorder>(263, 67, 0.13, 0.96);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeNormalizedFloat on integer types in 2D
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat", "",
|
||||
char, uchar, short, ushort,
|
||||
char1, uchar1, short1, ushort1,
|
||||
char2, uchar2, short2, ushort2,
|
||||
char4, uchar4, short4, ushort4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 23, 21") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeClamp>(23, 21, 0.4, -0.9);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 131, 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeClamp>(131, 263, 0.15, -0.34);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 67, 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 17, -0.3, -0.67);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 263, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(263, 67, 0.13, 0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 131, 263") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeBorder>(131, 263, 0.15, -0.34);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 23, 21") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeBorder>(23, 21, 0.4, -0.9);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 263, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(263, 67, 0.13, 0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 67, 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(67, 131, -0.3, -0.67);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeElementType and hipFilterModeLinear
|
||||
on float types in 2D,
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj2D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType float only", "",
|
||||
float, float1, float2, float4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 23, 17, 0., 0.") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 17, 0.79, 0.37);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 23, 17") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 17, -0.3, -0.67);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 263, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(263, 67, 0.13, 0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 263, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(263, 67, 0.13, 0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 67, 131") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(67, 131, -0.3, -0.67);
|
||||
}
|
||||
#else
|
||||
SUCCEED("Mipmaps are Supported only on windows, skipping the test.");
|
||||
#endif
|
||||
}
|
||||
|
||||
@@ -0,0 +1,522 @@
|
||||
/*
|
||||
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 <hip_array_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_texture_helper.hh>
|
||||
#include <algorithm>
|
||||
|
||||
#pragma clang diagnostic ignored "-Wunused-variable"
|
||||
#pragma clang diagnostic ignored "-Wunused-parameter"
|
||||
|
||||
//#define TEST_TEXTURE
|
||||
static constexpr bool printLog = false; // Print log for debugging
|
||||
|
||||
// Populate mipmap next level array
|
||||
template <typename T, hipTextureReadMode readMode>
|
||||
static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, hipTextureObject_t texIn,
|
||||
unsigned int width, unsigned int height, unsigned int depth, T* data = nullptr) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
float px = 1.0 / float(width);
|
||||
float py = 1.0 / float(height);
|
||||
float pz = 1.0 / float(depth);
|
||||
|
||||
if (x < width && y < height && z < depth) {
|
||||
if constexpr (readMode == hipReadModeElementType) {
|
||||
T t = tex3D<T>(texIn, x * px, y * py, z * pz);
|
||||
#ifdef TEST_TEXTURE
|
||||
printf("(%d/%u, %d/%u):(%f, %f)\n", x, width, y, height, t.x, t.y);
|
||||
#endif
|
||||
surf3Dwrite<T>(t, surfOut, x * sizeof(T), y, z);
|
||||
if (data) data[z* width * height + y * width + x] = t; // record it for later verification
|
||||
}
|
||||
if constexpr (readMode == hipReadModeNormalizedFloat) {
|
||||
float4 t = tex3D<float4>(texIn, x * px, y * py, z * pz);
|
||||
T tc = getTypeFromNormalizedFloat<T, float4>(t);
|
||||
surf3Dwrite<T>(tc, surfOut, x * sizeof(T), y, z);
|
||||
if (data) data[z * width * height + y * width + x] = tc;
|
||||
#ifdef TEST_TEXTURE
|
||||
printf(
|
||||
"populateMipmapNextLevelArray(%d/%u, %d/%u, %d/%u->%d): t.x=%f, t.y=%f, t.z=%f, t.w=%f, "
|
||||
"tc=%d\n",
|
||||
x, width, y, height, z, depth, z * width * height + y * width + x, t.x, t.y, t.z, t.w,
|
||||
(int)tc);
|
||||
#endif
|
||||
}
|
||||
// Users have freedom to use other methods to init level array
|
||||
// for example, use averge of surrounding pixes
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width,
|
||||
unsigned int height,
|
||||
unsigned int depth, float offsetX, float offsetY, float offsetZ, float lod,
|
||||
T* data = nullptr) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
float px = 1.0 / float(width);
|
||||
float py = 1.0 / float(height);
|
||||
float pz = 1.0 / float(depth);
|
||||
|
||||
if (x < width && y < height && z < depth) {
|
||||
T t = tex3DLod<T>(texMipmap, (x + offsetX) * px, (y + offsetY) * py, (z + offsetZ) * pz, lod);
|
||||
if (data) data[z * width * height + y * width + x] = t;
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T, hipTextureReadMode readMode = hipReadModeElementType,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void populateMipmaps(hipMipmappedArray_t mipmapArray, hipExtent size,
|
||||
std::vector<mipmapLevelArray<T>>& mipmapData) {
|
||||
size_t width = size.width;
|
||||
size_t height = size.height;
|
||||
size_t depth = size.depth;
|
||||
unsigned int level = 0;
|
||||
|
||||
while (width != 1 || height != 1 || depth != 1) {
|
||||
hipArray_t levelArray = nullptr, nextLevelArray = nullptr;
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&levelArray, mipmapArray, level));
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&nextLevelArray, mipmapArray, level + 1));
|
||||
|
||||
hipExtent levelArraySize {0, 0, 0};
|
||||
HIP_CHECK(hipArrayGetInfo(nullptr, &levelArraySize, nullptr, levelArray));
|
||||
if (levelArraySize.width != width || levelArraySize.height != height ||
|
||||
levelArraySize.depth != depth) {
|
||||
fprintf(stderr, "Level %u: size (%zu, %zu, %zu) != Expected size (%zu, %zu, %zu)\n", level,
|
||||
levelArraySize.width, levelArraySize.height, levelArraySize.depth, width, height,
|
||||
depth);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
width = width >> 1 ? width >> 1: 1;
|
||||
height = height >> 1 ? height >> 1 : 1;
|
||||
depth = depth >> 1 ? depth >> 1 : 1;
|
||||
|
||||
hipExtent nextLevelArraySize {0, 0, 0};
|
||||
HIP_CHECK(hipArrayGetInfo(nullptr, &nextLevelArraySize, nullptr, nextLevelArray));
|
||||
if (nextLevelArraySize.width != width || nextLevelArraySize.height != height ||
|
||||
nextLevelArraySize.depth != depth) {
|
||||
fprintf(stderr, "Next level %u: size (%zu, %zu, %zu) != Expected size (%zu, %zu, %zu)\n",
|
||||
level + 1, nextLevelArraySize.width, nextLevelArraySize.height, nextLevelArraySize.depth,
|
||||
width, height, depth);
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
hipTextureObject_t texIn;
|
||||
hipResourceDesc texRes;
|
||||
memset(&texRes, 0, sizeof(hipResourceDesc));
|
||||
texRes.resType = hipResourceTypeArray;
|
||||
texRes.res.array.array = levelArray;
|
||||
|
||||
hipTextureDesc texDescr;
|
||||
memset(&texDescr, 0, sizeof(hipTextureDesc));
|
||||
texDescr.normalizedCoords = 1; // To populate next level array smoothly
|
||||
texDescr.filterMode = filterMode;
|
||||
texDescr.addressMode[0] = addressMode;
|
||||
texDescr.addressMode[1] = addressMode;
|
||||
texDescr.addressMode[2] = addressMode;
|
||||
texDescr.readMode = readMode;
|
||||
HIP_CHECK(hipCreateTextureObject(&texIn, &texRes, &texDescr, NULL));
|
||||
|
||||
hipSurfaceObject_t surfOut;
|
||||
hipResourceDesc surfRes;
|
||||
memset(&surfRes, 0, sizeof(hipResourceDesc));
|
||||
surfRes.resType = hipResourceTypeArray;
|
||||
surfRes.res.array.array = nextLevelArray;
|
||||
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfOut, &surfRes));
|
||||
size_t size = width * height * depth * sizeof(T);
|
||||
mipmapLevelArray<T> data{nullptr, {width, height, depth}};
|
||||
HIP_CHECK(hipHostMalloc((void**)&data.data, size));
|
||||
memset(data.data, 0, size);
|
||||
|
||||
dim3 blockSize(16, 16, 4);
|
||||
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
|
||||
(height + blockSize.y - 1) / blockSize.y, (depth + blockSize.z - 1) / blockSize.z);
|
||||
|
||||
populateMipmapNextLevelArray<T, readMode>
|
||||
<<<gridSize, blockSize>>>(surfOut, texIn, width, height, depth, data.data);
|
||||
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfOut));
|
||||
HIP_CHECK(hipDestroyTextureObject(texIn));
|
||||
HIP_CHECK(hipFreeArray(levelArray));
|
||||
HIP_CHECK(hipFreeArray(nextLevelArray));
|
||||
mipmapData.push_back(data); // For later verification
|
||||
level++;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void verifyMipmapLevel(hipTextureObject_t texMipmap, T* data, size_t width, size_t height,
|
||||
size_t depth, float level, float offsetX, float offsetY, float offsetZ) {
|
||||
T* hOutput = nullptr;
|
||||
size_t size = width * height * depth * sizeof(T);
|
||||
HIP_CHECK(hipHostMalloc((void**)&hOutput, size));
|
||||
memset(hOutput, 0, size);
|
||||
|
||||
dim3 blockSize(16, 16, 4);
|
||||
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
|
||||
(height + blockSize.y - 1) / blockSize.y, (depth + blockSize.z - 1) / blockSize.z);
|
||||
|
||||
getMipmap<T><<<gridSize, blockSize>>>(texMipmap, width, height, depth, offsetX,
|
||||
offsetY, offsetZ, level, hOutput);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
for (unsigned int k = 0; k < depth; k++) {
|
||||
for (unsigned int j = 0; j < height; j++) {
|
||||
for (unsigned int i = 0; i < width; i++) {
|
||||
unsigned int index = k * width * height + j * width + i;
|
||||
T cpuExpected = getExpectedValue<T, addressMode, filterMode, false>(
|
||||
width, height, depth, i + offsetX, j + offsetY, k + offsetZ, data);
|
||||
T gpuOutput = hOutput[index];
|
||||
bool match = hipTextureSamplingVerify<T, filterMode, false>(gpuOutput, cpuExpected);
|
||||
if (!match) {
|
||||
WARN("Mismatch at (level " << level << ": " << i << ", " << j << ", " << k << "->"
|
||||
<< (i + offsetX) << ", " << (j + offsetY) << ", "
|
||||
<< (k + offsetZ) << ") GPU output : " << getString(gpuOutput)
|
||||
<< " CPU expected: " << getString(cpuExpected) << ", data["
|
||||
<< index << "]:" << getString(data[index]) << "\n");
|
||||
REQUIRE(false);
|
||||
} else if (printLog) {
|
||||
WARN("Matching at (level " << level << ": " << i << ", " << j << ", " << k << "->"
|
||||
<< (i + offsetX) << ", " << (j + offsetY) << ", "
|
||||
<< (k + offsetZ) << ") GPU output : " << getString(gpuOutput)
|
||||
<< " CPU expected: " << getString(cpuExpected) << ", data["
|
||||
<< index << "]:" << getString(data[index]) << "\n");
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipHostFree(hOutput));
|
||||
}
|
||||
|
||||
template <typename T, hipTextureReadMode readMode = hipReadModeElementType,
|
||||
hipTextureFilterMode filterMode = hipFilterModePoint,
|
||||
hipTextureAddressMode addressMode = hipAddressModeClamp>
|
||||
static void testMipmapTextureObj(size_t width, size_t height, size_t depth, float offsetX = 0.,
|
||||
float offsetY = 0., float offsetZ = 0.) {
|
||||
std::vector<mipmapLevelArray<T>> mipmapData;
|
||||
size_t size = width * height * depth * sizeof(T);
|
||||
mipmapLevelArray<T> data{nullptr, {width, height, depth}};
|
||||
HIP_CHECK(hipHostMalloc((void**)&data.data, size));
|
||||
memset(data.data, 0, size);
|
||||
for (int k = 0; k < depth; k++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int i = 0; i < width; i++) {
|
||||
if (isFloat<T>() && filterMode == hipFilterModeLinear) {
|
||||
/*
|
||||
* For linear sampling of images, the GPU does not use IEEE floating point types, it
|
||||
* uses lower-precision sampling optimized formats. Also, those formats often change
|
||||
* between GPU generations. So counting on IEEE precision and accuracy when doing linear
|
||||
* sampling is mistaken. To workaround this issue, we can initialize float pixels on a
|
||||
* retively smoothy surface.
|
||||
*/
|
||||
data.data[k * width * height + j * width + i] =
|
||||
T(float(i) * (float(i) - width + 1) * float(j) * (float(j) - height + 1) * float(k) *
|
||||
(float(k) - depth + 1));
|
||||
} else {
|
||||
initVal(data.data[k * width * height + j * width + i]); // Randomize initial values
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
mipmapData.push_back(data); // record level 0 data for later verification
|
||||
|
||||
// Get the max mipmap levels in terms of image size
|
||||
const unsigned int maxLevels = 1 + std::log2(std::max(width, std::max(height, depth)));
|
||||
|
||||
// create mipmap array
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
|
||||
hipMipmappedArray_t mipmapArray = nullptr;
|
||||
hipExtent extent { width, height, depth };
|
||||
HIP_CHECK(hipMallocMipmappedArray(&mipmapArray, &desc, extent, maxLevels));
|
||||
|
||||
// Initialize level 0
|
||||
hipArray_t levelArray;
|
||||
HIP_CHECK(hipGetMipmappedArrayLevel(&levelArray, mipmapArray, 0));
|
||||
hipMemcpy3DParms copyParams{};
|
||||
copyParams.srcPtr = make_hipPitchedPtr(data.data, width * sizeof(T),
|
||||
width, height);
|
||||
copyParams.dstArray = levelArray;
|
||||
copyParams.extent.width = width;
|
||||
copyParams.extent.height = height;
|
||||
copyParams.extent.depth = depth;
|
||||
copyParams.kind = hipMemcpyHostToDevice;
|
||||
HIP_CHECK(hipMemcpy3D(©Params));
|
||||
|
||||
// Populate other mipmap levels based on level 0
|
||||
populateMipmaps<T, readMode, filterMode, addressMode>(mipmapArray, extent, mipmapData);
|
||||
|
||||
if(maxLevels != mipmapData.size()) {
|
||||
fprintf(stderr, "maxLevels %u != mipmapData.size() %zu\n", maxLevels, mipmapData.size());
|
||||
REQUIRE(false);
|
||||
}
|
||||
hipResourceDesc resDescr;
|
||||
memset(&resDescr, 0, sizeof(hipResourceDesc));
|
||||
resDescr.resType = hipResourceTypeMipmappedArray; // For mipmap texture
|
||||
resDescr.res.mipmap.mipmap = mipmapArray;
|
||||
|
||||
hipTextureDesc texDescr;
|
||||
memset(&texDescr, 0, sizeof(hipTextureDesc));
|
||||
|
||||
texDescr.normalizedCoords = 1; // normalizedCoords must be 1 for mipmap array
|
||||
texDescr.filterMode = filterMode;
|
||||
texDescr.mipmapFilterMode = filterMode;
|
||||
texDescr.addressMode[0] = addressMode;
|
||||
texDescr.addressMode[1] = addressMode;
|
||||
texDescr.addressMode[2] = addressMode;
|
||||
// Ignored in AMD Hw sampler SRD, but cuda need it
|
||||
texDescr.maxMipmapLevelClamp = float(maxLevels - 1);
|
||||
texDescr.readMode = readMode;
|
||||
|
||||
hipTextureObject_t texMipmap = nullptr;
|
||||
HIP_CHECK(hipCreateTextureObject(&texMipmap, &resDescr, &texDescr, NULL));
|
||||
|
||||
for(unsigned int level = 0; level < mipmapData.size(); level++){
|
||||
mipmapLevelArray<T> &data = mipmapData.at(level);
|
||||
|
||||
if constexpr (hipReadModeNormalizedFloat == readMode) {
|
||||
typedef decltype(getNormalizedFloatType(*data.data)) TargetType;
|
||||
std::vector<TargetType> fData;
|
||||
fData.resize(data.e.width * data.e.height * data.e.depth);
|
||||
for (unsigned int k = 0; k < data.e.depth; k++) {
|
||||
for (unsigned int j = 0; j < data.e.height; j++) {
|
||||
for (unsigned int i = 0; i < data.e.width; i++) {
|
||||
unsigned int index = k * data.e.width * data.e.height + j * data.e.width + i;
|
||||
fData[index] = getNormalizedFloatType<T>(data.data[index]);
|
||||
#ifdef TEST_TEXTURE
|
||||
fprintf(stderr, "level=%u, (%u/%zu, %u/%zu, %u/%zu-->%u): data.data=%s, fData=%s\n",
|
||||
level, i, data.e.width, j, data.e.height, k, data.e.depth, index,
|
||||
getString(data.data[index]).c_str(), getString(fData[index]).c_str());
|
||||
#endif
|
||||
}
|
||||
}
|
||||
}
|
||||
verifyMipmapLevel<TargetType, filterMode, addressMode>(
|
||||
texMipmap, fData.data(), data.e.width, data.e.height, data.e.depth,
|
||||
level, offsetX, offsetY, offsetZ);
|
||||
|
||||
} else { // hipReadModeElementType == readMode
|
||||
verifyMipmapLevel<T, filterMode, addressMode>(
|
||||
texMipmap, data.data, data.e.width, data.e.height, data.e.depth, level,
|
||||
offsetX, offsetY, offsetZ);
|
||||
}
|
||||
HIP_CHECK(hipHostFree(data.data));
|
||||
memset(&data, 0, sizeof(data));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroyTextureObject(texMipmap));
|
||||
HIP_CHECK(hipFreeMipmappedArray(mipmapArray));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeElementType and hipFilterModePoint
|
||||
in 3D
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 23, 21, 47") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeClamp>(23, 21, 47, 0.4, -0.9, 0.77);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 67, 131, 99") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeClamp>(67, 131, 99, -0.3, -0.67, 0.49);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 131, 263, 31") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeBorder>(131, 263, 31, 0.15, -0.34, 0.85);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 263, 67, 17") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModePoint,
|
||||
hipAddressModeBorder>(263, 67, 17, 0.13, 0.96, -0.57);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeNormalizedFloat on integer types in 3D
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat", "",
|
||||
char, uchar, short, ushort,
|
||||
char1, uchar1, short1, ushort1,
|
||||
char2, uchar2, short2, ushort2,
|
||||
char4, uchar4, short4, ushort4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 23, 21, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeClamp>(23, 21, 67, 0.4, -0.9, 0.37);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeClamp 131, 263, 11") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeClamp>(131, 263, 11, 0.15, -0.34, 0.83);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 67, 131, 53") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 17, 53, -0.3, -0.67, 0.78);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 263, 67, 37") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(263, 67, 37, 0.13, 0.96, -0.96);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 131, 263, 11") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeBorder>(131, 263, 11, 0.15, -0.34, -0.11);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
|
||||
"hipAddressModeBorder 23, 21, 201") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModePoint,
|
||||
hipAddressModeBorder>(23, 21, 201, 0.4, -0.9, 0.54);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 263, 67, 51") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(263, 67, 51, 0.13, 0.96, 0.77);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 67, 131, 87") {
|
||||
testMipmapTextureObj<TestType, hipReadModeNormalizedFloat, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(67, 131, 87, -0.3, -0.67, -0.29);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - The suite will test following functions with hipReadModeElementType and hipFilterModeLinear
|
||||
on float types in 3D,
|
||||
creating MipMap array,
|
||||
getting level array,
|
||||
creating/initilizing texture and surface objects on level array,
|
||||
creating texture object on the mipmap array,
|
||||
verifing the texture object
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\texture\hipTextureMipmapObj3D.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - Host specific (WINDOWS)
|
||||
* - Textures supported on device
|
||||
* - HIP_VERSION >= 5.7
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType float only", "",
|
||||
float, float1, float2, float4) {
|
||||
CHECK_IMAGE_SUPPORT
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 23, 17, 301, 0., 0., 0.") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 17, 301, 0.19, 0.46, -0.9);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 23, 17, 243") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(23, 17, 243, -0.3, -0.67, 0.65);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeClamp 263, 67, 39") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeClamp>(263, 67, 39, 0.13, 0.96, 0.66);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 263, 67, 117") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(263, 67, 117, 0.13, 0.96, -0.93);
|
||||
}
|
||||
SECTION(
|
||||
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, "
|
||||
"hipAddressModeBorder 67, 131, 67") {
|
||||
testMipmapTextureObj<TestType, hipReadModeElementType, hipFilterModeLinear,
|
||||
hipAddressModeBorder>(67, 131, 67, -0.3, -0.67, 0.88);
|
||||
}
|
||||
}
|
||||
مرجع در شماره جدید
Block a user