From 930931383631f684825363c3cb25dcb45effa9b2 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 19 Jul 2022 19:26:56 +0530 Subject: [PATCH] SWDEV-333188 - Port image tests to catch2 (#2793) - Port image tests to catch2. - Disable 3D filter test for gfx9a. - Some other improvement - Temporarily disable linear filter tests on Windows Change-Id: Ieeddac9f8c50aa9b6b12ca9d8fd792a51973b8ca [ROCm/hip-tests commit: 0fc518d281bbf97e13901c7fe0242f262ac700eb] --- .../catch/include/hip_test_common.hh | 10 +- .../catch/include/hip_texture_helper.hh | 227 ++++++++++++++++++ .../catch/unit/texture/CMakeLists.txt | 6 + .../catch/unit/texture/hipBindTex2DPitch.cc | 14 +- .../unit/texture/hipBindTexRef1DFetch.cc | 5 +- .../hipCreateTextureObject_ArgValidation.cc | 2 +- .../texture/hipCreateTextureObject_Array.cc | 4 +- .../texture/hipCreateTextureObject_Linear.cc | 2 +- .../texture/hipCreateTextureObject_Pitch2D.cc | 2 +- .../catch/unit/texture/hipGetChanDesc.cc | 12 +- .../texture/hipNormalizedFloatValueTex.cc | 10 +- .../unit/texture/hipSimpleTexture2DLayered.cc | 4 + .../catch/unit/texture/hipSimpleTexture3D.cc | 11 +- .../unit/texture/hipTex1DFetchCheckModes.cc | 31 ++- .../catch/unit/texture/hipTexObjPitch.cc | 12 +- .../unit/texture/hipTextureMipmapObj2D.cc | 10 +- .../unit/texture/hipTextureObj1DCheckModes.cc | 148 ++++++++++++ .../unit/texture/hipTextureObj1DFetch.cc | 12 +- .../catch/unit/texture/hipTextureObj2D.cc | 17 +- .../unit/texture/hipTextureObj2DCheckModes.cc | 160 ++++++++++++ .../unit/texture/hipTextureObj3DCheckModes.cc | 218 +++++++++++++++++ .../unit/texture/hipTextureObjFetchVector.cc | 49 ++-- .../catch/unit/texture/hipTextureRef2D.cc | 4 + 23 files changed, 853 insertions(+), 117 deletions(-) create mode 100644 projects/hip-tests/catch/include/hip_texture_helper.hh create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc create mode 100644 projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index fe272f7ace..24f0a1c7dd 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -203,7 +203,7 @@ static inline int RAND_R(unsigned* rand_seed) { inline bool isImageSupported() { int imageSupport = 1; -#ifdef __HIP_PLATFORM_AMD__ +#if HT_AMD int device; HIP_CHECK(hipGetDevice(&device)); HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device)); @@ -276,8 +276,6 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB // This must be called in the beginning of image test app's main() to indicate whether image // is supported. -#define checkImageSupport() \ - if (!HipTest::isImageSupported()) { \ - printf("Texture is not support on the device. Skipped.\n"); \ - return; \ - } +#define CHECK_IMAGE_SUPPORT \ + if (!HipTest::isImageSupported()) \ + { INFO("Texture is not support on the device. Skipped."); return; } diff --git a/projects/hip-tests/catch/include/hip_texture_helper.hh b/projects/hip-tests/catch/include/hip_texture_helper.hh new file mode 100644 index 0000000000..80d57f8a38 --- /dev/null +++ b/projects/hip-tests/catch/include/hip_texture_helper.hh @@ -0,0 +1,227 @@ +#pragma once + +#define HIP_SAMPLING_VERIFY_EPSILON 0.00001 +// The internal precision varies by the GPU family and sometimes within the family. +// Thus the following threshold is subject to change. +#define HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD 0.05 // 5% for filter mode +#define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD 0.1 + +template +bool hipTextureSamplingVerify(const type outputData, const type expected) { + bool testResult = false; + if (fMode == hipFilterModePoint) { + testResult = outputData == expected; + } else if (fMode == hipFilterModeLinear) { + const type mean = (fabs(outputData) + fabs(expected)) / 2; + const type diff = fabs(outputData - expected); + const type ratio = diff / (mean + HIP_SAMPLING_VERIFY_EPSILON); + if (ratio <= HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) { + testResult = true; + } else if (diff <= HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD) { + // Some small outputs have big ratio due to float operation difference of ALU and GPU + testResult = true; + } + } + return testResult; +} + +// Simulate CTS static AddressingTable sAddressingTable +template +void hipTextureGetAddress(int &value, const int maxValue) +{ + switch(addressMode) + { + case hipAddressModeClamp: + value = value < 0 ? 0 + : (value > maxValue - 1 ? maxValue - 1 : value); + break; + case hipAddressModeBorder: + value = value < -1 ? -1 + : (value > maxValue ? maxValue : value); + break; + default: + break; + } +} + +// Simulate logics in CTS read_image_pixel_float(). +// x, y and z must be returned by hipTextureGetAddress() +template +float hipTextureGetValue(const float *data, const int x, const int width, + const int y = 0, const int height = 0,const int z = 0, const int depth = 0) { + float result = std::numeric_limits::lowest(); + switch (addressMode) { + case hipAddressModeClamp: + if (width > 0) { + if (height == 0 && depth == 0) { + result = data[x]; // 1D + } else if (depth == 0) { + result = data[y * width + x]; // 2D + } else { + result = data[z * width * height + y * width + x]; // 3D + } + } + break; + case hipAddressModeBorder: + if (width > 0) { + if (height == 0 && depth == 0) { + result = (x >= 0 && x < width) ? data[x] : 0; // 1D + } else if (depth == 0) { + result = (x >= 0 && x < width && y >= 0 && y < height) ? + data[y * width + x] : 0; // 2D + } else { + result = (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) ? + data[z * width * height + y * width + x] : 0; // 3D + } + } + break; + default: + break; + } + return result; +} + +template +float getExpectedValue(const int width, float x, const float *data) { + float result = std::numeric_limits::lowest(); + switch (filterMode) { + case hipFilterModePoint: { + int i1 = static_cast(floor(x)); + hipTextureGetAddress < addressMode > (i1, width); + result = hipTextureGetValue < addressMode > (data, i1, width); + } + break; + case hipFilterModeLinear: { + x -= 0.5; + int i1 = static_cast(floor(x)); + int i2 = i1 + 1; + float a = x - i1; + hipTextureGetAddress < addressMode > (i1, width); + hipTextureGetAddress < addressMode > (i2, width); + + float t1 = hipTextureGetValue < addressMode > (data, i1, width); + float t2 = hipTextureGetValue < addressMode > (data, i2, width); + + return (1 - a) * t1 + a * t2; + } + break; + } + return result; +} + +template +float getExpectedValue(const int width, const int height, float x, float y, const float *data) { + float result = std::numeric_limits::lowest(); + switch (filterMode) { + case hipFilterModePoint: { + int i1 = static_cast(floor(x)); + int j1 = static_cast(floor(y)); + hipTextureGetAddress < addressMode > (i1, width); + hipTextureGetAddress < addressMode > (j1, height); + result = hipTextureGetValue < addressMode > (data, i1, width, j1, height); + } + break; + case hipFilterModeLinear: { + x -= 0.5; + y -= 0.5; + + int i1 = static_cast(floor(x)); + int j1 = static_cast(floor(y)); + + int i2 = i1 + 1; + int j2 = j1 + 1; + + float a = x - i1; + float b = y - j1; + + hipTextureGetAddress < addressMode > (i1, width); + hipTextureGetAddress < addressMode > (i2, width); + hipTextureGetAddress < addressMode > (j1, height); + hipTextureGetAddress < addressMode > (j2, height); + + float t11 = hipTextureGetValue < addressMode + > (data, i1, width, j1, height); + float t21 = hipTextureGetValue < addressMode + > (data, i2, width, j1, height); + float t12 = hipTextureGetValue < addressMode + > (data, i1, width, j2, height); + float t22 = hipTextureGetValue < addressMode + > (data, i2, width, j2, height); + + result = (1 - a) * (1 - b) * t11 + a * (1 - b) * t21 + (1 - a) * b * t12 + + a * b * t22; + } + break; + } + return result; +} + +template +float getExpectedValue(const int width, const int height, const int depth, + float x, float y, float z, const float *data) { + float result = std::numeric_limits::lowest(); + switch (filterMode) { + case hipFilterModePoint: { + int i1 = static_cast(floor(x)); + int j1 = static_cast(floor(y)); + int k1 = static_cast(floor(z)); + + hipTextureGetAddress < addressMode > (i1, width); + hipTextureGetAddress < addressMode > (j1, height); + hipTextureGetAddress < addressMode > (k1, depth); + + result = hipTextureGetValue < addressMode > (data, i1, width, j1, height, k1, depth); + } + break; + case hipFilterModeLinear: { + x -= 0.5; + y -= 0.5; + z -= 0.5; + + int i1 = static_cast(floor(x)); + int j1 = static_cast(floor(y)); + int k1 = static_cast(floor(z)); + + int i2 = i1 + 1; + int j2 = j1 + 1; + int k2 = k1 + 1; + + float a = x - i1; + float b = y - j1; + float c = z - k1; + + hipTextureGetAddress < addressMode > (i1, width); + hipTextureGetAddress < addressMode > (i2, width); + hipTextureGetAddress < addressMode > (j1, height); + hipTextureGetAddress < addressMode > (j2, height); + hipTextureGetAddress < addressMode > (k1, depth); + hipTextureGetAddress < addressMode > (k2, depth); + + float t111 = hipTextureGetValue < addressMode + > (data, i1, width, j1, height, k1, depth); + float t211 = hipTextureGetValue < addressMode + > (data, i2, width, j1, height, k1, depth); + float t121 = hipTextureGetValue < addressMode + > (data, i1, width, j2, height, k1, depth); + float t112 = hipTextureGetValue < addressMode + > (data, i1, width, j1, height, k2, depth); + float t122 = hipTextureGetValue < addressMode + > (data, i1, width, j2, height, k2, depth); + float t212 = hipTextureGetValue < addressMode + > (data, i2, width, j1, height, k2, depth); + float t221 = hipTextureGetValue < addressMode + > (data, i2, width, j2, height, k1, depth); + float t222 = hipTextureGetValue < addressMode + > (data, i2, width, j2, height, k2, depth); + + result = + (1 - a) * (1 - b) * (1 - c) * t111 + a * (1 - b) * (1 - c) * t211 + + (1 - a) * b * (1 - c) * t121 + a * b * (1 - c) * t221 + + (1 - a) * (1 - b) * c * t112 + a * (1 - b) * c * t212 + + (1 - a) * b * c * t122 + a * b * c * t222; + + } + break; + } + return result; +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index e9c24a8226..686699a197 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -37,6 +37,12 @@ set(TEST_SRC hipGetChanDesc.cc hipTexObjPitch.cc hipTextureObj1DFetch.cc + hipBindTex2DPitch.cc + hipBindTexRef1DFetch.cc + hipTex1DFetchCheckModes.cc + hipTextureObj1DCheckModes.cc + hipTextureObj2DCheckModes.cc + hipTextureObj3DCheckModes.cc ) hip_add_exe_to_target(NAME TextureTest diff --git a/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc b/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc index c1d82ebfbe..2a7aab6fa6 100644 --- a/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc +++ b/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc @@ -28,15 +28,19 @@ texture tex; // texture object is a kernel argument static __global__ void texture2dCopyKernel(TYPE_t* dst) { - int x = threadIdx.x + blockIdx.x * blockDim.x; - int y = threadIdx.y + blockIdx.y * blockDim.y; - if ( (x < SIZE_W) && (y < SIZE_H) ) { - dst[SIZE_W*y+x] = tex2D(tex, x, y); - } +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = threadIdx.x + blockIdx.x * blockDim.x; + int y = threadIdx.y + blockIdx.y * blockDim.y; + if ( (x < SIZE_W) && (y < SIZE_H) ) { + dst[SIZE_W*y+x] = tex2D(tex, x, y); + } +#endif } TEST_CASE("Unit_hipBindTexture2D_Pitch") { + CHECK_IMAGE_SUPPORT + TYPE_t* B; TYPE_t* A; TYPE_t* devPtrB; diff --git a/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc b/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc index e714958163..4e96abc806 100644 --- a/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc +++ b/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc @@ -24,14 +24,17 @@ THE SOFTWARE. texture tex; static __global__ void kernel(float *out) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; if (x < N) { out[x] = tex1Dfetch(tex, x); } +#endif } - TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") { + CHECK_IMAGE_SUPPORT + float *texBuf; float val[N], output[N]; size_t offset = 0; diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc index c7570aea03..0ce2d51f2d 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc @@ -25,7 +25,7 @@ THE SOFTWARE. * Validate argument list of texture object api. */ TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { - checkImageSupport(); + CHECK_IMAGE_SUPPORT float *texBuf; hipError_t ret; diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc index da42c7878e..366547d655 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Array.cc @@ -23,7 +23,7 @@ THE SOFTWARE. * Validates Array Resource texture object with negative/functional tests. */ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { - checkImageSupport(); + CHECK_IMAGE_SUPPORT hipError_t ret; hipResourceDesc resDesc; @@ -49,7 +49,7 @@ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { * with negative/functional tests. */ TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") { - checkImageSupport(); + CHECK_IMAGE_SUPPORT hipError_t ret; hipResourceDesc resDesc; diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc index 0e2caf6a80..fb12586d37 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Linear.cc @@ -26,7 +26,7 @@ THE SOFTWARE. * Validates Linear Resource texture object with negative/functional tests. */ TEST_CASE("Unit_hipCreateTextureObject_LinearResource") { - checkImageSupport(); + CHECK_IMAGE_SUPPORT float *texBuf; hipError_t ret; diff --git a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc index 7fc593708c..fff56d7dd7 100644 --- a/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc @@ -29,7 +29,7 @@ THE SOFTWARE. * Validates Pitch2D Resource texture object with negative and functional tests */ TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { - checkImageSupport(); + CHECK_IMAGE_SUPPORT hipError_t ret; hipResourceDesc resDesc; diff --git a/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc b/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc index e5c37412ce..b7585059c8 100644 --- a/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc +++ b/projects/hip-tests/catch/unit/texture/hipGetChanDesc.cc @@ -24,19 +24,11 @@ THE SOFTWARE. TEST_CASE("Unit_hipGetChannelDesc_CreateAndGet") { + CHECK_IMAGE_SUPPORT + 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)); diff --git a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc index 8113117a64..a6eceef38d 100644 --- a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc +++ b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -142,16 +142,10 @@ static void runTest_hipTextureFilterMode() { } TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") { + CHECK_IMAGE_SUPPORT + #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); diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc index 0c08c0f6d2..fdf28458d2 100644 --- a/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc @@ -27,13 +27,17 @@ texture tex2DL; __global__ void simpleKernelLayeredArray(T* outputData, int width, int height, int layer) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT 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); +#endif } TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") { + CHECK_IMAGE_SUPPORT + constexpr int SIZE = 512; constexpr int num_layers = 5; constexpr unsigned int width = SIZE; diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc index 1ff46a1f2e..3d59101d51 100644 --- a/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc @@ -110,15 +110,8 @@ static void runSimpleTexture3D_Check(int width, int height, int depth, } 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 + CHECK_IMAGE_SUPPORT + for ( int i = 1; i < 25; i++ ) { runSimpleTexture3D_Check(i, i, i, &texf); runSimpleTexture3D_Check(i+1, i, i, &texi); diff --git a/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc index eb3501191c..c3249411b6 100644 --- a/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc @@ -23,9 +23,11 @@ THE SOFTWARE. #define offset 3 static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < (N - offset)) val[k] = tex1Dfetch(obj, k+offset); +#endif } @@ -84,7 +86,8 @@ static void runTest(hipTextureAddressMode addressMode, for (int i = 0; i < (N - offset); i++) { if (output[i] != val[i + offset]) { - INFO("Output not matching at index " << i); + INFO("Mismatch found at output[" << i << "]:" << output[i] << + " val[" << i + offset << "]:" << val[i + offset]); REQUIRE(false); } } @@ -103,16 +106,18 @@ static void runTest(hipTextureAddressMode addressMode, 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); - } + CHECK_IMAGE_SUPPORT + + 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); + } } diff --git a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc index 22a58e986d..727d8a64f2 100644 --- a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc +++ b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc @@ -38,21 +38,13 @@ static __global__ void texture2dCopyKernel(hipTextureObject_t texObj, TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int, unsigned char, int16_t, char, unsigned int) { + CHECK_IMAGE_SUPPORT + 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++) { diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc index 7a55f87ca4..7545f2a073 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc @@ -121,14 +121,8 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int #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 + CHECK_IMAGE_SUPPORT + #ifdef _WIN32 for (auto& hw : hw_vector) { for (auto& mip : mip_vector) { diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc new file mode 100644 index 0000000000..7ff602ab79 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc @@ -0,0 +1,148 @@ +/* +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 tex1DKernel(float *outputData, hipTextureObject_t textureObject, + int width, float offsetX) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int x = blockIdx.x * blockDim.x + threadIdx.x; + outputData[x] = tex1D(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX); +#endif +} + +template +void runTest(const int width, const float offsetX) { + //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, offsetX=%f)\n", __FUNCTION__, + // addressMode, filterMode, normalizedCoords, width, offsetX); + unsigned int size = width * sizeof(float); + float *hData = (float*) malloc(size); + memset(hData, 0, size); + for (int j = 0; j < width; j++) { + hData[j] = j; + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc( + 32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray *hipArray; + hipMallocArray(&hipArray, &channelDesc, width); + + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), 1, 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] = addressMode; + texDesc.filterMode = filterMode; + texDesc.readMode = hipReadModeElementType; + texDesc.normalizedCoords = normalizedCoords; + + // Create texture object + hipTextureObject_t textureObject = 0; + HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL)); + + float *dData = nullptr; + hipMalloc((void**) &dData, size); + + dim3 dimBlock(16, 1, 1); + dim3 dimGrid((width + dimBlock.x - 1)/ dimBlock.x, 1, 1); + + hipLaunchKernelGGL(tex1DKernel, dimGrid, dimBlock, 0, 0, dData, + textureObject, width, offsetX); + + hipDeviceSynchronize(); + + float *hOutputData = (float*) malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + bool result = true; + for (int j = 0; j < width; j++) { + float expectedValue = getExpectedValue(width, offsetX + j, hData); + if (!hipTextureSamplingVerify(hOutputData[j], expectedValue)) { + INFO("Mismatch at " << offsetX + j << ":" << hOutputData[j] << + " expected:" << expectedValue); + result = false; + break; + } + } + + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); + free(hData); + free(hOutputData); + REQUIRE(result); +} + +TEST_CASE("Unit_hipTextureObj1DCheckModes") { + CHECK_IMAGE_SUPPORT + +#ifdef _WIN32 + INFO("Unit_hipTextureObj1DCheckModes skipped on Windows"); + return; +#endif + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest(256, -3); + runTest(256, 4); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest(256, -8.5); + runTest(256, 12.5); + } +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc index a52b8baaf4..a7e25c6a40 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc @@ -22,14 +22,18 @@ THE SOFTWARE. #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); - } +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT + int k = blockIdx.x * blockDim.x + threadIdx.x; + if (k < N) { + val[k] = tex1Dfetch(obj, k); + } +#endif } TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") { + CHECK_IMAGE_SUPPORT + // Allocating the required buffer on gpu device float *texBuf, *texBufOut; float val[N], output[N]; diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc index 1836981e41..48efb13bb1 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc @@ -22,22 +22,15 @@ THE SOFTWARE. __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); + 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 + CHECK_IMAGE_SUPPORT + constexpr int SIZE = 256; constexpr unsigned int width = SIZE; constexpr unsigned int height = SIZE; diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc new file mode 100644 index 0000000000..34696b32e3 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc @@ -0,0 +1,160 @@ +/* +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 tex2DKernel(float *outputData, hipTextureObject_t textureObject, + int width, int height, float offsetX, + float offsetY) { +#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, + normalizedCoords ? (x + offsetX) / width : x + offsetX, + normalizedCoords ? (y + offsetY) / height : y + offsetY); +#endif +} + +template +void runTest(const int width, const int height, const float offsetX, const float offsetY) { + //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, offsetX=%f, offsetY=%f)\n", + // __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, offsetX, offsetY); + unsigned int size = width * height * sizeof(float); + float *hData = (float*) malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + int index = i * width + j; + hData[index] = index; + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc( + 32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray *hipArray; + hipMallocArray(&hipArray, &channelDesc, width, height); + + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(float), width * sizeof(float), height, 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] = addressMode; + texDesc.addressMode[1] = addressMode; + texDesc.filterMode = filterMode; + texDesc.readMode = hipReadModeElementType; + texDesc.normalizedCoords = normalizedCoords; + + // Create texture object + hipTextureObject_t textureObject = 0; + HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL)); + + float *dData = nullptr; + hipMalloc((void**) &dData, size); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dimGrid, dimBlock, 0, 0, dData, + textureObject, width, height, offsetX, offsetY); + + hipDeviceSynchronize(); + + float *hOutputData = (float*) malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + bool result = true; + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + int index = i * width + j; + float expectedValue = getExpectedValue(width, height, + offsetX + j, offsetY + i, hData); + if (!hipTextureSamplingVerify(hOutputData[index], expectedValue)) { + INFO("Mismatch at (" << offsetX + j << ", " << offsetY + i << "):" << + hOutputData[index] << " expected:" << expectedValue); + result = false; + goto line1; + } + } + } +line1: + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); + free(hData); + free(hOutputData); + REQUIRE(result); +} + +TEST_CASE("Unit_hipTextureObj2DCheckModes") { + CHECK_IMAGE_SUPPORT + +#ifdef _WIN32 + INFO("Unit_hipTextureObj2DCheckModes skipped on Windows"); + return; +#endif + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest(256, 256, -3.9, 6.1); + runTest(256, 256, 4.4, -7.0); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest(256, 256, -8.5, 2.9); + runTest(256, 256, 12.5, 6.7); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest(256, 256, -0.4, -0.4); + runTest(256, 256, 4, 14.6); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest(256, 256, -0.4, 0.4); + runTest(256, 256, 12.5, 23.7); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -3, 8.9); + runTest(256, 256, 4, -0.1); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -8.5, 15.9); + runTest(256, 256, 12.5, -17.9); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -3, 5.8); + runTest(256, 256, 4, 9.1); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -8.5, 6.6); + runTest(256, 256, 12.5, 0.01); + } +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc new file mode 100644 index 0000000000..df5132972b --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc @@ -0,0 +1,218 @@ +/* +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 + +bool isGfx90a = false; + +template +__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject, + int width, int height, int depth, float offsetX, + float offsetY, float offsetZ) { +#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; + outputData[z * width * height + y * width + x] = tex3D(textureObject, + normalizedCoords ? (x + offsetX) / width : x + offsetX, + normalizedCoords ? (y + offsetY) / height : y + offsetY, + normalizedCoords ? (z + offsetZ) / depth : z + offsetZ); +#endif +} + +template +void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY, + const float offsetZ) { + //printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, offsetX=%f, offsetY=%f, offsetZ=%f)\n", + // __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height, + // depth, offsetX, offsetY, offsetZ); + bool result = true; + unsigned int size = width * height * depth * sizeof(float); + float *hData = (float*) 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++) { + int index = i * width * height + j * width + k; + hData[index] = index; + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *arr; + + HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + 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(float), width, height); + myparms.dstArray = arr; + 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 = arr; + + // Specify texture object parameters + hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = addressMode; + texDesc.addressMode[1] = addressMode; + texDesc.addressMode[2] = addressMode; + texDesc.filterMode = filterMode; + texDesc.readMode = hipReadModeElementType; + texDesc.normalizedCoords = normalizedCoords; + + // Create texture object + hipTextureObject_t textureObject = 0; + hipError_t res = hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + if (res != hipSuccess) { + hipFreeArray(arr); + free(hData); + if (res == hipErrorNotSupported && isGfx90a) { + printf("gfx90a doesn't support 3D linear filter! Skipped!\n"); + } else { + result = false; + } + REQUIRE(result); + return; + } + + float *dData = nullptr; + hipMalloc((void**) &dData, size); + hipMemset(dData, 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); + + hipLaunchKernelGGL(tex3DKernel, dimGrid, dimBlock, 0, 0, dData, + textureObject, width, height, depth, offsetX, offsetY, offsetZ); + + hipDeviceSynchronize(); + + float *hOutputData = (float*) malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + 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; + float expectedValue = getExpectedValue( + width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData); + + if (!hipTextureSamplingVerify(hOutputData[index], expectedValue)) { + INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i << "):" << + hOutputData[index] << " expected:" << expectedValue); + result = false; + goto line1; + } + } + } + } +line1: + hipDestroyTextureObject(textureObject); + free(hOutputData); + hipFree(dData); + hipFreeArray(arr); + free(hData); + REQUIRE(result); + +} + +TEST_CASE("Unit_hipTextureObj3DCheckModes") { + CHECK_IMAGE_SUPPORT +#ifdef _WIN32 + INFO("Unit_hipTextureObj3DCheckModes skipped on Windows"); + return; +#endif + + int device = 0; + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + if (!strncmp(props.gcnArchName, "gfx90a", strlen("gfx90a"))) { + isGfx90a = true; + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest + (256, 256, 256, -3.9, 6.1, 9.5); + runTest + (256, 256, 256, 4.4, -7.0, 5.3); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest + (256, 256, 256, -8.5, 2.9, 5.8); + runTest + (256, 256, 256, 12.5, 6.7, 11.4); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") { + runTest + (256, 256, 256, -0.4, -0.4, -0.4); + runTest + (256, 256, 256, 4, 14.6, -0.3); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest + (256, 256, 256, 6.9, 7.4, 0.4); + runTest + (256, 256, 256, 12.5, 23.7, 0.34); + } + + SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest + (256, 256, 256, -3, 8.9, -4); + runTest + (256, 256, 256, 4, -0.1, 8.2); + } + + SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest + (256, 256, 256, -8.5, 15.9, 0.1); + runTest + (256, 256, 256, 12.5, -17.9, -0.35); + } + + SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest + (256, 256, 256, -3, 5.8, 0.89); + runTest + (256, 256, 256, 4, 9.1, 2.08); + } + + SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest + (256, 256, 256, -8.5, 6.6, 3.67); + runTest + (256, 256, 256, 12.5, 0.01, -9.9); + } +} diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc b/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc index d3faa5d5b0..616602038a 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc @@ -138,14 +138,12 @@ static inline bool isEqual(const T &val0, const T &val1) { } template -bool runTest(const char *description) { +bool runTest() { 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 @@ -199,46 +197,45 @@ bool runTest(const char *description) { HIP_CHECK(hipFree(texBuf)); HIP_CHECK(hipFree(texBufOut)); - printf(": %s\n", testResult ? "succeeded" : "failed"); REQUIRE(testResult == true); return testResult; } TEST_CASE("Unit_hipTextureFetch_vector") { - checkImageSupport(); + CHECK_IMAGE_SUPPORT // test for char - runTest("char1"); - runTest("char2"); - runTest("char4"); + runTest(); + runTest(); + runTest(); // test for uchar - runTest("uchar1"); - runTest("uchar2"); - runTest("uchar4"); + runTest(); + runTest(); + runTest(); // test for short - runTest("short1"); - runTest("short2"); - runTest("short4"); + runTest(); + runTest(); + runTest(); // test for ushort - runTest("ushort1"); - runTest("ushort2"); - runTest("ushort4"); + runTest(); + runTest(); + runTest(); // test for int - runTest("int1"); - runTest("int2"); - runTest("int4"); + runTest(); + runTest(); + runTest(); // test for unsigned int - runTest("uint1"); - runTest("uint2"); - runTest("uint4"); + runTest(); + runTest(); + runTest(); // test for float - runTest("float1"); - runTest("float2"); - runTest("float4"); + runTest(); + runTest(); + runTest(); } diff --git a/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc index 1a6600d099..a4abcd4975 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc @@ -22,12 +22,16 @@ THE SOFTWARE. texture tex; __global__ void tex2DKernel(float* outputData, 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(tex, x, y); +#endif } TEST_CASE("Unit_hipTextureRef2D_Check") { + CHECK_IMAGE_SUPPORT + constexpr int SIZE = 256; constexpr unsigned int width = SIZE; constexpr unsigned int height = SIZE;