From e05f98967e3ea7d81456c0d7743d1a3f817d7037 Mon Sep 17 00:00:00 2001 From: TomSang Date: Thu, 17 Feb 2022 16:48:45 -0500 Subject: [PATCH 1/2] SWDEV-306947 - Add mode combination tests (#2447) Address modes are valid for array rather than linear memory. Comment out codes of address test for linear memory. Add tests of mode combination(address, linear and normorlized coordinates) in 1D, 2D and 3D. Change-Id: I9e7e97d9bf619743c311f68f51d5bf115fdd534d --- tests/src/texture/hipBindTexRef1DFetch.cpp | 2 +- .../texture/hipNormalizedFloatValueTex.cpp | 12 +- tests/src/texture/hipTex1DFetchCheckModes.cpp | 21 +- tests/src/texture/hipTextureHelper.hpp | 227 ++++++++++++++++++ .../src/texture/hipTextureObj1DCheckModes.cpp | 121 ++++++++++ .../src/texture/hipTextureObj2DCheckModes.cpp | 136 +++++++++++ .../src/texture/hipTextureObj3DCheckModes.cpp | 153 ++++++++++++ tests/src/texture/simpleTexture3D.cpp | 5 +- 8 files changed, 653 insertions(+), 24 deletions(-) create mode 100644 tests/src/texture/hipTextureHelper.hpp create mode 100644 tests/src/texture/hipTextureObj1DCheckModes.cpp create mode 100644 tests/src/texture/hipTextureObj2DCheckModes.cpp create mode 100644 tests/src/texture/hipTextureObj3DCheckModes.cpp diff --git a/tests/src/texture/hipBindTexRef1DFetch.cpp b/tests/src/texture/hipBindTexRef1DFetch.cpp index b2ea94b0bc..2903bf46a4 100644 --- a/tests/src/texture/hipBindTexRef1DFetch.cpp +++ b/tests/src/texture/hipBindTexRef1DFetch.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_RUNTIME rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index 827ac69902..520e64a89b 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -32,8 +32,7 @@ THE SOFTWARE. #include "test_common.h" #include #define SIZE 10 -#define EPSILON 0.00001 -#define THRESH_HOLD 0.01 // For filter mode +#include "hipTextureHelper.hpp" static float getNormalizedValue(const float value, const hipChannelFormatDesc& desc) { @@ -104,8 +103,8 @@ bool textureVerifyFilterModeLinear(float *hOutputData, float *expected, size_t bool testResult = true; for (int i = 0; i < size; i++) { float mean = (fabs(expected[i]) + fabs(hOutputData[i])) / 2; - float ratio = fabs(expected[i] - hOutputData[i]) / (mean + EPSILON); - if (ratio > THRESH_HOLD) { + float ratio = fabs(expected[i] - hOutputData[i]) / (mean + HIP_SAMPLING_VERIFY_EPSILON); + if (ratio > HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) { printf("mismatch at output[%d]:%f expected[%d]:%f, ratio:%f\n", i, hOutputData[i], i, expected[i], ratio); testResult = false; @@ -131,7 +130,7 @@ bool textureTest(texture *tex) { hipChannelFormatDesc desc = hipCreateChannelDesc(); hipArray_t dData; - HIPCHECK(hipMallocArray(&dData, &desc, SIZE, 1, hipArrayDefault)); + HIPCHECK(hipMallocArray(&dData, &desc, SIZE)); T hData[] = {65, 66, 67, 68, 69, 70, 71, 72, 73, 74}; HIPCHECK(hipMemcpy2DToArray(dData, 0, 0, hData, sizeof(T)*SIZE, sizeof(T)*SIZE, 1, hipMemcpyHostToDevice)); @@ -196,7 +195,8 @@ int main(int argc, char** argv) status = runTest(); } else if(textureFilterMode == 1) { printf("Test hipFilterModeLinear\n"); - printf("THRESH_HOLD:%f, EPSILON:%f\n", THRESH_HOLD, EPSILON); + printf("THRESH_HOLD:%f, EPSILON:%f\n", HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD, + HIP_SAMPLING_VERIFY_EPSILON); status = runTest(); } else { printf("Wrong argument!\n"); diff --git a/tests/src/texture/hipTex1DFetchCheckModes.cpp b/tests/src/texture/hipTex1DFetchCheckModes.cpp index 856eaeb05a..433185c124 100644 --- a/tests/src/texture/hipTex1DFetchCheckModes.cpp +++ b/tests/src/texture/hipTex1DFetchCheckModes.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_RUNTIME rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ @@ -38,9 +38,9 @@ int runTest(hipTextureAddressMode, hipTextureFilterMode); int main(int argc, char **argv) { int testResult = runTest(hipAddressModeClamp,hipFilterModePoint); - testResult = runTest(hipAddressModeClamp,hipFilterModeLinear); - testResult = runTest(hipAddressModeWrap,hipFilterModePoint); - testResult = runTest(hipAddressModeWrap,hipFilterModeLinear); + testResult = testResult & runTest(hipAddressModeClamp,hipFilterModeLinear); + testResult = testResult & runTest(hipAddressModeBorder,hipFilterModePoint); + testResult = testResult & runTest(hipAddressModeBorder,hipFilterModeLinear); if(testResult) { passed(); } else { @@ -84,7 +84,6 @@ int runTest(hipTextureAddressMode addressMode, hipTextureFilterMode filterMode) texDesc.readMode = hipReadModeElementType; texDesc.addressMode[0] = addressMode; - texDesc.addressMode[1] = addressMode; texDesc.filterMode = filterMode; texDesc.normalizedCoords = false; @@ -107,16 +106,12 @@ int runTest(hipTextureAddressMode addressMode, hipTextureFilterMode filterMode) break; } } - if(testResult){ - for(int i = N-offset; i < N; i++){ - if (output[i] != 0){ - testResult = 0; - break; - } - } - } + // For hipResourceTypeLinear, reading of out-of-boundary address is undefined! + // So we won't verify those data + HIPCHECK(hipDestroyTextureObject(texObj)); HIPCHECK(hipFree(texBuf)); HIPCHECK(hipFree(texBufOut)); + printf("%s(addressMode %d, filterMode %d) %s\n", __FUNCTION__, addressMode, filterMode, testResult ? "succeed" : "failed"); return testResult; } diff --git a/tests/src/texture/hipTextureHelper.hpp b/tests/src/texture/hipTextureHelper.hpp new file mode 100644 index 0000000000..26f3675834 --- /dev/null +++ b/tests/src/texture/hipTextureHelper.hpp @@ -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; + } +} + +// Simlate 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/tests/src/texture/hipTextureObj1DCheckModes.cpp b/tests/src/texture/hipTextureObj1DCheckModes.cpp new file mode 100644 index 0000000000..fbcf14cf62 --- /dev/null +++ b/tests/src/texture/hipTextureObj1DCheckModes.cpp @@ -0,0 +1,121 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include +#include +#include + +#include +#include "test_common.h" +#include "hipTextureHelper.hpp" + +template +__global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject, + int width, float offsetX) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + outputData[x] = tex1D(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX); +} + +template +bool 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); + bool testResult = true; + 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); + + HIPCHECK(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; + hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + + float *dData = NULL; + 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); + + for (int j = 0; j < width; j++) { + float expectedValue = getExpectedValue(width, offsetX + j, hData); + if (!hipTextureSamplingVerify(hOutputData[j], expectedValue)) { + printf("mismatched [ %d ]:%f ----%f\n", j, hOutputData[j], expectedValue); + testResult = false; + break; + } + } + + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); + free(hData); + free(hOutputData); + printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed"); + return testResult; +} + +int main(int argc, char **argv) { + bool testResult = true; + testResult = testResult && runTest(256, -3); + testResult = testResult && runTest(256, 4); + + testResult = testResult && runTest(256, -8.5); + testResult = testResult && runTest(256, 12.5); + + testResult = testResult && runTest(256, -3); + testResult = testResult && runTest(256, 4); + + testResult = testResult && runTest(256, -8.5); + testResult = testResult && runTest(256, 12.5); + + testResult = testResult && runTest(256, -3); + testResult = testResult && runTest(256, 4); + + testResult = testResult && runTest(256, -8.5); + testResult = testResult && runTest(256, 12.5); + + testResult = testResult && runTest(256, -3); + testResult = testResult && runTest(256, 4); + + testResult = testResult && runTest(256, -8.5); + testResult = testResult && runTest(256, 12.5); + + if (testResult) { + passed(); + } else { + exit (EXIT_FAILURE); + } +} diff --git a/tests/src/texture/hipTextureObj2DCheckModes.cpp b/tests/src/texture/hipTextureObj2DCheckModes.cpp new file mode 100644 index 0000000000..1c79ff59d3 --- /dev/null +++ b/tests/src/texture/hipTextureObj2DCheckModes.cpp @@ -0,0 +1,136 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include +#include +#include + +#include +#include "test_common.h" +#include "hipTextureHelper.hpp" + +template +__global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject, + int width, int height, float offsetX, + float offsetY) { + 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); +} + +template +bool 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); + bool testResult = true; + 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); + + HIPCHECK(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; + hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + + float *dData = NULL; + 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); + + 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)) { + printf("mismatched [ %d %d ]:%f ----%f\n", j, i, hOutputData[index], expectedValue); + testResult = false; + goto line1; + } + } + } +line1: + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); + free(hData); + free(hOutputData); + printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed"); + return testResult; +} + +int main(int argc, char **argv) { + bool testResult = true; + + testResult = testResult && runTest(256, 256, -3.9, 6.1); + testResult = testResult && runTest(256, 256, 4.4, -7.0); + + testResult = testResult && runTest(256, 256, -8.5, 2.9); + testResult = testResult && runTest(256, 256, 12.5, 6.7); + + testResult = testResult && runTest(256, 256, -0.4, -0.4); + testResult = testResult && runTest(256, 256, 4, 14.6); + + // The following two cases have quite big deviation on Cpu and Gpu in 2D, so comment them out temporarily. + testResult = testResult && runTest(256, 256, -0.4, 0.4); + testResult = testResult && runTest(256, 256, 12.5, 23.7); + + testResult = testResult && runTest(256, 256, -3, 8.9); + testResult = testResult && runTest(256, 256, 4, -0.1); + + testResult = testResult && runTest(256, 256, -8.5, 15.9); + testResult = testResult && runTest(256, 256, 12.5, -17.9); + + testResult = testResult && runTest(256, 256, -3, 5.8); + testResult = testResult && runTest(256, 256, 4, 9.1); + + // The following two cases have quite big deviation on Cpu and Gpu in 2D, so comment them out temporarily. + testResult = testResult && runTest(256, 256, -8.5, 6.6); + testResult = testResult && runTest(256, 256, 12.5, 0.01); + + if (testResult) { + passed(); + } else { + exit (EXIT_FAILURE); + } +} diff --git a/tests/src/texture/hipTextureObj3DCheckModes.cpp b/tests/src/texture/hipTextureObj3DCheckModes.cpp new file mode 100644 index 0000000000..3c71ee2af8 --- /dev/null +++ b/tests/src/texture/hipTextureObj3DCheckModes.cpp @@ -0,0 +1,153 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include +#include +#include + +#include +#include "test_common.h" +#include "hipTextureHelper.hpp" + + +template +__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject, + int width, int height, int depth, float offsetX, + float offsetY, float offsetZ) { + 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 * depth + y * width + x] = tex3D(textureObject, + normalizedCoords ? (x + offsetX) / width : x + offsetX, + normalizedCoords ? (y + offsetY) / height : y + offsetY, + normalizedCoords ? (z + offsetZ) / depth : z + offsetZ); +} + +template +bool 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 testResult = 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 * depth + j * width + k; + hData[index] = index; + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *arr; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); + hipMemcpy3DParms myparms = {0}; + 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; + + HIPCHECK(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; + hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + + float *dData = NULL; + 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 * depth + j * width + k; + float expectedValue = getExpectedValue( + width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData); + + if (!hipTextureSamplingVerify(hOutputData[index], expectedValue)) { + printf("mismatched [ %d %d %d]:%f ----%f\n", k, j, i, hOutputData[index], expectedValue); + testResult = false; + goto line1; + } + } + } + } +line1: + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(arr); + free(hData); + free(hOutputData); + printf("%s %s\n", __FUNCTION__, testResult ? "succeeded":"failed"); + return testResult; +} + +int main(int argc, char **argv) { + bool testResult = true; + + testResult = testResult && runTest(256, 256, 256, -3.9, 6.1, 9.5); + testResult = testResult && runTest(256, 256, 256, 4.4, -7.0, 5.3); + + testResult = testResult && runTest(256, 256, 256, -8.5, 2.9, 5.8); + testResult = testResult && runTest(256, 256, 256, 12.5, 6.7, 11.4); + + testResult = testResult && runTest(256, 256, 256, -0.4, -0.4, -0.4); + testResult = testResult && runTest(256, 256, 256, 4, 14.6, -0.3); + + testResult = testResult && runTest(256, 256, 256, 6.9, 7.4, 0.4); + testResult = testResult && runTest(256, 256, 256, 12.5, 23.7, 0.34); + + testResult = testResult && runTest(256, 256, 256, -3, 8.9, -4); + testResult = testResult && runTest(256, 256, 256, 4, -0.1, 8.2); + + testResult = testResult && runTest(256, 256, 256, -8.5, 15.9, 0.1); + testResult = testResult && runTest(256, 256, 256, 12.5, -17.9, -0.35); + + testResult = testResult && runTest(256, 256, 256, -3, 5.8, 0.89); + testResult = testResult && runTest(256, 256, 256, 4, 9.1, 2.08); + + testResult = testResult && runTest(256, 256, 256, -8.5, 6.6, 3.67); + testResult = testResult && runTest(256, 256, 256, 12.5, 0.01, -9.9); + + if (testResult) { + passed(); + } else { + exit (EXIT_FAILURE); + } +} diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp index 06f2a31c97..b825577546 100644 --- a/tests/src/texture/simpleTexture3D.cpp +++ b/tests/src/texture/simpleTexture3D.cpp @@ -87,11 +87,8 @@ void runTest(int width,int height,int depth,texture Date: Fri, 18 Feb 2022 07:24:18 +0530 Subject: [PATCH 2/2] Bump HIP version to 5.1 --- bin/hipvars.pm | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipvars.pm b/bin/hipvars.pm index 961f45bb68..9a24969ea5 100644 --- a/bin/hipvars.pm +++ b/bin/hipvars.pm @@ -26,7 +26,7 @@ use Cwd; use File::Basename; $HIP_BASE_VERSION_MAJOR = "5"; -$HIP_BASE_VERSION_MINOR = "0"; +$HIP_BASE_VERSION_MINOR = "1"; $HIP_BASE_VERSION_PATCH = "0"; #---