From f13ebc4e4e8aa980a1258d9c165f779d67160068 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 11 Mar 2022 16:53:48 +0530 Subject: [PATCH] SWDEV-321872 - [catch2][dtest] Texture tests migrated to catch2 (#2536) Change-Id: I4e29001812e2310680c6c24028667eb560194e87 [ROCm/hip-tests commit: 3f4fb28faf2d9be805f8806fda916d3acd6949e1] --- projects/hip-tests/catch/CMakeLists.txt | 2 +- .../catch/unit/texture/CMakeLists.txt | 6 + .../texture/hipNormalizedFloatValueTex.cc | 166 ++++++++++++++++++ .../unit/texture/hipSimpleTexture2DLayered.cc | 103 +++++++++++ .../catch/unit/texture/hipSimpleTexture3D.cc | 127 ++++++++++++++ .../unit/texture/hipTextureMipmapObj2D.cc | 153 ++++++++++++++++ .../catch/unit/texture/hipTextureObj2D.cc | 113 ++++++++++++ .../catch/unit/texture/hipTextureRef2D.cc | 87 +++++++++ 8 files changed, 756 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc diff --git a/projects/hip-tests/catch/CMakeLists.txt b/projects/hip-tests/catch/CMakeLists.txt index 0bc5231260..ed978a1b6a 100644 --- a/projects/hip-tests/catch/CMakeLists.txt +++ b/projects/hip-tests/catch/CMakeLists.txt @@ -85,7 +85,7 @@ set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake) if(HIP_PLATFORM MATCHES "amd" AND HIP_COMPILER MATCHES "clang") - add_compile_options(-Wall -Wextra -pedantic -Werror) + add_compile_options(-Wall -Wextra -pedantic -Werror -Wno-deprecated) endif() cmake_policy(PUSH) diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index 30be9023aa..25ea95ef89 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -25,6 +25,12 @@ set(TEST_SRC hipCreateTextureObject_Pitch2D.cc hipCreateTextureObject_Array.cc hipTextureObjFetchVector.cc + hipNormalizedFloatValueTex.cc + hipTextureObj2D.cc + hipSimpleTexture3D.cc + hipTextureRef2D.cc + hipSimpleTexture2DLayered.cc + hipTextureMipmapObj2D.cc ) hip_add_exe_to_target(NAME TextureTest diff --git a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc new file mode 100644 index 0000000000..e9cf002eae --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -0,0 +1,166 @@ +/* +Copyright (c) 2022 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 + +#define SIZE 10 +#define EPSILON 0.00001 +#define THRESH_HOLD 0.01 // For filter mode + +static float getNormalizedValue(const float value, + const hipChannelFormatDesc& desc) { + if ((desc.x == 8) && (desc.f == hipChannelFormatKindSigned)) + return (value / SCHAR_MAX); + if ((desc.x == 8) && (desc.f == hipChannelFormatKindUnsigned)) + return (value / UCHAR_MAX); + if ((desc.x == 16) && (desc.f == hipChannelFormatKindSigned)) + return (value / SHRT_MAX); + if ((desc.x == 16) && (desc.f == hipChannelFormatKindUnsigned)) + return (value / USHRT_MAX); + return value; +} + +texture texc; +texture texuc; + +template +__global__ void normalizedValTextureTest(unsigned int numElements, + float* pDst) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + unsigned int elementID = hipThreadIdx_x; + if (elementID >= numElements) + return; + float coord = elementID/static_cast(numElements); + if (std::is_same::value) + pDst[elementID] = tex1D(texc, coord); + else if (std::is_same::value) + pDst[elementID] = tex1D(texuc, coord); +#endif +} + +static void textureVerifyFilterModePoint(float *hOutputData, + float *expected, int size) { + for (int i = 0; i < size; i++) { + if ((hOutputData[i] == expected[i]) + || (i >= 1 && hOutputData[i] == expected[i - 1]) || // round down + (i < (size - 1) && hOutputData[i] == expected[i + 1])) { // round up + continue; + } + INFO("Mismatch at output[" << i << "]:" << hOutputData[i] << + " expected[" << i << "]:" << expected[i]); + if (i >= 1) { + INFO(", expected[" << i - 1 << "]:" << expected[i - 1]); + } + if (i < (size - 1)) { + INFO(", expected[" << i + 1 << "]:" << expected[i + 1]); + } + REQUIRE(false); + } +} + +static void textureVerifyFilterModeLinear(float *hOutputData, + float *expected, int size) { + for (int i = 0; i < size; i++) { + float mean = (fabs(expected[i]) + fabs(hOutputData[i])) / 2; + float ratio = fabs(expected[i] - hOutputData[i]) / (mean + EPSILON); + if (ratio > THRESH_HOLD) { + INFO("Mismatch found at output[" << i << "]:" << hOutputData[i] << + " expected[" << i << "]:" << expected[i] << ", ratio:" << ratio); + REQUIRE(false); + } + } +} + +template +static void textureVerify(float *hOutputData, float *expected, size_t size) { + if (fMode == hipFilterModePoint) { + textureVerifyFilterModePoint(hOutputData, expected, size); + } else if (fMode == hipFilterModeLinear) { + textureVerifyFilterModeLinear(hOutputData, expected, size); + } +} + +template +static void textureTest(texture *tex) { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + hipArray_t dData; + HIP_CHECK(hipMallocArray(&dData, &desc, SIZE, 1, hipArrayDefault)); + + T hData[] = {65, 66, 67, 68, 69, 70, 71, 72, 73, 74}; + HIP_CHECK(hipMemcpy2DToArray(dData, 0, 0, hData, sizeof(T) * SIZE, + sizeof(T) * SIZE, 1, hipMemcpyHostToDevice)); + + tex->normalized = true; + tex->channelDesc = desc; + tex->filterMode = fMode; + HIP_CHECK(hipBindTextureToArray(tex, dData, &desc)); + + float *dOutputData = NULL; + HIP_CHECK(hipMalloc(&dOutputData, sizeof(float) * SIZE)); + REQUIRE(dOutputData != nullptr); + + hipLaunchKernelGGL(normalizedValTextureTest, dim3(1, 1, 1), + dim3(SIZE, 1, 1), 0, 0, SIZE, dOutputData); + + float *hOutputData = new float[SIZE]; + REQUIRE(hOutputData != nullptr); + HIP_CHECK(hipMemcpy(hOutputData, dOutputData, (sizeof(float) * SIZE), + hipMemcpyDeviceToHost)); + + float expected[SIZE]; + for (int i = 0; i < SIZE; i++) { + expected[i] = getNormalizedValue(static_cast(hData[i]), desc); + } + textureVerify(hOutputData, expected, SIZE); + + HIP_CHECK(hipFreeArray(dData)); + HIP_CHECK(hipFree(dOutputData)); + delete [] hOutputData; +} + +template +static void runTest_hipTextureFilterMode() { + textureTest(&texc); + textureTest(&texuc); +} + +TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") { +#if HT_AMD + int imageSupport{}; + HIP_CHECK(hipDeviceGetAttribute(&imageSupport, + hipDeviceAttributeImageSupport, 0)); + if (!imageSupport) { + INFO("Texture is not supported on the device. Test is skipped"); + return; + } + hipDeviceProp_t props; + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipGetDeviceProperties(&props, 0)); + INFO("Device :: " << props.name); + INFO("Arch - AMD GPU :: " << props.gcnArch); +#endif + + SECTION("hipNormalizedFloatValueTexture for hipFilterModePoint") { + runTest_hipTextureFilterMode(); + } + SECTION("hipNormalizedFloatValueTexture for hipFilterModeLinear") { + runTest_hipTextureFilterMode(); + } +} diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc new file mode 100644 index 0000000000..0c08c0f6d2 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc @@ -0,0 +1,103 @@ +/* +Copyright (c) 2022 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 + +typedef float T; + +// Texture reference for 2D Layered texture +texture tex2DL; + +__global__ void simpleKernelLayeredArray(T* outputData, + int width, int height, int layer) { + unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; + unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[layer * width * height + y * width + x] = tex2DLayered(tex2DL, + x, y, layer); +} + +TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") { + constexpr int SIZE = 512; + constexpr int num_layers = 5; + constexpr unsigned int width = SIZE; + constexpr unsigned int height = SIZE; + constexpr unsigned int size = width * height * num_layers * sizeof(T); + + T* hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + + for (unsigned int layer = 0; layer < num_layers; layer++) { + for (int i = 0; i < static_cast(width * height); i++) { + hData[layer * width * height + i] = i; + } + } + hipChannelFormatDesc channelDesc; + // Allocate array and copy image data + channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, + hipChannelFormatKindFloat); + hipArray *arr; + + HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, + make_hipExtent(width, height, num_layers), hipArrayLayered)); + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width , height, num_layers); + // myparms.kind = hipMemcpyHostToDevice; + HIP_CHECK(hipMemcpy3D(&myparms)); + + // set texture parameters + tex2DL.addressMode[0] = hipAddressModeWrap; + tex2DL.addressMode[1] = hipAddressModeWrap; + tex2DL.filterMode = hipFilterModePoint; + tex2DL.normalized = false; + + // Bind the array to the texture + HIP_CHECK(hipBindTextureToArray(tex2DL, arr, channelDesc)); + + // Allocate device memory for result + T* dData = nullptr; + HIP_CHECK(hipMalloc(&dData, size)); + REQUIRE(dData != nullptr); + + dim3 dimBlock(8, 8, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + for (unsigned int layer = 0; layer < num_layers; layer++) + hipLaunchKernelGGL(simpleKernelLayeredArray, dimGrid, dimBlock, 0, 0, + dData, width, height, layer); + HIP_CHECK(hipDeviceSynchronize()); + + // Allocate mem for the result on host side + T *hOutputData = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, size); + + // copy result from device to host + HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + HipTest::checkArray(hData, hOutputData, width, height, num_layers); + + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(arr)); + free(hData); + free(hOutputData); +} diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc new file mode 100644 index 0000000000..1ff46a1f2e --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc @@ -0,0 +1,127 @@ +/* +Copyright (c) 2022 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 + +// Texture reference for 3D texture +texture texf; +texture texi; +texture texc; + +template +__global__ void simpleKernel3DArray(T* outputData, int width, + int height, int depth) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + if (std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texf, k, j, i); + else if (std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texi, k, j, i); + else if (std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texc, k, j, i); + } + } + } +#endif +} + +template +static void runSimpleTexture3D_Check(int width, int height, int depth, + texture *tex) { + unsigned int size = width * height * depth * sizeof(T); + T* hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *arr; + + HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, + make_hipExtent(width, height, depth), hipArrayDefault)); + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + // set texture parameters + tex->addressMode[0] = hipAddressModeWrap; + tex->addressMode[1] = hipAddressModeWrap; + tex->filterMode = hipFilterModePoint; + tex->normalized = false; + + // Bind the array to the texture + HIP_CHECK(hipBindTextureToArray(*tex, arr, channelDesc)); + + // Allocate device memory for result + T* dData = nullptr; + HIP_CHECK(hipMalloc(&dData, size)); + REQUIRE(dData != nullptr); + + hipLaunchKernelGGL(simpleKernel3DArray, dim3(1, 1, 1), dim3(1, 1, 1), + 0, 0, dData, width, height, depth); + HIP_CHECK(hipDeviceSynchronize()); + + // Allocate mem for the result on host side + T *hOutputData = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, size); + + // copy result from device to host + HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + HipTest::checkArray(hData, hOutputData, width, height, depth); + + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(arr)); + free(hData); + free(hOutputData); +} + +TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") { +#if HT_AMD + int imageSupport{}; + HIP_CHECK(hipDeviceGetAttribute(&imageSupport, + hipDeviceAttributeImageSupport, 0)); + if (!imageSupport) { + INFO("Texture is not supported on the device. Test is skipped"); + return; + } +#endif + for ( int i = 1; i < 25; i++ ) { + runSimpleTexture3D_Check(i, i, i, &texf); + runSimpleTexture3D_Check(i+1, i, i, &texi); + runSimpleTexture3D_Check(i, i+1, i, &texc); + } +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc new file mode 100644 index 0000000000..2bbefe3179 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc @@ -0,0 +1,153 @@ +/* +Copyright (c) 2022 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 + +// Height Width Vector +std::vector hw_vector = {2048, 1024, 512, 256, 64}; +std::vector mip_vector = {8, 4, 2, 1}; + +__global__ void tex2DKernel(float* outputData, + hipTextureObject_t textureObject, + int width, float level) { +#ifndef __gfx90a__ +#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); +#endif +#endif +} + +#ifdef _WIN32 // MipMap is currently supported only on windows +static void runMipMapTest(unsigned int width, unsigned int height, + unsigned int mipmap_level) { + INFO("Width: " << width << "Height: " << height << "mip: " << mipmap_level); + + // 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* 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; + } + } + + 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; + + hipMipmappedArray* mip_array_ptr; + HIP_CHECK(hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc, + 2 * mipmap_level)); + + hipArray *hipArray = nullptr; + HIP_CHECK(hipMipmappedArrayGetLevel(&hipArray, mip_array_ptr, mipmap_level)); + HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, + hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // 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; + + // Create texture object + hipTextureObject_t textureObject = 0; + HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, + &texDesc, nullptr)); + + float* dData = nullptr; + HIP_CHECK(hipMalloc(&dData, size)); + REQUIRE(dData != nullptr); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, + textureObject, width, (2 * mipmap_level)); + hipDeviceSynchronize(); + + 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]); + REQUIRE(false); + } + } + } + HIP_CHECK(hipDestroyTextureObject(textureObject)); + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(hipArray)); + free(hData); +} +#endif + +TEST_CASE("Unit_hipTextureMipmapObj2D_Check") { +#if HT_AMD + int imageSupport{}; + HIP_CHECK(hipDeviceGetAttribute(&imageSupport, + hipDeviceAttributeImageSupport, 0)); + if (!imageSupport) { + INFO("Texture is not supported on the device. Test is skipped"); + return; + } +#endif +#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); + } + } + } +#else + SUCCEED("Mipmaps are Supported only on windows, skipping the test."); +#endif +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc new file mode 100644 index 0000000000..1836981e41 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc @@ -0,0 +1,113 @@ +/* +Copyright (c) 2022 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 + +__global__ void tex2DKernel(float* outputData, + hipTextureObject_t textureObject, int width) { +#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] = tex2D(textureObject, x, y); +#endif +} + +TEST_CASE("Unit_hipTextureObj2D_Check") { +#if HT_AMD + int imageSupport{}; + HIP_CHECK(hipDeviceGetAttribute(&imageSupport, + hipDeviceAttributeImageSupport, 0)); + if (!imageSupport) { + INFO("Texture is not supported on the device. Test is skipped"); + return; + } +#endif + constexpr int SIZE = 256; + constexpr unsigned int width = SIZE; + constexpr unsigned int height = SIZE; + constexpr unsigned int size = width * height * sizeof(float); + unsigned int i, j; + + float* dData = nullptr; + HIP_CHECK(hipMalloc(&dData, size)); + REQUIRE(dData != nullptr); + + float* hOutputData = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, size); + + 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; + } + } + + hipChannelFormatDesc channelDesc = + hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray* hipArray; + HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height)); + HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, + hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // 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; + + // Create texture object + hipTextureObject_t textureObject = 0; + HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, + &texDesc, nullptr)); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), + 0, 0, dData, textureObject, width); + + HIP_CHECK(hipDeviceSynchronize()); + 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]); + REQUIRE(false); + } + } + } + + HIP_CHECK(hipDestroyTextureObject(textureObject)); + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(hipArray)); + free(hData); +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc new file mode 100644 index 0000000000..1a6600d099 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc @@ -0,0 +1,87 @@ +/* +Copyright (c) 2022 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 + +texture tex; + +__global__ void tex2DKernel(float* outputData, int width) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[y * width + x] = tex2D(tex, x, y); +} + +TEST_CASE("Unit_hipTextureRef2D_Check") { + constexpr int SIZE = 256; + constexpr unsigned int width = SIZE; + constexpr unsigned int height = SIZE; + constexpr unsigned int size = width * height * sizeof(float); + unsigned int i, j; + + 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; + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, + hipChannelFormatKindFloat); + hipArray* hipArray; + HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height)); + HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, + hipMemcpyHostToDevice)); + + tex.addressMode[0] = hipAddressModeWrap; + tex.addressMode[1] = hipAddressModeWrap; + tex.filterMode = hipFilterModePoint; + tex.normalized = 0; + + HIP_CHECK(hipBindTextureToArray(tex, hipArray, channelDesc)); + + float* dData = nullptr; + HIP_CHECK(hipMalloc(&dData, size)); + REQUIRE(dData != nullptr); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, + dData, width); + hipDeviceSynchronize(); + + 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]); + REQUIRE(false); + } + } + } + HIP_CHECK(hipUnbindTexture(tex)); + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(hipArray)); +}