diff --git a/catch/include/hip_texture_helper.hh b/catch/include/hip_texture_helper.hh index 80d57f8a38..d0c9db0dc9 100644 --- a/catch/include/hip_texture_helper.hh +++ b/catch/include/hip_texture_helper.hh @@ -1,4 +1,5 @@ #pragma once +#include #define HIP_SAMPLING_VERIFY_EPSILON 0.00001 // The internal precision varies by the GPU family and sometimes within the family. @@ -6,15 +7,149 @@ #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) { +#if HT_NVIDIA +template +typename std::enable_if::type +inline __host__ __device__ operator+(const T &a, const T &b) +{ + return {a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w}; +} + +template +typename std::enable_if::type +inline __host__ __device__ operator-(const T &a, const T &b) +{ + return {a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w}; +} + +template +typename std::enable_if::type +inline __host__ __device__ operator==(const T &a, const T &b) +{ + return a.x == b.x && a.y == b.y && a.z == b.z && a.w == b.w; +} + +template +typename std::enable_if::type +inline __host__ __device__ operator*(const decltype(T::x) &a, const T &b) +{ + return {a * b.x, a * b.y, a * b.z, a * b.w}; +} + +template +typename std::enable_if::type +inline __host__ __device__ operator*=(T &a, const decltype(T::x) &b) +{ + a.x *= b; + a.y *= b; + a.z *= b; + a.w *= b; +} +#endif // HT_NVIDIA + +// See https://en.wikipedia.org/wiki/SRGB#Transformation +// From CIE 1931 color space to sRGB +inline float hipSRGBMap(float fc) { + double c = static_cast(fc); + +#if !defined(_WIN32) + if (std::isnan(c)) + c = 0.0; +#else + if (_isnan(c)) c = 0.0; +#endif + + if (c > 1.0) + c = 1.0; + else if (c < 0.0) + c = 0.0; + else if (c < 0.0031308) + c = 12.92 * c; + else + c = 1.055 * pow(c, 5.0 / 12.0) - 0.055; + + return static_cast(c); +} + +// From sRGB to CIE 1931 color space +inline float hipSRGBUnmap(float fc) { + double c = static_cast(fc); + + if (c <= 0.04045) + c = c / 12.92; + else + c = pow((c + 0.055) / 1.055, 2.4); + + return static_cast(c); +} + +inline float4 hipSRGBMap(float4 fc) { + fc.x = hipSRGBMap(fc.x); + fc.y = hipSRGBMap(fc.y); + fc.z = hipSRGBMap(fc.z); + // Alpha channel will keep unchanged + return fc; +} + +inline float4 hipSRGBUnmap(float4 fc) { + fc.x = hipSRGBUnmap(fc.x); + fc.y = hipSRGBUnmap(fc.y); + fc.z = hipSRGBUnmap(fc.z); + // Alpha channel will keep unchanged + return fc; +} + +template +typename std::enable_if::value == true, double>::type +hipFabs(const T &t) { + return fabs(t); +} + +template +typename std::enable_if::type +hipFabs(const T &t) { + return fabs(t.x); +} + +template +typename std::enable_if::type +hipFabs(const T &t) { + double x = static_cast(t.x); + double y = static_cast(t.y); + double s = x * x + y * y; + return sqrt(s); +} + +template +typename std::enable_if::type +hipFabs(const T &t) { + double x = static_cast(t.x); + double y = static_cast(t.y); + double z = static_cast(t.z); + double s = x * x + y * y + z * z; + return sqrt(s); +} + +template +typename std::enable_if::type +hipFabs(const T &t) { + double x = static_cast(t.x); + double y = static_cast(t.y); + double z = static_cast(t.z); + double w = static_cast(t.w); + double s = x * x + y * y + z * z + w * w; + return sqrt(s); +} + +template +bool hipTextureSamplingVerify(T outputData, T expected) { bool testResult = false; - if (fMode == hipFilterModePoint) { + if (fMode == hipFilterModePoint && !sRGB) { 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); + } else { + double mean = (hipFabs(outputData) + hipFabs(expected)) / 2; + double diff = hipFabs(outputData - expected); + double ratio = diff / (mean + HIP_SAMPLING_VERIFY_EPSILON); if (ratio <= HIP_SAMPLING_VERIFY_RELATIVE_THRESHOLD) { testResult = true; } else if (diff <= HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD) { @@ -46,10 +181,11 @@ void hipTextureGetAddress(int &value, const int maxValue) // 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(); +template +T hipTextureGetValue(const T *data, const int x, const int width, + const int y = 0, const int height = 0, const int z = 0, const int depth = 0) { + T result; + memset(&result, 0, sizeof(result)); switch (addressMode) { case hipAddressModeClamp: if (width > 0) { @@ -65,30 +201,35 @@ float hipTextureGetValue(const float *data, const int x, const int width, case hipAddressModeBorder: if (width > 0) { if (height == 0 && depth == 0) { - result = (x >= 0 && x < width) ? data[x] : 0; // 1D + if (x >= 0 && x < width) + result = data[x]; // 1D } else if (depth == 0) { - result = (x >= 0 && x < width && y >= 0 && y < height) ? - data[y * width + x] : 0; // 2D + if (x >= 0 && x < width && y >= 0 && y < height) + result = data[y * width + x]; // 2D } else { - result = (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) ? - data[z * width * height + y * width + x] : 0; // 3D + if (x >= 0 && x < width && y >= 0 && y < height && z >= 0 && z < depth) + result = data[z * width * height + y * width + x]; // 3D } } break; default: break; } + if (sRGB && std::is_same::value) { + result = hipSRGBUnmap(result); + } return result; } -template -float getExpectedValue(const int width, float x, const float *data) { - float result = std::numeric_limits::lowest(); +template +T getExpectedValue(const int width, float x, const T *data) { + T result; + memset(&result, 0, sizeof(result)); switch (filterMode) { case hipFilterModePoint: { int i1 = static_cast(floor(x)); hipTextureGetAddress < addressMode > (i1, width); - result = hipTextureGetValue < addressMode > (data, i1, width); + result = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width); } break; case hipFilterModeLinear: { @@ -99,8 +240,8 @@ float getExpectedValue(const int width, float x, const float *data) { hipTextureGetAddress < addressMode > (i1, width); hipTextureGetAddress < addressMode > (i2, width); - float t1 = hipTextureGetValue < addressMode > (data, i1, width); - float t2 = hipTextureGetValue < addressMode > (data, i2, width); + T t1 = hipTextureGetValue < T, addressMode, sRGB> (data, i1, width); + T t2 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width); return (1 - a) * t1 + a * t2; } @@ -109,16 +250,17 @@ float getExpectedValue(const int width, float x, const float *data) { return result; } -template -float getExpectedValue(const int width, const int height, float x, float y, const float *data) { - float result = std::numeric_limits::lowest(); +template +T getExpectedValue(const int width, const int height, float x, float y, const T *data) { + T result; + memset(&result, 0, sizeof(result)); 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); + result = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height); } break; case hipFilterModeLinear: { @@ -139,13 +281,13 @@ float getExpectedValue(const int width, const int height, float x, float y, cons hipTextureGetAddress < addressMode > (j1, height); hipTextureGetAddress < addressMode > (j2, height); - float t11 = hipTextureGetValue < addressMode + T t11 = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height); - float t21 = hipTextureGetValue < addressMode + T t21 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width, j1, height); - float t12 = hipTextureGetValue < addressMode + T t12 = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j2, height); - float t22 = hipTextureGetValue < addressMode + T t22 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width, j2, height); result = (1 - a) * (1 - b) * t11 + a * (1 - b) * t21 + (1 - a) * b * t12 @@ -156,10 +298,11 @@ float getExpectedValue(const int width, const int height, float x, float y, cons 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(); +template +T getExpectedValue(const int width, const int height, const int depth, + float x, float y, float z, const T *data) { + T result; + memset(&result, 0, sizeof(result)); switch (filterMode) { case hipFilterModePoint: { int i1 = static_cast(floor(x)); @@ -170,7 +313,7 @@ float getExpectedValue(const int width, const int height, const int depth, hipTextureGetAddress < addressMode > (j1, height); hipTextureGetAddress < addressMode > (k1, depth); - result = hipTextureGetValue < addressMode > (data, i1, width, j1, height, k1, depth); + result = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height, k1, depth); } break; case hipFilterModeLinear: { @@ -197,21 +340,21 @@ float getExpectedValue(const int width, const int height, const int depth, hipTextureGetAddress < addressMode > (k1, depth); hipTextureGetAddress < addressMode > (k2, depth); - float t111 = hipTextureGetValue < addressMode + T t111 = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height, k1, depth); - float t211 = hipTextureGetValue < addressMode + T t211 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width, j1, height, k1, depth); - float t121 = hipTextureGetValue < addressMode + T t121 = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j2, height, k1, depth); - float t112 = hipTextureGetValue < addressMode + T t112 = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j1, height, k2, depth); - float t122 = hipTextureGetValue < addressMode + T t122 = hipTextureGetValue < T, addressMode, sRGB > (data, i1, width, j2, height, k2, depth); - float t212 = hipTextureGetValue < addressMode + T t212 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width, j1, height, k2, depth); - float t221 = hipTextureGetValue < addressMode + T t221 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width, j2, height, k1, depth); - float t222 = hipTextureGetValue < addressMode + T t222 = hipTextureGetValue < T, addressMode, sRGB > (data, i2, width, j2, height, k2, depth); result = @@ -224,4 +367,4 @@ float getExpectedValue(const int width, const int height, const int depth, break; } return result; -} \ No newline at end of file +} diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index 686699a197..0d33c27a29 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -43,6 +43,7 @@ set(TEST_SRC hipTextureObj1DCheckModes.cc hipTextureObj2DCheckModes.cc hipTextureObj3DCheckModes.cc + hipTextureObj2DCheckSRGBModes.cc ) hip_add_exe_to_target(NAME TextureTest diff --git a/catch/unit/texture/hipTextureObj1DCheckModes.cc b/catch/unit/texture/hipTextureObj1DCheckModes.cc index a1c1ca2756..052967e2f0 100644 --- a/catch/unit/texture/hipTextureObj1DCheckModes.cc +++ b/catch/unit/texture/hipTextureObj1DCheckModes.cc @@ -31,7 +31,7 @@ __global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject, } template -void runTest(const int width, const float offsetX) { +static 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); @@ -82,7 +82,7 @@ void runTest(const int width, const float offsetX) { bool result = true; for (int j = 0; j < width; j++) { - float expectedValue = getExpectedValue(width, offsetX + j, hData); + float expectedValue = getExpectedValue(width, offsetX + j, hData); if (!hipTextureSamplingVerify(hOutputData[j], expectedValue)) { INFO("Mismatch at " << offsetX + j << ":" << hOutputData[j] << " expected:" << expectedValue); diff --git a/catch/unit/texture/hipTextureObj2DCheckModes.cc b/catch/unit/texture/hipTextureObj2DCheckModes.cc index a09946e008..19562f8966 100644 --- a/catch/unit/texture/hipTextureObj2DCheckModes.cc +++ b/catch/unit/texture/hipTextureObj2DCheckModes.cc @@ -35,7 +35,7 @@ __global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject, } template -void runTest(const int width, const int height, const float offsetX, const float offsetY) { +static 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); @@ -92,7 +92,7 @@ void runTest(const int width, const int height, const float offsetX, const float for (int i = 0; i < height; i++) { for (int j = 0; j < width; j++) { int index = i * width + j; - float expectedValue = getExpectedValue(width, height, + float expectedValue = getExpectedValue(width, height, offsetX + j, offsetY + i, hData); if (!hipTextureSamplingVerify(hOutputData[index], expectedValue)) { INFO("Mismatch at (" << offsetX + j << ", " << offsetY + i << "):" << @@ -128,7 +128,7 @@ TEST_CASE("Unit_hipTextureObj2DCheckModes") { runTest(256, 256, 12.5, 6.7); } - SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { + SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") { runTest(256, 256, -0.4, -0.4); runTest(256, 256, 4, 14.6); } diff --git a/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc b/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc new file mode 100644 index 0000000000..5f44e2d7e9 --- /dev/null +++ b/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc @@ -0,0 +1,242 @@ +/* +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 tex2DRGBAKernel(float4 *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 +static 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); + constexpr float uCharMax = UCHAR_MAX; + unsigned int size = width * height * sizeof(uchar4); + uchar4 *hData = (uchar4*) 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].x = static_cast(j); + hData[index].y = static_cast(i); + hData[index].z = static_cast(index); + hData[index].w = static_cast(i + j); + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray *hipArray; + HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height)); + + HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, width * sizeof(uchar4), width * sizeof(uchar4), 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 = hipReadModeNormalizedFloat; + texDesc.normalizedCoords = normalizedCoords; + texDesc.sRGB = sRGB ? 1 : 0; + + // Create texture object + hipTextureObject_t textureObject = 0; + HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL)); + + float4 *dData = nullptr; + size = width * height * sizeof(float4); + HIP_CHECK(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(tex2DRGBAKernel, dimGrid, dimBlock, 0, 0, dData, + textureObject, width, height, offsetX, offsetY); + + HIP_CHECK(hipDeviceSynchronize()); + + float4 *hInputData = (float4*) malloc(size); // CPU expected values + float4 *hOutputData = (float4*) malloc(size); // GPU output values + memset(hInputData, 0, size); + memset(hOutputData, 0, size); + + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + int index = i * width + j; + hInputData[index].x = hData[index].x / uCharMax; + hInputData[index].y = hData[index].y / uCharMax; + hInputData[index].z = hData[index].z / uCharMax; + hInputData[index].w = hData[index].w / uCharMax; + } + } + HIP_CHECK(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; + //printf("(%d, %d): hInputData=(%f, %f, %f, %f), hOutputData=(%f, %f, %f, %f)\n", i, j, + // hInputData[index].x, hInputData[index].y, hInputData[index].z, hInputData[index].w, + // hOutputData[index].x, hOutputData[index].y, hOutputData[index].z, hOutputData[index].w); + + float4 cpuExpected = getExpectedValue(width, height, + offsetX + j, offsetY + i, hInputData); + float4 gpuOutput = hOutputData[index]; + if (sRGB) { + // CTS will map to sRGP before comparison, so we do so + cpuExpected = hipSRGBMap(cpuExpected); + gpuOutput = hipSRGBMap(gpuOutput); + } + // Convert from [0, 1] back to [0, 255] + gpuOutput *= uCharMax; + cpuExpected *= uCharMax; + if (!hipTextureSamplingVerify(gpuOutput, cpuExpected)) { + WARN("Mismatch at (" << offsetX + j << ", " << offsetY + i << ") GPU output : " << + gpuOutput.x << ", " << + gpuOutput.y << ", " << + gpuOutput.z << ", " << + gpuOutput.w << ", " << + " CPU expected: " << + cpuExpected.x << ", " << + cpuExpected.y << ", " << + cpuExpected.z << ", " << + cpuExpected.w << "\n"); + result = false; + goto line1; + } + } + } + +line1: + HIP_CHECK(hipDestroyTextureObject(textureObject)); + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(hipArray)); + free(hData); + free(hOutputData); + free(hInputData); + REQUIRE(result); +} + +TEST_CASE("Unit_hipTextureObj2DCheckRGBAModes") { + CHECK_IMAGE_SUPPORT + + SECTION("RGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest(256, 256, -3.9, 6.1); + runTest(256, 256, 4.4, -7.0); + } + + SECTION("RGBA 2D hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest(256, 256, -8.5, 2.9); + runTest(256, 256, 12.5, 6.7); + } + + SECTION("RGBA 2D hipAddressModeClamp, hipFilterModeLinear, regularCoords") { + runTest(256, 256, -0.4, -0.4); + runTest(256, 256, 4, 14.6); + } + + SECTION("RGBA 2D hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest(256, 256, -0.4, 0.4); + runTest(256, 256, 12.5, 23.7); + } + + SECTION("RGBA 2D hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -3, 8.9); + runTest(256, 256, 4, -0.1); + } + + SECTION("RGBA 2D hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -8.5, 15.9); + runTest(256, 256, 12.5, -17.9); + } + + SECTION("RGBA 2D hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -3, 5.8); + runTest(256, 256, 4, 9.1); + } + + SECTION("RGBA 2D hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -8.5, 6.6); + runTest(256, 256, 12.5, 0.01); + } +} + + +TEST_CASE("Unit_hipTextureObj2DCheckSRGBAModes") { + CHECK_IMAGE_SUPPORT + + SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") { + runTest(256, 256, -3.9, 6.1); + runTest(256, 256, 4.4, -7.0); + } + + SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModePoint, regularCoords") { + runTest(256, 256, -8.5, 2.9); + runTest(256, 256, 12.5, 6.7); + } + + SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModeLinear, regularCoords") { + runTest(256, 256, -0.4, -0.4); + runTest(256, 256, 4, 14.6); + } + + SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModeLinear, regularCoords") { + runTest(256, 256, -0.4, 0.4); + runTest(256, 256, 12.5, 23.7); + } + + SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -3, 8.9); + runTest(256, 256, 4, -0.1); + } + + SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModePoint, normalizedCoords") { + runTest(256, 256, -8.5, 15.9); + runTest(256, 256, 12.5, -17.9); + } + + SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -3, 5.8); + runTest(256, 256, 4, 9.1); + } + + SECTION("SRGBA 2D hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") { + runTest(256, 256, -8.5, 6.6); + runTest(256, 256, 12.5, 0.01); + } +} diff --git a/catch/unit/texture/hipTextureObj3DCheckModes.cc b/catch/unit/texture/hipTextureObj3DCheckModes.cc index eb54713763..d219763a91 100644 --- a/catch/unit/texture/hipTextureObj3DCheckModes.cc +++ b/catch/unit/texture/hipTextureObj3DCheckModes.cc @@ -39,7 +39,7 @@ __global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject, } template -void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY, +static 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, @@ -124,7 +124,7 @@ void runTest(const int width, const int height, const int depth, const float off 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( + float expectedValue = getExpectedValue( width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData); if (!hipTextureSamplingVerify(hOutputData[index], expectedValue)) {