diff --git a/projects/hip/tests/catch/unit/texture/CMakeLists.txt b/projects/hip/tests/catch/unit/texture/CMakeLists.txt index 2472abf163..30be9023aa 100644 --- a/projects/hip/tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/texture/CMakeLists.txt @@ -24,6 +24,7 @@ set(TEST_SRC hipCreateTextureObject_Linear.cc hipCreateTextureObject_Pitch2D.cc hipCreateTextureObject_Array.cc + hipTextureObjFetchVector.cc ) hip_add_exe_to_target(NAME TextureTest diff --git a/projects/hip/tests/catch/unit/texture/hipTextureObjFetchVector.cc b/projects/hip/tests/catch/unit/texture/hipTextureObjFetchVector.cc new file mode 100644 index 0000000000..8fc6c6730a --- /dev/null +++ b/projects/hip/tests/catch/unit/texture/hipTextureObjFetchVector.cc @@ -0,0 +1,240 @@ +/* +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 +#include + +template +__global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) { + int k = blockIdx.x * blockDim.x + threadIdx.x; + if (k < N) { + val[k] = tex1Dfetch(obj, k); + } +} + +template +static inline __host__ __device__ constexpr int rank() { + return sizeof(T) / sizeof(decltype(T::x)); +} + +template +static inline T getRandom() { + double r = 0; + if (std::is_signed < T > ::value) { + r = (std::rand() - RAND_MAX / 2.0) / (RAND_MAX / 2.0 + 1.); + } else { + r = std::rand() / (RAND_MAX + 1.); + } + return static_cast(std::numeric_limits < T > ::max() * r); +} + +template< + typename T, + typename std::enable_if() == 1>::type* = nullptr> +static inline void initVal(T &val) { + val.x = getRandom(); +} + +template< + typename T, + typename std::enable_if() == 2>::type* = nullptr> +static inline void initVal(T &val) { + val.x = getRandom(); + val.y = getRandom(); +} + +template< + typename T, + typename std::enable_if() == 4>::type* = nullptr> +static inline void initVal(T &val) { + val.x = getRandom(); + val.y = getRandom(); + val.z = getRandom(); + val.w = getRandom(); +} + +template< + typename T, + typename std::enable_if() == 1>::type* = nullptr> +static inline void printVector(T &val) { + using B = decltype(T::x); + constexpr bool isChar = std::is_same::value + || std::is_same::value; + std::cout << "("; + std::cout << (isChar ? static_cast(val.x) : val.x); + std::cout << ")"; +} + +template< + typename T, + typename std::enable_if() == 2>::type* = nullptr> +static inline void printVector(T &val) { + using B = decltype(T::x); + constexpr bool isChar = std::is_same::value + || std::is_same::value; + std::cout << "("; + std::cout << (isChar ? static_cast(val.x) : val.x); + std::cout << ", " << (isChar ? static_cast(val.y) : val.y); + std::cout << ")"; +} + +template< + typename T, + typename std::enable_if() == 4>::type* = nullptr> +static inline void printVector(T &val) { + using B = decltype(T::x); + constexpr bool isChar = std::is_same::value + || std::is_same::value; + std::cout << "("; + std::cout << (isChar ? static_cast(val.x) : val.x); + std::cout << ", " << (isChar ? static_cast(val.y) : val.y); + std::cout << ", " << (isChar ? static_cast(val.z) : val.z); + std::cout << ", " << (isChar ? static_cast(val.w) : val.w); + std::cout << ")"; +} + +template< + typename T, + typename std::enable_if() == 1>::type* = nullptr> +static inline bool isEqual(const T &val0, const T &val1) { + return val0.x == val1.x; +} + +template< + typename T, + typename std::enable_if() == 2>::type* = nullptr> +static inline bool isEqual(const T &val0, const T &val1) { + return val0.x == val1.x && + val0.y == val1.y; +} + +template< + typename T, + typename std::enable_if() == 4>::type* = nullptr> +static inline bool isEqual(const T &val0, const T &val1) { + return val0.x == val1.x && + val0.y == val1.y && + val0.z == val1.z && + val0.w == val1.w; +} + +template +bool runTest(const char *description) { + const int N = 1024; + bool testResult = true; + // Allocating the required buffer on gpu device + T *texBuf, *texBufOut; + T val[N], output[N]; + printf("%s<%s>(): size: %zu, %zu\n", __FUNCTION__, description, + sizeof(T), sizeof(decltype(T::x))); + + memset(output, 0, sizeof(output)); + std::srand(std::time(nullptr)); // use current time as seed for random generator + + for (int i = 0; i < N; i++) { + initVal(val[i]); + } + + HIP_CHECK(hipMalloc(&texBuf, N * sizeof(T))); + HIP_CHECK(hipMalloc(&texBufOut, N * sizeof(T))); + HIP_CHECK(hipMemcpy(texBuf, val, N * sizeof(T), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemset(texBufOut, 0, N * sizeof(T))); + hipResourceDesc resDescLinear; + + memset(&resDescLinear, 0, sizeof(resDescLinear)); + resDescLinear.resType = hipResourceTypeLinear; + resDescLinear.res.linear.devPtr = texBuf; + resDescLinear.res.linear.desc = hipCreateChannelDesc(); + resDescLinear.res.linear.sizeInBytes = N * sizeof(T); + + hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + texDesc.addressMode[0] = hipAddressModeClamp; + + // 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) / dimBlock.x, 1, 1); + + hipLaunchKernelGGL(tex1dKernelFetch, dimGrid, dimBlock, 0, 0, texBufOut, + texObj, N); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(output, texBufOut, N * sizeof(T), hipMemcpyDeviceToHost)); + + for (int i = 0; i < N; i++) { + if (!isEqual(output[i], val[i])) { + std::cout << "output[" << i << "]= "; + printVector(output[i]); + std::cout << ", expected val[" << i << "]= "; + printVector(val[i]); + std::cout << "\n"; + testResult = false; + break; + } + } + HIP_CHECK(hipDestroyTextureObject(texObj)); + HIP_CHECK(hipFree(texBuf)); + HIP_CHECK(hipFree(texBufOut)); + + printf(": %s\n", testResult ? "succeeded" : "failed"); + REQUIRE(testResult == true); + return testResult; +} + +TEST_CASE("Unit_hipTextureFetch_vector") { + // test for char + runTest("char1"); + runTest("char2"); + runTest("char4"); + + // test for uchar + runTest("uchar1"); + runTest("uchar2"); + runTest("uchar4"); + + // test for short + runTest("short1"); + runTest("short2"); + runTest("short4"); + + // test for ushort + runTest("ushort1"); + runTest("ushort2"); + runTest("ushort4"); + + // test for int + runTest("int1"); + runTest("int2"); + runTest("int4"); + + // test for unsigned int + runTest("uint1"); + runTest("uint2"); + runTest("uint4"); + + // test for float + runTest("float1"); + runTest("float2"); + runTest("float4"); +}