diff --git a/tests/src/test_common.cpp b/tests/src/test_common.cpp index 05aa71da4e..03464fcba5 100644 --- a/tests/src/test_common.cpp +++ b/tests/src/test_common.cpp @@ -37,6 +37,7 @@ char memsetD8val = 0xDE; int iterations = 1; unsigned blocksPerCU = 6; // to hide latency unsigned threadsPerBlock = 256; +int textureFilterMode = 0; // 0: hipFilterModePoint; 1: hipFilterModeLinear int p_gpuDevice = 0; unsigned p_verbose = 0; int p_tests = -1; /*which tests to run. Interpretation is left to each test. default:all*/ @@ -176,11 +177,16 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg) { failed("Bad memsetD8val argument"); } memsetD8val = ex; + } else if (!strcmp(arg, "--textureFilterMode")) { + int mode; + if (++i >= argc || !HipTest::parseInt(argv[i], &mode)) { + failed("Bad textureFilterMode argument"); + } + textureFilterMode = mode; } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) { failed("Bad iterations argument"); } - } else if (!strcmp(arg, "--gpu") || (!strcmp(arg, "-gpuDevice")) || (!strcmp(arg, "-g"))) { if (++i >= argc || !HipTest::parseInt(argv[i], &p_gpuDevice)) { failed("Bad gpuDevice argument"); diff --git a/tests/src/test_common.h b/tests/src/test_common.h index 40acdf4bc1..b6dcba4538 100755 --- a/tests/src/test_common.h +++ b/tests/src/test_common.h @@ -153,6 +153,7 @@ extern char memsetD8val; extern int iterations; extern unsigned blocksPerCU; extern unsigned threadsPerBlock; +extern int textureFilterMode; extern int p_gpuDevice; extern unsigned p_verbose; extern int p_tests; diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index 09a9246c01..f2f1937132 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -22,14 +22,18 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../test_common.cpp - * TEST: %t + * // Test hipFilterModePoint + * TEST: %t --textureFilterMode 0 + * // Test hipFilterModeLinear + * TEST: %t --textureFilterMode 1 * HIT_END */ #include "test_common.h" #include #define SIZE 10 -#define THRESH_HOLD 0.02 +#define EPSILON 0.00001 +#define THRESH_HOLD 0.01 // For filter mode static float getNormalizedValue(const float value, const hipChannelFormatDesc& desc) { @@ -70,7 +74,57 @@ __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) pDst[elementID] = tex1D(texus, coord); } -template +bool textureVerifyFilterModePoint(float *hOutputData, float *expected, size_t size) { + bool testResult = true; + for (int i = 0; i < size; i++) { + if ((hOutputData[i] == expected[i]) + || (i >= 1 && hOutputData[i] == expected[i - 1]) || // round down + (i < (size - 1) && hOutputData[i] == expected[i + 1])) // round up + { + continue; + } + printf("mismatch at output[%d]:%f expected[%d]:%f", i, hOutputData[i], i, + expected[i]); + if (i >= 1) { + printf(", expected[%d]:%f", i - 1, expected[i - 1]); + } + if (i < (size - 1)) { + printf(", expected[%d]:%f", i + 1, expected[i + 1]); + } + printf("\n"); + testResult = false; + break; + } + return testResult; +} + +bool textureVerifyFilterModeLinear(float *hOutputData, float *expected, size_t size) { + 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) { + printf("mismatch at output[%d]:%f expected[%d]:%f, ratio:%f\n", i, + hOutputData[i], i, expected[i], ratio); + testResult = false; + break; + } + } + return testResult; +} + +template +bool textureVerify(float *hOutputData, float *expected, size_t size) { + bool testResult = true; + if (fMode == hipFilterModePoint) { + testResult = textureVerifyFilterModePoint(hOutputData, expected, size); + } else if (fMode == hipFilterModeLinear) { + testResult = textureVerifyFilterModeLinear(hOutputData, expected, size); + } + return testResult; +} + +template bool textureTest(texture *tex) { hipChannelFormatDesc desc = hipCreateChannelDesc(); @@ -82,6 +136,7 @@ bool textureTest(texture *tex) tex->normalized = true; tex->channelDesc = desc; + tex->filterMode = fMode; HIPCHECK(hipBindTextureToArray(tex, dData, &desc)); float *dOutputData = NULL; @@ -92,19 +147,11 @@ bool textureTest(texture *tex) float *hOutputData = new float[SIZE]; HIPCHECK(hipMemcpy(hOutputData, dOutputData, (sizeof(float)*SIZE), hipMemcpyDeviceToHost)); - bool testResult = true; - for(int i = 0; i < SIZE; i++) - { - float expected = getNormalizedValue(float(hData[i]), desc); - float mean = (fabs(expected) + fabs(hOutputData[i])) / 2; - float ratio = fabs(expected - hOutputData[i]) / mean; - if(ratio > THRESH_HOLD) - { - printf("mismatch at index:%d output:%f expected:%f, ratio:%f\n", i, hOutputData[i], expected, ratio); - testResult = false; - break; - } + float expected[SIZE]; + for(int i = 0; i < SIZE; i++) { + expected[i] = getNormalizedValue(float(hData[i]), desc); } + bool testResult = textureVerify(hOutputData, expected, SIZE); HIPCHECK(hipFreeArray(dData)); HIPCHECK(hipFree(dOutputData)); @@ -112,10 +159,21 @@ bool textureTest(texture *tex) return testResult; } +template +bool runTest() { + bool status = true; + status &= textureTest(&texc); + status &= textureTest(&texuc); + status &= textureTest(&texs); + status &= textureTest(&texus); + return status; +} + int main(int argc, char** argv) { + HipTest::parseStandardArguments(argc, argv, true); int device = 0; - bool status = true; + bool status = false; HIPCHECK(hipSetDevice(device)); hipDeviceProp_t props; HIPCHECK(hipGetDeviceProperties(&props, device)); @@ -124,12 +182,18 @@ int main(int argc, char** argv) std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl; #endif - printf("THRESH_HOLD:%f\n", THRESH_HOLD); - - status &= textureTest (&texc); - status &= textureTest (&texuc); - status &= textureTest (&texs); - status &= textureTest(&texus); + if(textureFilterMode == 0) { + printf("Test hipFilterModePoint\n"); + status = runTest(); + } else if(textureFilterMode == 1) { + printf("Test hipFilterModeLinear\n"); + printf("THRESH_HOLD:%f, EPSILON:%f\n", THRESH_HOLD, EPSILON); + status = runTest(); + } else { + printf("Wrong argument!\n"); + printf("hipNormalizedFloatValueTex --textureFilterMode 0 for hipFilterModePoint\n"); + printf("hipNormalizedFloatValueTex --textureFilterMode 1 for hipFilterModeLinear\n"); + } if(status){ passed();