From d7e34e85346b63fc38e8d47e967e2d4bdacc4dbd Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 10 May 2022 05:11:37 +0530 Subject: [PATCH] SWDEV-321874 - [catch2][dtest] Texture dtests migrated from HIT to catch2 framework (#2576) Change-Id: Ice382bf552fd0a9b41133f2a6ab333e3672b7c28 --- tests/catch/unit/texture/CMakeLists.txt | 3 + tests/catch/unit/texture/hipBindTex2DPitch.cc | 74 +++++++++++ .../unit/texture/hipBindTexRef1DFetch.cc | 77 ++++++++++++ .../unit/texture/hipTex1DFetchCheckModes.cc | 118 ++++++++++++++++++ 4 files changed, 272 insertions(+) create mode 100644 tests/catch/unit/texture/hipBindTex2DPitch.cc create mode 100644 tests/catch/unit/texture/hipBindTexRef1DFetch.cc create mode 100644 tests/catch/unit/texture/hipTex1DFetchCheckModes.cc diff --git a/tests/catch/unit/texture/CMakeLists.txt b/tests/catch/unit/texture/CMakeLists.txt index 773bd309f0..e9c24a8226 100644 --- a/tests/catch/unit/texture/CMakeLists.txt +++ b/tests/catch/unit/texture/CMakeLists.txt @@ -31,6 +31,9 @@ set(TEST_SRC hipTextureRef2D.cc hipSimpleTexture2DLayered.cc hipTextureMipmapObj2D.cc + hipBindTex2DPitch.cc + hipBindTexRef1DFetch.cc + hipTex1DFetchCheckModes.cc hipGetChanDesc.cc hipTexObjPitch.cc hipTextureObj1DFetch.cc diff --git a/tests/catch/unit/texture/hipBindTex2DPitch.cc b/tests/catch/unit/texture/hipBindTex2DPitch.cc new file mode 100644 index 0000000000..eff6a9f939 --- /dev/null +++ b/tests/catch/unit/texture/hipBindTex2DPitch.cc @@ -0,0 +1,74 @@ +/* +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 8 +#define SIZE_W 12 +#define TYPE_t float + +texture tex; + +// texture object is a kernel argument +static __global__ void texture2dCopyKernel(TYPE_t* dst) { + int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; + if ( (x < SIZE_W) && (y < SIZE_H) ) { + dst[SIZE_W*y+x] = tex2D(tex, x, y); + } +} + + +TEST_CASE("Unit_hipBindTexture2D_Pitch") { + TYPE_t* B; + TYPE_t* A; + TYPE_t* devPtrB; + TYPE_t* devPtrA; + + B = new TYPE_t[SIZE_H*SIZE_W]; + A = new TYPE_t[SIZE_H*SIZE_W]; + for (size_t i = 1; i <= (SIZE_H * SIZE_W); i++) { + A[i-1] = i; + } + + size_t devPitchA, tex_ofs; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtrA), &devPitchA, + SIZE_W*sizeof(TYPE_t), SIZE_H)); + HIP_CHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); + + tex.normalized = false; + HIP_CHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, + SIZE_W, SIZE_H, devPitchA)); + HIP_CHECK(hipMalloc(reinterpret_cast(&devPtrB), + SIZE_W*sizeof(TYPE_t)*SIZE_H)); + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(4, 4, 1), dim3(32, 32, 1), + 0, 0, devPtrB); + hipDeviceSynchronize(); + HIP_CHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, + SIZE_W*sizeof(TYPE_t), SIZE_W*sizeof(TYPE_t), + SIZE_H, hipMemcpyDeviceToHost)); + HipTest::checkArray(A, B, SIZE_H, SIZE_W); + delete []A; + delete []B; + hipFree(devPtrA); + hipFree(devPtrB); +} diff --git a/tests/catch/unit/texture/hipBindTexRef1DFetch.cc b/tests/catch/unit/texture/hipBindTexRef1DFetch.cc new file mode 100644 index 0000000000..e714958163 --- /dev/null +++ b/tests/catch/unit/texture/hipBindTexRef1DFetch.cc @@ -0,0 +1,77 @@ +/* +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 + +texture tex; + +static __global__ void kernel(float *out) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + if (x < N) { + out[x] = tex1Dfetch(tex, x); + } +} + + +TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") { + float *texBuf; + float val[N], output[N]; + size_t offset = 0; + float *devBuf; + for (int i = 0; i < N; i++) { + val[i] = i; + output[i] = 0.0; + } + hipChannelFormatDesc chanDesc = + hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + + HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float))); + HIP_CHECK(hipMalloc(&devBuf, N * sizeof(float))); + HIP_CHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice)); + + tex.addressMode[0] = hipAddressModeClamp; + tex.addressMode[1] = hipAddressModeClamp; + tex.filterMode = hipFilterModePoint; + tex.normalized = 0; + + HIP_CHECK(hipBindTexture(&offset, tex, reinterpret_cast(texBuf), + chanDesc, N * sizeof(float))); + HIP_CHECK(hipGetTextureAlignmentOffset(&offset, &tex)); + + dim3 dimBlock(64, 1, 1); + dim3 dimGrid(N / dimBlock.x, 1, 1); + + hipLaunchKernelGGL(kernel, dim3(dimGrid), dim3(dimBlock), 0, 0, devBuf); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(output, devBuf, 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(hipUnbindTexture(&tex)); + HIP_CHECK(hipFree(texBuf)); + HIP_CHECK(hipFree(devBuf)); +} diff --git a/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc b/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc new file mode 100644 index 0000000000..eb3501191c --- /dev/null +++ b/tests/catch/unit/texture/hipTex1DFetchCheckModes.cc @@ -0,0 +1,118 @@ +/* +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 16 +#define offset 3 + +static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { + int k = blockIdx.x * blockDim.x + threadIdx.x; + if (k < (N - offset)) + val[k] = tex1Dfetch(obj, k+offset); +} + + +static void runTest(hipTextureAddressMode addressMode, + hipTextureFilterMode filterMode) { + hipCtx_t HipContext; + hipDevice_t HipDevice; + int deviceID = 0; + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipDeviceGet(&HipDevice, deviceID)); + HIP_CHECK(hipCtxCreate(&HipContext, 0, HipDevice)); + + // 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; + 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; + texDesc.addressMode[0] = addressMode; + texDesc.addressMode[1] = addressMode; + texDesc.filterMode = filterMode; + texDesc.normalizedCoords = false; + + // Creating texture object + hipTextureObject_t texObj = 0; + HIP_CHECK(hipCreateTextureObject(&texObj, &resDescLinear, &texDesc, NULL)); + + dim3 dimBlock(1, 1, 1); + dim3 dimGrid(N, 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 - offset); i++) { + if (output[i] != val[i + offset]) { + INFO("Output not matching at index " << i); + REQUIRE(false); + } + } + + for (int i = (N - offset); i < N; i++) { + if (output[i] != 0) { + INFO("Output found to be updated at index " << i); + REQUIRE(false); + } + } + + HIP_CHECK(hipDestroyTextureObject(texObj)); + HIP_CHECK(hipFree(texBuf)); + HIP_CHECK(hipFree(texBufOut)); +} + + +TEST_CASE("Unit_tex1Dfetch_CheckModes") { + SECTION("hipAddressModeClamp AND hipFilterModePoint") { + runTest(hipAddressModeClamp, hipFilterModePoint); + } + SECTION("hipAddressModeClamp AND hipFilterModeLinear") { + runTest(hipAddressModeClamp, hipFilterModeLinear); + } + SECTION("hipAddressModeWrap AND hipFilterModePoint") { + runTest(hipAddressModeWrap, hipFilterModePoint); + } + SECTION("hipAddressModeWrap AND hipFilterModeLinear") { + runTest(hipAddressModeWrap, hipFilterModeLinear); + } +}