diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index bd4d765863..3eccab939b 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -208,8 +208,6 @@ "Unit_hipMemcpyDtoDAsync_Negative_Parameters", "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===", "Unit_hipGetChannelDesc_Negative_Parameters", - "Unit_hipTextureMipmapRef2D_Positive_Check", - "Unit_hipTextureMipmapRef2D_Negative_Parameters", "=== SWDEV-430116:Below tests failed in stress test on 27/10/23 ===", "Unit_hipFreeAsync_negative", "Unit_hipLaunchHostFunc_multistreams", diff --git a/catch/unit/texture/hipBindTextureToMipmappedArray.cc b/catch/unit/texture/hipBindTextureToMipmappedArray.cc index 40f5370c38..3004cdd835 100644 --- a/catch/unit/texture/hipBindTextureToMipmappedArray.cc +++ b/catch/unit/texture/hipBindTextureToMipmappedArray.cc @@ -21,6 +21,7 @@ THE SOFTWARE. #include #include +#if (!HT_NVIDIA) || (CUDA_VERSION < CUDA_12000) /** * @addtogroup hipBindTextureToMipmappedArray hipBindTextureToMipmappedArray * @{ @@ -29,15 +30,16 @@ THE SOFTWARE. * hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc)` - * Binds a mipmapped array to a texture. */ -#if CUDA_VERSION < CUDA_12000 texture texRef; // MipMap is currently supported only on windows #if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) -__global__ void tex2DKernel(float* outputData, int width, float level) { +__global__ void tex2DKernel(float* outputData, int width, int height, float level) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; - outputData[y * width + x] = tex2DLod(texRef, x, y, level); + float u = x / (float)width; + float v = y / (float)height; + outputData[y * width + x] = tex2DLod(texRef, u, v, level); } static void runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) { @@ -77,7 +79,7 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int texRef.addressMode[0] = hipAddressModeWrap; texRef.addressMode[1] = hipAddressModeWrap; texRef.filterMode = hipFilterModePoint; - texRef.normalized = 0; + texRef.normalized = 1; // Bind the array to the texture HIP_CHECK(hipBindTextureToMipmappedArray(&texRef, mip_array_ptr, &channelDesc)); @@ -90,7 +92,8 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int dim3 dimBlock(16, 16, 1); dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); - hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, mipmap_level); + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, + width, height, mipmap_level); HIP_CHECK(hipGetLastError()); HIP_CHECK(hipDeviceSynchronize()); @@ -136,7 +139,7 @@ TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") { // Height Width Vector std::vector hw_vec = {2048, 1024, 512, 256, 64}; std::vector mip_vec = {8, 4, 2, 1}; -#ifdef _WIN32 +#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) for (auto& hw : hw_vec) { for (auto& mip : mip_vec) { if ((hw / static_cast(pow(2, (mip * 2)))) > 0) { @@ -145,7 +148,8 @@ TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") { } } #else - SUCCEED("Mipmaps are Supported only on windows, skipping the test."); + SUCCEED("Mipmaps are Supported only on windows on devices with image support," + " skipping the test."); #endif } @@ -171,7 +175,7 @@ TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") { TEST_CASE("Unit_hipTextureMipmapRef2D_Negative_Parameters") { CHECK_IMAGE_SUPPORT -#ifdef _WIN32 +#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) unsigned int width = 64; unsigned int height = 64; unsigned int mipmap_level = 1; @@ -205,8 +209,8 @@ TEST_CASE("Unit_hipTextureMipmapRef2D_Negative_Parameters") { HIP_CHECK(hipFreeMipmappedArray(mip_array_ptr)); #else - SUCCEED("Mipmaps are Supported only on windows, skipping the test."); + SUCCEED("Mipmaps are Supported only on windows on devices with image support," + " skipping the test."); #endif } - -#endif +#endif \ No newline at end of file