diff --git a/projects/hip-tests/catch/include/hip_array_common.hh b/projects/hip-tests/catch/include/hip_array_common.hh index a0b5d96622..e9b438cf93 100644 --- a/projects/hip-tests/catch/include/hip_array_common.hh +++ b/projects/hip-tests/catch/include/hip_array_common.hh @@ -128,7 +128,12 @@ template< const std::string getString(const T& t) { std::ostringstream os; - os<< "(" << t.x << ")"; + if constexpr (std::is_same::value || + std::is_same::value) { + os << "(" << static_cast(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::value || + std::is_same::value) { + os << "(" << static_cast(t.x) << ", " << static_cast(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::value || + std::is_same::value) { + os << "(" << static_cast(t.x) << ", " << static_cast(t.y) << ", " << + static_cast(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::value || + std::is_same::value) { + os << "(" << static_cast(t.x) << ", " << static_cast(t.y) << ", " << + static_cast(t.z) << ", " << static_cast(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::value || + std::is_same::value) { + os << static_cast(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(std::numeric_limits < T > ::max() * r); + if constexpr (std::is_floating_point::value) { + // Restrict any float within (-1000, 1000) + // to prevent too big float value that would make caculation sick + return static_cast(r * 1000.); + } else { + return static_cast(std::numeric_limits::max() * r); + } } template< @@ -214,3 +247,62 @@ template< static inline void initVal(T &val) { val = getRandom(); } + +/*Convert normalized floatx to typex*/ +template inline __device__ T getTypeFromNormalizedFloat(const F &f) { + T t; + if constexpr (std::is_scalar::value) + t = static_cast(f.x * std::numeric_limits::max()); + else { + if constexpr (rank() > 0) + t.x = static_cast(f.x * std::numeric_limits::max()); + if constexpr (rank() > 1) + t.y = static_cast(f.y * std::numeric_limits::max()); + if constexpr (rank() > 2) + t.z = static_cast(f.z * std::numeric_limits::max()); + if constexpr (rank() > 3) + t.w = static_cast(f.w * std::numeric_limits::max()); + } + return t; +} + +/*Convert typex to normalized floatx*/ +template +inline auto getNormalizedFloatType(const T &t) { + if constexpr (std::is_scalar::value) + return static_cast(t) / std::numeric_limits::max(); + else { + if constexpr (rank() == 1) { + float1 f{static_cast(t.x) / std::numeric_limits::max()}; + return f; + } + if constexpr (rank() == 2) { + float2 f{static_cast(t.x) / std::numeric_limits::max(), + static_cast(t.y) / std::numeric_limits::max()}; + return f; + } + if constexpr (rank() == 3) { + float3 f{static_cast(t.x) / std::numeric_limits::max(), + static_cast(t.y) / std::numeric_limits::max(), + static_cast(t.z) / std::numeric_limits::max()}; + return f; + } + if constexpr (rank() == 4) { + float4 f{static_cast(t.x) / std::numeric_limits::max(), + static_cast(t.y) / std::numeric_limits::max(), + static_cast(t.z) / std::numeric_limits::max(), + static_cast(t.w) / std::numeric_limits::max()}; + return f; + } + } +} + +/*Check if T is floatx*/ +template inline bool constexpr isFloat() { + if constexpr (std::is_scalar::value) + return std::is_floating_point::value; + else { + return std::is_floating_point::value; + } + return false; +} diff --git a/projects/hip-tests/catch/include/hip_texture_helper.hh b/projects/hip-tests/catch/include/hip_texture_helper.hh index 461966abfa..81dd07e09e 100644 --- a/projects/hip-tests/catch/include/hip_texture_helper.hh +++ b/projects/hip-tests/catch/include/hip_texture_helper.hh @@ -49,6 +49,11 @@ inline __host__ __device__ operator*=(T &a, const decltype(T::x) &b) } #endif // HT_NVIDIA +template 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(fc); @@ -216,7 +221,7 @@ T hipTextureGetValue(const T *data, const int x, const int width, default: break; } - if (sRGB && std::is_same::value) { + if constexpr (sRGB && std::is_same::value) { result = hipSRGBUnmap(result); } return result; diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index 9c7bbfa4a6..c4760d1e69 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -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) diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc new file mode 100644 index 0000000000..b27176621b --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc @@ -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 +#include +#include +#include +#include + +#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 +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(texIn, x * px); +#ifdef TEST_TEXTURE + printf("(%d/%u):(%f)\n", x, width, t.x; +#endif + surf1Dwrite(t, surfOut, x * sizeof(T)); + if (data) data[x] = t; // record it for later verification + } + if constexpr (readMode == hipReadModeNormalizedFloat) { + float4 t = tex1D(texIn, x * px); + T tc = getTypeFromNormalizedFloat(t); + surf1Dwrite(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 +__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(texMipmap, (x + offsetX) * px, lod); + if (data) data[x] = t; + } +#endif +} + +template +static void populateMipmaps(hipMipmappedArray_t mipmapArray, hipExtent size, + std::vector>& 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 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 + <<>>(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 +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<<>>(texMipmap, width, offsetX, level, hOutput); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipGetLastError()); + + for (unsigned int i = 0; i < width; i++) { + T cpuExpected = getExpectedValue(width, i + offsetX, data); + T gpuOutput = hOutput[i]; + bool match = hipTextureSamplingVerify(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 +static void testMipmapTextureObj(size_t width, float offsetX = 0.) { + std::vector> mipmapData; + size_t size = width * sizeof(T); + mipmapLevelArray 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() && 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(); + 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(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 &data = mipmapData.at(level); + + if constexpr (hipReadModeNormalizedFloat == readMode) { + typedef decltype(getNormalizedFloatType(*data.data)) TargetType; + std::vector fData; + fData.resize(data.e.width); + for (unsigned int i = 0; i < data.e.width; i++) { + fData[i] = getNormalizedFloatType(data.data[i]); + } + verifyMipmapLevel( + texMipmap, fData.data(), data.e.width, level, offsetX); + } else { // hipReadModeElementType == readMode + verifyMipmapLevel( + 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(23, 0.49); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeClamp 67") { + testMipmapTextureObj(67, -0.3); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeBorder 131") { + testMipmapTextureObj(131, 0.15); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeBorder 263") { + testMipmapTextureObj(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(23, -0.9); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeClamp 131") { + testMipmapTextureObj(131, 0.15); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeClamp 67") { + testMipmapTextureObj(67, -0.3); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeClamp 263") { + testMipmapTextureObj(263, 0.13); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeBorder 131") { + testMipmapTextureObj(131, -0.34); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeBorder 23") { + testMipmapTextureObj(23, 0.4); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeBorder 263") { + testMipmapTextureObj(263, 0.96); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeBorder 67") { + testMipmapTextureObj(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(23, 0.7); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeClamp 23") { + testMipmapTextureObj(23, -0.67); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeClamp 263") { + testMipmapTextureObj(263, 0.13); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeBorder 131") { + testMipmapTextureObj(131, 0.96); + } + SECTION( + "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeBorder 67") { + testMipmapTextureObj(67, -0.97); + } +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc index 35f4557f42..dadeeab2af 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc @@ -19,6 +19,16 @@ THE SOFTWARE. #pragma clang diagnostic ignored "-Wunused-parameter" #include +#include +#include +#include +#include + +#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 hw_vector = {2048, 1024, 512, 256, 64}; -std::vector 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 +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(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(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, surfOut, x * sizeof(T), y); + if (data) data[y * width + x] = t; // record it for later verification + } + if constexpr (readMode == hipReadModeNormalizedFloat) { + float4 t = tex2D(texIn, x * px, y * py); + T tc = getTypeFromNormalizedFloat(t); + surf2Dwrite(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 +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(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(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 +static void populateMipmaps(hipMipmappedArray_t mipmapArray, hipExtent size, + std::vector>& 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 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 + <<>>(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 +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<<>>(texMipmap, width, height, offsetX, + offsetY, level, hOutput); HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipGetLastError()); - float* hOutputData = reinterpret_cast(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( + width, height, i + offsetX, j + offsetY, data); + T gpuOutput = hOutput[index]; + bool match = hipTextureSamplingVerify(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 +static void testMipmapTextureObj(size_t width, size_t height, float offsetX = 0., + float offsetY = 0.) { + std::vector> mipmapData; + size_t size = width * height * sizeof(T); + mipmapLevelArray 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() && 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(); + 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(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 &data = mipmapData.at(level); + + if constexpr (hipReadModeNormalizedFloat == readMode) { + typedef decltype(getNormalizedFloatType(*data.data)) TargetType; + std::vector 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(data.data[index]); + } + } + verifyMipmapLevel( + texMipmap, fData.data(), data.e.width, data.e.height, level, offsetX, offsetY); + + } else { // hipReadModeElementType == readMode + verifyMipmapLevel( + 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(pow(2, (mip * 2)))) > 0) { - runMipMapTest(hw, hw, mip); - } - } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeClamp 23, 21") { + testMipmapTextureObj(23, 21, 0.4, -0.9); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeClamp 67, 131") { + testMipmapTextureObj(67, 131, -0.3, -0.67); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeBorder 131, 263") { + testMipmapTextureObj(131, 263, 0.15, -0.34); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeBorder 263, 67") { + testMipmapTextureObj(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(23, 21, 0.4, -0.9); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeClamp 131, 263") { + testMipmapTextureObj(131, 263, 0.15, -0.34); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeClamp 67, 131") { + testMipmapTextureObj(23, 17, -0.3, -0.67); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeClamp 263, 67") { + testMipmapTextureObj(263, 67, 0.13, 0.96); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeBorder 131, 263") { + testMipmapTextureObj(131, 263, 0.15, -0.34); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeBorder 23, 21") { + testMipmapTextureObj(23, 21, 0.4, -0.9); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeBorder 263, 67") { + testMipmapTextureObj(263, 67, 0.13, 0.96); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeBorder 67, 131") { + testMipmapTextureObj(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(23, 17, 0.79, 0.37); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeClamp 23, 17") { + testMipmapTextureObj(23, 17, -0.3, -0.67); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeClamp 263, 67") { + testMipmapTextureObj(263, 67, 0.13, 0.96); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeBorder 263, 67") { + testMipmapTextureObj(263, 67, 0.13, 0.96); + } + SECTION( + "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeBorder 67, 131") { + testMipmapTextureObj(67, 131, -0.3, -0.67); } -#else - SUCCEED("Mipmaps are Supported only on windows, skipping the test."); -#endif } diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc new file mode 100644 index 0000000000..fede084e0d --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc @@ -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 +#include +#include +#include +#include + +#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 +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(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, 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(texIn, x * px, y * py, z * pz); + T tc = getTypeFromNormalizedFloat(t); + surf3Dwrite(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 +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(texMipmap, (x + offsetX) * px, (y + offsetY) * py, (z + offsetZ) * pz, lod); + if (data) data[z * width * height + y * width + x] = t; + } +#endif +} + +template +static void populateMipmaps(hipMipmappedArray_t mipmapArray, hipExtent size, + std::vector>& 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 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 + <<>>(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 +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<<>>(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( + width, height, depth, i + offsetX, j + offsetY, k + offsetZ, data); + T gpuOutput = hOutput[index]; + bool match = hipTextureSamplingVerify(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 +static void testMipmapTextureObj(size_t width, size_t height, size_t depth, float offsetX = 0., + float offsetY = 0., float offsetZ = 0.) { + std::vector> mipmapData; + size_t size = width * height * depth * sizeof(T); + mipmapLevelArray 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() && 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(); + 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(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 &data = mipmapData.at(level); + + if constexpr (hipReadModeNormalizedFloat == readMode) { + typedef decltype(getNormalizedFloatType(*data.data)) TargetType; + std::vector 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(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( + texMipmap, fData.data(), data.e.width, data.e.height, data.e.depth, + level, offsetX, offsetY, offsetZ); + + } else { // hipReadModeElementType == readMode + verifyMipmapLevel( + 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(23, 21, 47, 0.4, -0.9, 0.77); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeClamp 67, 131, 99") { + testMipmapTextureObj(67, 131, 99, -0.3, -0.67, 0.49); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeBorder 131, 263, 31") { + testMipmapTextureObj(131, 263, 31, 0.15, -0.34, 0.85); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, " + "hipAddressModeBorder 263, 67, 17") { + testMipmapTextureObj(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(23, 21, 67, 0.4, -0.9, 0.37); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeClamp 131, 263, 11") { + testMipmapTextureObj(131, 263, 11, 0.15, -0.34, 0.83); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeClamp 67, 131, 53") { + testMipmapTextureObj(23, 17, 53, -0.3, -0.67, 0.78); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeClamp 263, 67, 37") { + testMipmapTextureObj(263, 67, 37, 0.13, 0.96, -0.96); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeBorder 131, 263, 11") { + testMipmapTextureObj(131, 263, 11, 0.15, -0.34, -0.11); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " + "hipAddressModeBorder 23, 21, 201") { + testMipmapTextureObj(23, 21, 201, 0.4, -0.9, 0.54); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeBorder 263, 67, 51") { + testMipmapTextureObj(263, 67, 51, 0.13, 0.96, 0.77); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModeLinear, " + "hipAddressModeBorder 67, 131, 87") { + testMipmapTextureObj(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(23, 17, 301, 0.19, 0.46, -0.9); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeClamp 23, 17, 243") { + testMipmapTextureObj(23, 17, 243, -0.3, -0.67, 0.65); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeClamp 263, 67, 39") { + testMipmapTextureObj(263, 67, 39, 0.13, 0.96, 0.66); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeBorder 263, 67, 117") { + testMipmapTextureObj(263, 67, 117, 0.13, 0.96, -0.93); + } + SECTION( + "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, " + "hipAddressModeBorder 67, 131, 67") { + testMipmapTextureObj(67, 131, 67, -0.3, -0.67, 0.88); + } +}