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"; #--- 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