diff --git a/catch/include/hip_array_common.hh b/catch/include/hip_array_common.hh index fd6f094f8d..a0b5d96622 100644 --- a/catch/include/hip_array_common.hh +++ b/catch/include/hip_array_common.hh @@ -81,4 +81,136 @@ struct vector_info : type_and_size_and_format {}; template <> struct vector_info - : type_and_size_and_format {}; \ No newline at end of file + : type_and_size_and_format {}; + +template < + typename T, + typename std::enable_if::value == false>::type* = nullptr> +static inline __host__ __device__ constexpr int rank() { + return sizeof(T) / sizeof(decltype(T::x)); +} + +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< + typename T, + typename std::enable_if::value>::type* = nullptr> +static inline bool isEqual(const T &val0, const T &val1) { + return val0 == val1; +} + +template< + typename T, + typename std::enable_if() == 1>::type* = nullptr> +const std::string getString(const T& t) +{ + std::ostringstream os; + os<< "(" << t.x << ")"; + return os.str(); +} + +template< + typename T, + typename std::enable_if() == 2>::type* = nullptr> +const std::string getString(const T& t) +{ + std::ostringstream os; + os<< "(" << t.x << ", " << t.y << ")"; + return os.str(); +} + +template< + typename T, + typename std::enable_if() == 3>::type* = nullptr> +const std::string getString(const T& t) +{ + std::ostringstream os; + os<< "(" << t.x << ", " << t.y << ", " << t.z << ")"; + return os.str(); +} + +template< + typename T, + typename std::enable_if() == 4>::type* = nullptr> +const std::string getString(const T& t) +{ + std::ostringstream os; + os<< "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")"; + return os.str(); +} + +template< + typename T, + typename std::enable_if::value>::type* = nullptr> +std::string getString(const T& t) +{ + std::ostringstream os; + os << t; + return os.str(); +} + +template +static inline T getRandom() { + double r = 0; + if (std::is_signed::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::value>::type* = nullptr> +static inline void initVal(T &val) { + val = getRandom(); +} diff --git a/catch/include/hip_texture_helper.hh b/catch/include/hip_texture_helper.hh index 82404a5471..461966abfa 100644 --- a/catch/include/hip_texture_helper.hh +++ b/catch/include/hip_texture_helper.hh @@ -8,6 +8,8 @@ #define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD 0.1 #if HT_NVIDIA +typedef unsigned char uchar; + template typename std::enable_if::type inline __host__ __device__ operator+(const T &a, const T &b) diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 31bd37da9b..c26af21b02 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -28,6 +28,7 @@ add_subdirectory(occupancy) add_subdirectory(device) add_subdirectory(printf) add_subdirectory(texture) +add_subdirectory(surface) add_subdirectory(streamperthread) add_subdirectory(kernel) add_subdirectory(multiThread) diff --git a/catch/unit/surface/CMakeLists.txt b/catch/unit/surface/CMakeLists.txt new file mode 100644 index 0000000000..4bc01e0794 --- /dev/null +++ b/catch/unit/surface/CMakeLists.txt @@ -0,0 +1,30 @@ +# 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. + +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipSurfaceObj1D.cc + hipSurfaceObj2D.cc + hipSurfaceObj3D.cc +) + +hip_add_exe_to_target(NAME SurfaceTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) \ No newline at end of file diff --git a/catch/unit/surface/hipSurfaceObj1D.cc b/catch/unit/surface/hipSurfaceObj1D.cc new file mode 100644 index 0000000000..5068e770b6 --- /dev/null +++ b/catch/unit/surface/hipSurfaceObj1D.cc @@ -0,0 +1,317 @@ +/* +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 + +template +__global__ void +surf1DKernelR(hipSurfaceObject_t surfaceObject, + T* outputData, int width) +{ +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + if (x < width) { + surf1Dread(outputData + x, surfaceObject, x * sizeof(T)); + } +#endif +} + +template +__global__ void +surf1DKernelW(hipSurfaceObject_t surfaceObject, + T* inputData, int width) +{ +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + if (x < width) { + surf1Dwrite(inputData[x], surfaceObject, x * sizeof(T)); + } +#endif +} + +template +__global__ void +surf1DKernelRW(hipSurfaceObject_t surfaceObject, + hipSurfaceObject_t outputSurfObj, int width) +{ +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + if (x < width) { + T data; + surf1Dread(&data, surfaceObject, x * sizeof(T)); + surf1Dwrite(data, outputSurfObj, x * sizeof(T)); + } +#endif +} + +template +static void runTestR(const int width) +{ + unsigned int size = width * sizeof(T); + T *hData = (T*) malloc (size); + memset(hData, 0, size); + for (int j = 0; j < width; j++) + { + initVal(hData[j]); + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + + hipArray *hipArray = nullptr; + HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore)); + + HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + T *hOutputData = nullptr; + HIP_CHECK(hipHostMalloc((void**)&hOutputData, size)); + memset(hOutputData, 0, size); + + dim3 dimBlock (16, 1, 1); + dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1); + + surf1DKernelR<<>>(surfaceObject, hOutputData, width); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int j = 0; j < width; j++) { + if (!isEqual(hData[j], hOutputData[j])) { + printf("Difference [ %d ]:%s ----%s\n", j, + getString(hData[j]).c_str(), getString(hOutputData[j]).c_str()); + REQUIRE(false); + } + } + + HIP_CHECK(hipDestroySurfaceObject(surfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + free(hData); + HIP_CHECK(hipHostFree(hOutputData)); + REQUIRE(true); +} + +template +static void runTestW(const int width) +{ + unsigned int size = width * sizeof(T); + T *hData = nullptr; + HIP_CHECK(hipHostMalloc((void**)&hData, size)); + memset(hData, 0, size); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + + hipArray *hipArray = nullptr; + HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore)); + + HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + for (int j = 0; j < width; j++) + { + initVal(hData[j]); + } + + dim3 dimBlock (16, 1, 1); + dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1); + + surf1DKernelW<<>>(surfaceObject, hData, width); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + T *hOutputData = (T*) malloc (size); + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpyFromArray(hOutputData, hipArray, 0, 0, size, hipMemcpyDeviceToHost)); + + for (int j = 0; j < width; j++) { + if (!isEqual(hData[j], hOutputData[j])) { + printf("Difference [ %d ]:%s ----%s\n", j, + getString(hData[j]).c_str(), getString(hOutputData[j]).c_str()); + REQUIRE(false); + } + } + + HIP_CHECK(hipDestroySurfaceObject(surfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipHostFree(hData)); + free(hOutputData); + REQUIRE(true); +} + + +template +static void runTestRW(const int width) +{ + unsigned int size = width * sizeof(T); + T *hData = (T*) malloc (size); + memset(hData, 0, size); + for (int j = 0; j < width; j++) + { + initVal(hData[j]); + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + + hipArray *hipArray = nullptr, *hipOutArray = nullptr; + HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore)); + + HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + HIP_CHECK(hipMallocArray(&hipOutArray, &channelDesc, width, 0, hipArraySurfaceLoadStore)); + + hipResourceDesc resOutDesc; + memset(&resOutDesc, 0, sizeof(resOutDesc)); + resOutDesc.resType = hipResourceTypeArray; + resOutDesc.res.array.array = hipOutArray; + + hipSurfaceObject_t outSurfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject (&outSurfaceObject, &resOutDesc)); + + dim3 dimBlock (16, 1, 1); + dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1); + + surf1DKernelRW<<>>(surfaceObject, outSurfaceObject, width); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + T *hOutputData = (T*) malloc (size); + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpyFromArray(hOutputData, hipOutArray, 0, 0, size, hipMemcpyDeviceToHost)); + + for (int j = 0; j < width; j++) { + if (!isEqual(hData[j], hOutputData[j])) { + printf("Difference [ %d ]:%s ----%s\n", j, + getString(hData[j]).c_str(), getString(hOutputData[j]).c_str()); + REQUIRE(false); + } + } + + HIP_CHECK(hipDestroySurfaceObject(surfaceObject)); + HIP_CHECK(hipDestroySurfaceObject(outSurfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipFreeArray(hipOutArray)); + free(hData); + free(hOutputData); + REQUIRE(true); +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_R", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj1D_type_R - 31") { + runTestR(31); + } + + SECTION("Unit_hipSurfaceObj1D_type_R - 67") { + runTestR(67); + } + + SECTION("Unit_hipSurfaceObj1D_type_R - 131") { + runTestR(131); + } + + SECTION("Unit_hipSurfaceObj1D_type_R - 263") { + runTestR(263); + } +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_W", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj1D_type_W - 31") { + runTestW(31); + } + + SECTION("Unit_hipSurfaceObj1D_type_W - 63") { + runTestW(63); + } + + SECTION("Unit_hipSurfaceObj1D_type_W - 131") { + runTestW(131); + } + + SECTION("Unit_hipSurfaceObj1D_type_W - 263") { + runTestW(263); + } +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_RW", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj1D_type_RW - 23") { + runTestRW(23); + } + + SECTION("Unit_hipSurfaceObj1D_type_RW - 67") { + runTestRW(67); + } + + SECTION("Unit_hipSurfaceObj1D_type_RW - 131") { + runTestRW(131); + } + + SECTION("Unit_hipSurfaceObj1D_type_RW - 263") { + runTestRW(263); + } +} diff --git a/catch/unit/surface/hipSurfaceObj2D.cc b/catch/unit/surface/hipSurfaceObj2D.cc new file mode 100644 index 0000000000..22f143e7df --- /dev/null +++ b/catch/unit/surface/hipSurfaceObj2D.cc @@ -0,0 +1,364 @@ +/* +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 + +#define LOG_DATA 0 + +template +__global__ void +surf2DKernelR(hipSurfaceObject_t surfaceObject, + T* outputData, int width, int height) +{ +#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; + if (x < width && y < height) { + surf2Dread(outputData + y * width + x, surfaceObject, x * sizeof(T), y); + } +#endif +} + +template +__global__ void +surf2DKernelW(hipSurfaceObject_t surfaceObject, + T* inputData, int width, int height) +{ +#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; + if (x < width && y < height) { + surf2Dwrite(inputData[y * width + x], surfaceObject, x * sizeof(T), y); + } +#endif +} + +template +__global__ void +surf2DKernelRW(hipSurfaceObject_t surfaceObject, + hipSurfaceObject_t outputSurfObj, int width, int height) +{ +#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; + if (x < width && y < height) { + T data; + surf2Dread(&data, surfaceObject, x * sizeof(T), y); + surf2Dwrite(data, outputSurfObj, x * sizeof(T), y); + } +#endif +} + +template +static void runTestR(const int width, const int height) +{ + unsigned int size = width * height * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) + { + for (int j = 0; j < width; j++) + { + initVal(hData[i * width + j]); + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray = nullptr; + HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height, + hipArraySurfaceLoadStore)); + + // Need set source pitch, but we don't have any padding here + const size_t spitch = width * sizeof(T); + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height, + hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + T* hOutputData = nullptr; + HIP_CHECK(hipHostMalloc((void**)&hOutputData, size)); + memset(hOutputData, 0, size); + + dim3 dimBlock (16, 16, 1); + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1); + surf2DKernelR<<>>(surfaceObject, hOutputData, width, height); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + int index = i * width + j; + if (!isEqual(hData[index], hOutputData[index])) { + printf("Difference [ %d %d ]:%s ----%s\n", i, j, + getString(hData[index]).c_str(), getString(hOutputData[index]).c_str()); + REQUIRE(false); + } + } + } + + HIP_CHECK(hipDestroySurfaceObject(surfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + free(hData); + HIP_CHECK(hipHostFree(hOutputData)); + REQUIRE(true); +} + +template +static void runTestW(const int width, const int height) +{ + unsigned int size = width * height * sizeof(T); + T* hData = nullptr; + HIP_CHECK(hipHostMalloc((void**)&hData, size)); + memset(hData, 0, size); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray = nullptr; + HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height, + hipArraySurfaceLoadStore)); + + // Need set source pitch, but we don't have any padding here + const size_t spitch = width * sizeof(T); + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height, + hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + for (int i = 0; i < height; i++) + { + for (int j = 0; j < width; j++) + { + initVal(hData[i * width + j]); + } + } + + dim3 dimBlock (16, 16, 1); + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1); + surf2DKernelW<<>>(surfaceObject, hData, width, height); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + T* hOutputData = (T*) malloc(size); + + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpy2DFromArray(hOutputData, spitch, hipArray, 0, 0, spitch, + height, hipMemcpyDeviceToHost)); + + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + int index = i * width + j; + if (!isEqual(hData[index], hOutputData[index])) { + printf("Difference [ %d %d ]:%s ----%s\n", i, j, + getString(hData[index]).c_str(), getString(hOutputData[index]).c_str()); + REQUIRE(false); + } + } + } + + HIP_CHECK(hipDestroySurfaceObject(surfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipHostFree(hData)); + free(hOutputData); + REQUIRE(true); +} + +template +static void runTestRW(const int width, const int height) +{ + unsigned int size = width * height * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) + { + for (int j = 0; j < width; j++) + { + initVal(hData[i * width + j]); + } + } +#if LOG_DATA + printf ("hData: "); + for (int i = 0; i < 32; i++) + { + printf ("%s ", getString(hData[i]).c_str()); + } + printf ("\n"); +#endif + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray = nullptr, *hipOutArray = nullptr; + HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height, + hipArraySurfaceLoadStore)); + + // Need set source pitch, but we don't have any padding here + const size_t spitch = width * sizeof(T); + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height, + hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + HIP_CHECK(hipMallocArray(&hipOutArray, &channelDesc, width, height, + hipArraySurfaceLoadStore)); + + hipResourceDesc resOutDesc; + memset(&resOutDesc, 0, sizeof(resOutDesc)); + resOutDesc.resType = hipResourceTypeArray; + resOutDesc.res.array.array = hipOutArray; + + hipSurfaceObject_t outSurfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject (&outSurfaceObject, &resOutDesc)); + + dim3 dimBlock (16, 16, 1); + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1); + surf2DKernelRW<<>>(surfaceObject, outSurfaceObject, width, height); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + T* hOutputData = (T*) malloc(size); + + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpy2DFromArray(hOutputData, spitch, hipOutArray, 0, 0, spitch, + height, hipMemcpyDeviceToHost)); + +#if LOG_DATA + printf ("dData: "); + for (int i = 0; i < 32; i++) + { + printf ("%s ", getString(hOutputData[i]).c_str()); + } + printf ("\n"); +#endif + + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + int index = i * width + j; + if (!isEqual(hData[index], hOutputData[index])) { + printf("Difference [ %d %d ]:%s ----%s\n", i, j, + getString(hData[index]).c_str(), getString(hOutputData[index]).c_str()); + REQUIRE(false); + } + } + } + + HIP_CHECK(hipDestroySurfaceObject(surfaceObject)); + HIP_CHECK(hipDestroySurfaceObject(outSurfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipFreeArray(hipOutArray)); + free(hData); + free(hOutputData); + REQUIRE(true); +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_R", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj2D_type_R - 23, 67") { + runTestR(23, 67); + } + + SECTION("Unit_hipSurfaceObj2D_type_R - 67, 23") { + runTestR(67, 23); + } + + SECTION("Unit_hipSurfaceObj2D_type_R - 131, 67") { + runTestR(131, 67); + } + + SECTION("Unit_hipSurfaceObj2D_type_R - 263, 131") { + runTestR(263, 131); + } +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_W", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj2D_type_W - 23, 67") { + runTestW(23, 67); + } + + SECTION("Unit_hipSurfaceObj2D_type_W - 67, 23") { + runTestW(67, 23); + } + + SECTION("Unit_hipSurfaceObj2D_type_W - 131, 67") { + runTestW(131, 67); + } + + SECTION("Unit_hipSurfaceObj2D_type_W - 263, 23") { + runTestW(263, 23); + } +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_RW", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj2D_type_RW - 23, 67") { + runTestRW(23, 67); + } + + SECTION("Unit_hipSurfaceObj2D_type_RW - 67, 131") { + runTestRW(67, 131); + } + + SECTION("Unit_hipSurfaceObj2D_type_RW - 131, 263") { + runTestRW(131, 263); + } + + SECTION("Unit_hipSurfaceObj2D_type_RW - 263, 67") { + runTestRW(263, 67); + } +} diff --git a/catch/unit/surface/hipSurfaceObj3D.cc b/catch/unit/surface/hipSurfaceObj3D.cc new file mode 100644 index 0000000000..804cc5fb81 --- /dev/null +++ b/catch/unit/surface/hipSurfaceObj3D.cc @@ -0,0 +1,400 @@ +/* +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 + +template +__global__ void +surf3DKernelR(hipSurfaceObject_t surfaceObject, + T* outputData, int width, int height, int depth) +{ +#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; + if (x < width && y < height && z < depth) { + surf3Dread(outputData + z * width * height + y * width + x, + surfaceObject, x * sizeof(T), y, z); + } +#endif +} + +template +__global__ void +surf3DKernelW(hipSurfaceObject_t surfaceObject, + T* inputData, int width, int height, int depth) +{ +#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; + if (x < width && y < height && z < depth) { + surf3Dwrite(inputData[z * width * height + y * width + x], + surfaceObject, x * sizeof(T), y, z); + } +#endif +} + +template +__global__ void +surf3DKernelRW(hipSurfaceObject_t surfaceObject, + hipSurfaceObject_t outputSurfObj, int width, int height, int depth) +{ +#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; + if (x < width && y < height && z < depth) { + T data; + surf3Dread(&data, surfaceObject, x * sizeof(T), y, z); + surf3Dwrite(data, outputSurfObj, x * sizeof(T), y, z); + } +#endif +} + +template +static void runTestR(const int width, const int height, const int depth) +{ + unsigned int size = width * height * depth * sizeof(T); + T *hData = (T*) malloc(size); + 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++) { + initVal(hData[i * width * height + j * width + k]); + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray = nullptr; + HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth), + hipArraySurfaceLoadStore)); + + hipMemcpy3DParms myparms; + memset(&myparms, 0, sizeof(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 = hipArray; + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + T *hOutputData = nullptr; + HIP_CHECK(hipHostMalloc((void**)&hOutputData, size)); + memset(hOutputData, 0, size); + + dim3 dimBlock(8, 8, 8); // 512 threads + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, + (depth + dimBlock.z - 1) / dimBlock.z); + + surf3DKernelR<<>>(surfaceObject, hOutputData, width, height, depth); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + int index = i * width * height + j * width + k; + if (!isEqual(hData[index], hOutputData[index])) { + printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k, + getString(hData[index]).c_str(), getString(hOutputData[index]).c_str()); + REQUIRE(false); + } + } + } + } + + HIP_CHECK(hipDestroySurfaceObject (surfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + free(hData); + HIP_CHECK(hipHostFree(hOutputData)); + REQUIRE(true); +} + +template +static void runTestW(const int width, const int height, const int depth) +{ + unsigned int size = width * height * depth * sizeof(T); + T *hData = nullptr; + HIP_CHECK(hipHostMalloc((void**)&hData, size)); + memset(hData, 0, size); + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray = nullptr; + HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth), + hipArraySurfaceLoadStore)); + + hipMemcpy3DParms myparms; + memset(&myparms, 0, sizeof(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 = hipArray; + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + initVal(hData[i * width * height + j * width + k]); + } + } + } + + dim3 dimBlock(8, 8, 8); // 512 threads + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, + (depth + dimBlock.z - 1) / dimBlock.z); + + surf3DKernelW<<>>(surfaceObject, hData, width, height, depth); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + T *hOutputData = (T*) malloc (size); + memset(hOutputData, 0, size); + + memset(&myparms, 0, sizeof(myparms)); + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcArray= hipArray; + myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height); + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + int index = i * width * height + j * width + k; + if (!isEqual(hData[index], hOutputData[index])) { + printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k, + getString(hData[index]).c_str(), getString(hOutputData[index]).c_str()); + REQUIRE(false); + } + } + } + } + + HIP_CHECK(hipDestroySurfaceObject (surfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipHostFree(hData)); + free(hOutputData); + REQUIRE(true); +} + +template +static void runTestRW(const int width, const int height, const int depth) +{ + unsigned int size = width * height * depth * sizeof(T); + T *hData = (T*) malloc(size); + 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++) { + initVal(hData[i * width * height + j * width + k]); + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray = nullptr, *hipOutArray = nullptr; + HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth), + hipArraySurfaceLoadStore)); + + hipMemcpy3DParms myparms; + memset(&myparms, 0, sizeof(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 = hipArray; + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Create surface object + hipSurfaceObject_t surfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc)); + + HIP_CHECK(hipMalloc3DArray(&hipOutArray, &channelDesc, make_hipExtent(width, height, depth), + hipArraySurfaceLoadStore)); + + hipResourceDesc resOutDesc; + memset(&resOutDesc, 0, sizeof(resOutDesc)); + resOutDesc.resType = hipResourceTypeArray; + resOutDesc.res.array.array = hipOutArray; + + hipSurfaceObject_t outSurfaceObject = 0; + HIP_CHECK(hipCreateSurfaceObject(&outSurfaceObject, &resOutDesc)); + + dim3 dimBlock(8, 8, 8); // 512 threads + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, + (depth + dimBlock.z - 1) / dimBlock.z); + + surf3DKernelRW<<>>(surfaceObject, outSurfaceObject, width, height, depth); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + T *hOutputData = (T*) malloc (size); + memset(hOutputData, 0, size); + + memset(&myparms, 0, sizeof(myparms)); + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcArray= hipOutArray; + myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height); + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + int index = i * width * height + j * width + k; + if (!isEqual(hData[index], hOutputData[index])) { + printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k, + getString(hData[index]).c_str(), getString(hOutputData[index]).c_str()); + REQUIRE(false); + } + } + } + } + + HIP_CHECK(hipDestroySurfaceObject (surfaceObject)); + HIP_CHECK(hipDestroySurfaceObject (outSurfaceObject)); + HIP_CHECK(hipFreeArray(hipArray)); + HIP_CHECK(hipFreeArray(hipOutArray)); + free(hData); + free(hOutputData); + REQUIRE(true); +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_R", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj3D_type_R - 31, 67, 131") { + runTestR(31, 67, 131); + } + + SECTION("Unit_hipSurfaceObj3D_type_R - 67, 31, 263") { + runTestR(67, 31, 263); + } + + SECTION("Unit_hipSurfaceObj3D_type_R - 131, 131, 67") { + runTestR(131, 131, 67); + } + + SECTION("Unit_hipSurfaceObj3D_type_R - 263, 131, 263") { + runTestR(263, 131, 263); + } +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_W", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj3D_type_W - 31, 67, 131") { + runTestW(31, 67, 131); + } + + SECTION("Unit_hipSurfaceObj3D_type_W - 67, 67, 31") { + runTestW(67, 67, 31); + } + + SECTION("Unit_hipSurfaceObj3D_type_W - 131, 131, 67") { + runTestW(131, 131, 67); + } + + SECTION("Unit_hipSurfaceObj3D_type_W - 263, 131, 263") { + runTestW(263, 131, 263); + } +} + +TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_RW", "", + 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 + auto err = hipGetLastError(); // reset last err due to previous negative tests + + SECTION("Unit_hipSurfaceObj3D_type_RW - 31, 31, 67") { + runTestRW(31, 31, 67); + } + + SECTION("Unit_hipSurfaceObj3D_type_RW - 67, 67, 31") { + runTestRW(67, 67, 31); + } + + SECTION("Unit_hipSurfaceObj3D_type_RW - 131, 67, 263") { + runTestRW(131, 67, 263); + } + + SECTION("Unit_hipSurfaceObj3D_type_RW - 263, 131, 263") { + runTestRW(263, 131, 263); + } +} diff --git a/catch/unit/texture/hipTextureObjFetchVector.cc b/catch/unit/texture/hipTextureObjFetchVector.cc index b965ac74e8..d31e5fbd7f 100644 --- a/catch/unit/texture/hipTextureObjFetchVector.cc +++ b/catch/unit/texture/hipTextureObjFetchVector.cc @@ -18,6 +18,7 @@ THE SOFTWARE. */ #include +#include #include #include @@ -31,47 +32,6 @@ __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) { #endif } -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> @@ -112,31 +72,6 @@ static inline void printVector(T &val) { 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 int N = 1024; @@ -144,7 +79,7 @@ bool runTest() { // Allocating the required buffer on gpu device T *texBuf, *texBufOut; T val[N], output[N]; - + hipGetLastError(); // Clear err due to negative tests memset(output, 0, sizeof(output)); std::srand(std::time(nullptr)); // use current time as seed for random generator