From 209cbd78710df02aa2a9a374a90eb4db6a9df321 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 1 Apr 2022 09:01:58 +0530 Subject: [PATCH] SWDEV-321874 - [catch2][dtest] Texture tests migrated to catch2 framework (#2577) Change-Id: I5d88438d2bf6d20d2ebde8e6ae0cbd1d27630045 [ROCm/hip-tests commit: a977b9c8ce0cef95ba15aa10a43543ca34bb5ae2] --- .../catch/unit/texture/CMakeLists.txt | 3 + .../catch/unit/texture/hipGetChanDesc.cc | 53 +++++++++ .../catch/unit/texture/hipTexObjPitch.cc | 106 ++++++++++++++++++ .../unit/texture/hipTextureObj1DFetch.cc | 84 ++++++++++++++ 4 files changed, 246 insertions(+) create mode 100644 projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index 25ea95ef89..773bd309f0 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -31,6 +31,9 @@ set(TEST_SRC hipTextureRef2D.cc hipSimpleTexture2DLayered.cc hipTextureMipmapObj2D.cc + hipGetChanDesc.cc + hipTexObjPitch.cc + hipTextureObj1DFetch.cc ) hip_add_exe_to_target(NAME TextureTest diff --git a/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc b/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc new file mode 100644 index 0000000000..e5c37412ce --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc @@ -0,0 +1,53 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +#define R 8 // rows, height +#define C 8 // columns, width + + +TEST_CASE("Unit_hipGetChannelDesc_CreateAndGet") { + hipChannelFormatDesc chan_test, chan_desc; + hipArray *hipArray; + +#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 + + chan_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned); + HIP_CHECK(hipMallocArray(&hipArray, &chan_desc, C, R, 0)); + HIP_CHECK(hipGetChannelDesc(&chan_test, hipArray)); + + if ((chan_test.x != 32) || (chan_test.y != 0) + || (chan_test.z != 0) || (chan_test.f != 0)) { + INFO("Mismatch observed : " << chan_test.x << chan_test.y + << chan_test.z << chan_test.f); + REQUIRE(false); + } + + + HIP_CHECK(hipFreeArray(hipArray)); +} diff --git a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc new file mode 100644 index 0000000000..22a58e986d --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc @@ -0,0 +1,106 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#define SIZE_H 20 +#define SIZE_W 179 + +// texture object is a kernel argument +template +static __global__ void texture2dCopyKernel(hipTextureObject_t texObj, + TYPE_t* dst) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + for (int i = 0; i < SIZE_H; i++) + for (int j = 0; j < SIZE_W; j++) + dst[SIZE_W*i+j] = tex2D(texObj, j, i); + __syncthreads(); +#endif +} + + +TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int, + unsigned char, int16_t, char, unsigned int) { + TestType* B; + TestType* A; + TestType* devPtrB; + TestType* devPtrA; + +#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 + + B = new TestType[SIZE_H*SIZE_W]; + A = new TestType[SIZE_H*SIZE_W]; + for (size_t i=1; i <= (SIZE_H*SIZE_W); i++) { + A[i-1] = i; + } + + size_t devPitchA; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtrA), &devPitchA, + SIZE_W*sizeof(TestType), SIZE_H)); + HIP_CHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TestType), + SIZE_W*sizeof(TestType), SIZE_H, hipMemcpyHostToDevice)); + + // Use the texture object + hipResourceDesc texRes; + memset(&texRes, 0, sizeof(texRes)); + texRes.resType = hipResourceTypePitch2D; + texRes.res.pitch2D.devPtr = devPtrA; + texRes.res.pitch2D.height = SIZE_H; + texRes.res.pitch2D.width = SIZE_W; + texRes.res.pitch2D.pitchInBytes = devPitchA; + texRes.res.pitch2D.desc = hipCreateChannelDesc(); + + hipTextureDesc texDescr; + memset(&texDescr, 0, sizeof(texDescr)); + texDescr.normalizedCoords = false; + texDescr.filterMode = hipFilterModePoint; + texDescr.mipmapFilterMode = hipFilterModePoint; + texDescr.addressMode[0] = hipAddressModeClamp; + texDescr.addressMode[1] = hipAddressModeClamp; + texDescr.addressMode[2] = hipAddressModeClamp; + texDescr.readMode = hipReadModeElementType; + + hipTextureObject_t texObj; + HIP_CHECK(hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL)); + + HIP_CHECK(hipMalloc(reinterpret_cast(&devPtrB), + SIZE_W*sizeof(TestType)*SIZE_H)); + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, + texObj, devPtrB); + + HIP_CHECK(hipMemcpy2D(B, SIZE_W*sizeof(TestType), devPtrB, + SIZE_W*sizeof(TestType), SIZE_W*sizeof(TestType), + SIZE_H, hipMemcpyDeviceToHost)); + + HipTest::checkArray(A, B, SIZE_H, SIZE_W); + delete []A; + delete []B; + hipFree(devPtrA); + hipFree(devPtrB); +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc new file mode 100644 index 0000000000..a52b8baaf4 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc @@ -0,0 +1,84 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +#define N 512 + +static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { + int k = blockIdx.x * blockDim.x + threadIdx.x; + if (k < N) { + val[k] = tex1Dfetch(obj, k); + } +} + + +TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") { + // Allocating the required buffer on gpu device + float *texBuf, *texBufOut; + float val[N], output[N]; + + for (int i = 0; i < N; i++) { + val[i] = (i + 1) * (i + 1); + output[i] = 0.0; + } + + HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float))); + HIP_CHECK(hipMalloc(&texBufOut, N * sizeof(float))); + HIP_CHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemset(texBufOut, 0, N * sizeof(float))); + hipResourceDesc resDescLinear; + + memset(&resDescLinear, 0, sizeof(resDescLinear)); + resDescLinear.resType = hipResourceTypeLinear; + resDescLinear.res.linear.devPtr = texBuf; + resDescLinear.res.linear.desc = + hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + resDescLinear.res.linear.sizeInBytes = N * sizeof(float); + + hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + + // Creating texture object + hipTextureObject_t texObj = 0; + HIP_CHECK(hipCreateTextureObject(&texObj, &resDescLinear, &texDesc, NULL)); + + dim3 dimBlock(64, 1, 1); + dim3 dimGrid(N / dimBlock.x, 1, 1); + + hipLaunchKernelGGL(tex1dKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, + texBufOut, texObj); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(output, texBufOut, N * sizeof(float), + hipMemcpyDeviceToHost)); + + for (int i = 0; i < N; i++) { + if (output[i] != val[i]) { + INFO("Mismatch at index : " << i << ", output[i] " << output[i] + << ", val[i] " << val[i]); + REQUIRE(false); + } + } + + HIP_CHECK(hipDestroyTextureObject(texObj)); + HIP_CHECK(hipFree(texBuf)); + HIP_CHECK(hipFree(texBufOut)); +}