diff --git a/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh index 606435a088..685072277e 100644 --- a/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh +++ b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh @@ -33,7 +33,7 @@ constexpr size_t ChannelToRead = 1; template __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height, bool textureGather) { - #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT // Calculate normalized texture coordinates const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -50,7 +50,7 @@ __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t wid output[y * width + x] = tex2D(texObj, u, v); } } - #endif +#endif } template void checkDataIsAscending(const std::vector& hostData) { diff --git a/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc b/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc index 29686b2ab6..7beae3d2bb 100644 --- a/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc +++ b/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc @@ -75,6 +75,11 @@ static void ArrayCreate_DiffSizes(int gpu) { TEST_CASE("Unit_hipArrayCreate_DiffSizes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + ArrayCreate_DiffSizes(0); HIP_CHECK_THREAD_FINALIZE(); } @@ -87,6 +92,11 @@ and verifies the hipArrayCreate API with small and big chunks data TEST_CASE("Unit_hipArrayCreate_MultiThread") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + std::vector threadlist; int devCnt = 0; @@ -188,7 +198,7 @@ void testArrayAsTexture(hipArray_t array, const size_t width, const size_t heigh std::fill(std::begin(hostData), std::end(hostData), 0); HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost)); -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT checkDataIsAscending(hostData); #endif @@ -203,6 +213,11 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_happy", "", uint, int, int4, ushort, sho char4, float, float2, float4) { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + using vec_info = vector_info; DriverContext ctx; @@ -228,6 +243,11 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort uchar2, char4, float, float2, float4) { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + using vec_info = vector_info; DriverContext ctx; @@ -288,6 +308,11 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; desc.Format = driverFormats[0]; @@ -304,6 +329,11 @@ TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { TEST_CASE("Unit_hipArrayCreate_Nullptr") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + DriverContext ctx; SECTION("Null array") { HIP_ARRAY_DESCRIPTOR desc; @@ -324,6 +354,11 @@ TEST_CASE("Unit_hipArrayCreate_Nullptr") { TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; desc.Format = GENERATE(from_range(std::begin(driverFormats), std::end(driverFormats))); @@ -342,6 +377,11 @@ TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") { TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; diff --git a/projects/hip-tests/catch/unit/surface/surf1D.cc b/projects/hip-tests/catch/unit/surface/surf1D.cc index 62dbf71351..58764e5887 100644 --- a/projects/hip-tests/catch/unit/surface/surf1D.cc +++ b/projects/hip-tests/catch/unit/surface/surf1D.cc @@ -35,7 +35,7 @@ THE SOFTWARE. template __global__ void surf1DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, int width) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; if (x < width) { surf1Dread(outputData + x, surfaceObject, x * sizeof(T)); @@ -45,7 +45,7 @@ __global__ void surf1DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, i template __global__ void surf1DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, int width) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; if (x < width) { surf1Dwrite(inputData[x], surfaceObject, x * sizeof(T)); @@ -56,7 +56,7 @@ __global__ void surf1DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, in template __global__ void surf1DKernelRW(hipSurfaceObject_t surfaceObject, hipSurfaceObject_t outputSurfObj, int width) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; if (x < width) { T data; @@ -247,6 +247,11 @@ TEMPLATE_TEST_CASE("Unit_surf1Dread_Positive_Basic", "", char, uchar, short, ush uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67, 131, 263); runTestR(width); } @@ -268,6 +273,11 @@ TEMPLATE_TEST_CASE("Unit_surf1Dwrite_Positive_Basic", "", char, uchar, short, us uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67, 131, 263); runTestW(width); } @@ -289,6 +299,11 @@ TEMPLATE_TEST_CASE("Unit_surf1D_Positive_ReadWrite", "", char, uchar, short, ush uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67, 131, 263); runTestRW(width); } diff --git a/projects/hip-tests/catch/unit/surface/surf2D.cc b/projects/hip-tests/catch/unit/surface/surf2D.cc index 477951f816..9814a1e840 100644 --- a/projects/hip-tests/catch/unit/surface/surf2D.cc +++ b/projects/hip-tests/catch/unit/surface/surf2D.cc @@ -38,7 +38,7 @@ THE SOFTWARE. template __global__ void surf2DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, int width, int height) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { @@ -50,7 +50,7 @@ __global__ void surf2DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, i template __global__ void surf2DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, int width, int height) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { @@ -62,7 +62,7 @@ __global__ void surf2DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, in template __global__ void surf2DKernelRW(hipSurfaceObject_t surfaceObject, hipSurfaceObject_t outputSurfObj, int width, int height) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { @@ -288,6 +288,11 @@ TEMPLATE_TEST_CASE("Unit_surf2Dread_Positive_Basic", "", char, uchar, short, ush uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67); const int height = GENERATE(131, 263); runTestR(width, height); @@ -310,6 +315,11 @@ TEMPLATE_TEST_CASE("Unit_surf2Dwrite_Positive_Basic", "", char, uchar, short, us uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67); const int height = GENERATE(131, 263); runTestW(width, height); @@ -332,6 +342,11 @@ TEMPLATE_TEST_CASE("Unit_surf2D_Positive_ReadWrite", "", char, uchar, short, ush uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67); const int height = GENERATE(131, 263); runTestRW(width, height); diff --git a/projects/hip-tests/catch/unit/surface/surf3D.cc b/projects/hip-tests/catch/unit/surface/surf3D.cc index 7728335082..2a4f8b2a75 100644 --- a/projects/hip-tests/catch/unit/surface/surf3D.cc +++ b/projects/hip-tests/catch/unit/surface/surf3D.cc @@ -36,7 +36,7 @@ THE SOFTWARE. template __global__ void surf3DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, int width, int height, int depth) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -49,7 +49,7 @@ __global__ void surf3DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, i template __global__ void surf3DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, int width, int height, int depth) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -62,7 +62,7 @@ __global__ void surf3DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, in template __global__ void surf3DKernelRW(hipSurfaceObject_t surfaceObject, hipSurfaceObject_t outputSurfObj, int width, int height, int depth) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -334,6 +334,11 @@ TEMPLATE_TEST_CASE("Unit_surf3Dread_Positive_Basic", "", char, uchar, short, ush uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67); const int height = GENERATE(131, 263); const int depth = GENERATE(4, 11); @@ -357,6 +362,11 @@ TEMPLATE_TEST_CASE("Unit_surf3Dwrite_Positive_Basic", "", char, uchar, short, us uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67); const int height = GENERATE(131, 263); const int depth = GENERATE(4, 11); @@ -380,6 +390,11 @@ TEMPLATE_TEST_CASE("Unit_surf3D_Positive_ReadWrite", "", char, uchar, short, ush uint4, float4) { CHECK_IMAGE_SUPPORT; +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + const int width = GENERATE(31, 67); const int height = GENERATE(131, 263); const int depth = GENERATE(4, 11); diff --git a/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc b/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc index 3d01d2429d..79aa118c73 100644 --- a/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc +++ b/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc @@ -32,7 +32,7 @@ texture tex; // texture object is a kernel argument static __global__ void texture2dCopyKernel(TYPE_t* dst) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = threadIdx.x + blockIdx.x * blockDim.x; int y = threadIdx.y + blockIdx.y * blockDim.y; if ( (x < SIZE_W) && (y < SIZE_H) ) { @@ -45,6 +45,11 @@ static __global__ void texture2dCopyKernel(TYPE_t* dst) { TEST_CASE("Unit_hipBindTexture2D_Pitch") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + TYPE_t* B; TYPE_t* A; TYPE_t* devPtrB; diff --git a/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc b/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc index 9de21ecbf7..47c3aa8978 100644 --- a/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc +++ b/projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc @@ -27,7 +27,7 @@ THE SOFTWARE. texture tex; static __global__ void kernel(float *out) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; if (x < N) { out[x] = tex1Dfetch(tex, x); @@ -38,6 +38,11 @@ static __global__ void kernel(float *out) { TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + float *texBuf; float val[N], output[N]; size_t offset = 0; diff --git a/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc b/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc index 767089cbb4..25680da7bf 100644 --- a/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc +++ b/projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc @@ -33,7 +33,7 @@ THE SOFTWARE. texture texRef; // MipMap is currently supported only on windows -#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) +#if (defined(_WIN32) && !__HIP_NO_IMAGE_SUPPORT) __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; @@ -136,10 +136,16 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int */ TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + // Height Width Vector std::vector hw_vec = {2048, 1024, 512, 256, 64}; std::vector mip_vec = {8, 4, 2, 1}; -#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) +#if (defined(_WIN32) && !__HIP_NO_IMAGE_SUPPORT) for (auto& hw : hw_vec) { for (auto& mip : mip_vec) { if ((hw / static_cast(pow(2, (mip * 2)))) > 0) { @@ -175,7 +181,12 @@ TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") { TEST_CASE("Unit_hipTextureMipmapRef2D_Negative_Parameters") { CHECK_IMAGE_SUPPORT -#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT)) +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + +#if defined(_WIN32) unsigned int width = 64; unsigned int height = 64; unsigned int mipmap_level = 1; diff --git a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc index a8134e1ae2..056f59f3ef 100644 --- a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc +++ b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -46,7 +46,7 @@ texture texuc; template __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT unsigned int elementID = threadIdx.x; if (elementID >= numElements) return; @@ -149,6 +149,11 @@ static void runTest_hipTextureFilterMode() { TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + #if HT_AMD hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, 0)); diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture1DLayered.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture1DLayered.cc index 7b3de8d235..850779026f 100644 --- a/projects/hip-tests/catch/unit/texture/hipSimpleTexture1DLayered.cc +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture1DLayered.cc @@ -27,7 +27,7 @@ THE SOFTWARE. template __global__ void simpleKernelLayered1DArray(hipTextureObject_t tex, TestType* outputData, unsigned int width, unsigned int layer) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; outputData[layer * width + x] = tex1DLayered(tex, x, layer); #endif @@ -62,6 +62,12 @@ TEMPLATE_TEST_CASE("Unit_Layered1DTexture_Check_HostBufferToFromLayered1DArray", char2, uchar2, short2, ushort2, int2, uint2, float2, char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + constexpr int SIZE = 512; constexpr int num_layers = 5; constexpr unsigned int width = SIZE; @@ -205,6 +211,12 @@ TEMPLATE_TEST_CASE("Unit_Layered1DTexture_Check_DeviceBufferToFromLayered1DArray char2, uchar2, short2, ushort2, int2, uint2, float2, char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + constexpr int SIZE = 512; constexpr int num_layers = 5; constexpr unsigned int width = SIZE; diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc index 5075a73d9a..4e2cb22ef3 100644 --- a/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc @@ -28,7 +28,7 @@ template __global__ void simpleKernelLayered2DArray(hipTextureObject_t tex, TestType* outputData, unsigned int width, unsigned int height, unsigned int layer) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; outputData[layer * width * height + y * width + x] = tex2DLayered(tex, x, y, layer); @@ -64,6 +64,12 @@ TEMPLATE_TEST_CASE("Unit_Layered2DTexture_Check_HostBufferToFromLayered2DArray", char2, uchar2, short2, ushort2, int2, uint2, float2, char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + constexpr int SIZE = 512; constexpr int num_layers = 5; constexpr unsigned int width = SIZE; @@ -210,6 +216,12 @@ TEMPLATE_TEST_CASE("Unit_Layered2DTexture_Check_DeviceBufferToFromLayered2DArray char2, uchar2, short2, ushort2, int2, uint2, float2, char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + constexpr int SIZE = 512; constexpr int num_layers = 5; constexpr unsigned int width = SIZE; diff --git a/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc b/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc index 4c397aec9f..486e8f0d5b 100644 --- a/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc +++ b/projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc @@ -31,7 +31,7 @@ texture texc; template __global__ void simpleKernel3DArray(T* outputData, int width, int height, int depth) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT for (int i = 0; i < depth; i++) { for (int j = 0; j < height; j++) { for (int k = 0; k < width; k++) { @@ -116,6 +116,11 @@ static void runSimpleTexture3D_Check(int width, int height, int depth, TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + for ( int i = 1; i < 25; i++ ) { runSimpleTexture3D_Check(i, i, i, &texf); runSimpleTexture3D_Check(i+1, i, i, &texi); diff --git a/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc index 0954e770f5..216ba061db 100644 --- a/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc @@ -24,7 +24,7 @@ THE SOFTWARE. #define offset 3 static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < (N - offset)) val[k] = tex1Dfetch(obj, k+offset); @@ -110,6 +110,11 @@ static void runTest(hipTextureAddressMode addressMode, TEST_CASE("Unit_tex1Dfetch_CheckModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("hipAddressModeClamp AND hipFilterModePoint") { runTest(hipAddressModeClamp, hipFilterModePoint); } diff --git a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc index 1cc0d8efc1..7af30821a5 100644 --- a/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc +++ b/projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc @@ -37,7 +37,7 @@ THE SOFTWARE. template static __global__ void texture2dCopyKernel(hipTextureObject_t texObj, TYPE_t* dst) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT for (int i = 0; i < SIZE_H; i++) for (int j = 0; j < SIZE_W; j++) dst[SIZE_W*i+j] = tex2D(texObj, j, i); @@ -61,6 +61,11 @@ TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int, unsigned char, int16_t, char, unsigned int) { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + TestType* B; TestType* A; TestType* devPtrB; diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc index b27176621b..5e79c3c273 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc @@ -33,7 +33,7 @@ static constexpr bool printLog = false; // Print log for debugging template static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, hipTextureObject_t texIn, unsigned int width, T* data = nullptr) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; float px = 1.0 / float(width); @@ -61,7 +61,7 @@ static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, template __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width, float offsetX, float lod, T* data = nullptr) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; float px = 1.0 / float(width); @@ -316,6 +316,11 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType", char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, " "hipAddressModeClamp 23") { @@ -366,6 +371,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloa char2, uchar2, short2, ushort2, char4, uchar4, short4, ushort4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " "hipAddressModeClamp 23") { @@ -438,6 +449,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloa TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType float only", "", float, float1, float2, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, " "hipAddressModeClamp 23, 0.") { diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc index 420ba2d9aa..4393a19802 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc @@ -40,7 +40,7 @@ static constexpr bool printLog = false; // Print log for debugging template static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, hipTextureObject_t texIn, unsigned int width, unsigned int height, T* data = nullptr) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -71,7 +71,7 @@ static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, template static __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width, unsigned int height, float offsetX, float offsetY, float lod, T* data = nullptr) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -349,6 +349,11 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType", char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, " "hipAddressModeClamp 23, 21") { @@ -399,6 +404,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloa char2, uchar2, short2, ushort2, char4, uchar4, short4, ushort4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " "hipAddressModeClamp 23, 21") { @@ -471,6 +482,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloa TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType float only", "", float, float1, float2, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, " "hipAddressModeClamp 23, 17, 0., 0.") { diff --git a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc index fede084e0d..872b1b5335 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc @@ -33,7 +33,7 @@ static constexpr bool printLog = false; // Print log for debugging template static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, hipTextureObject_t texIn, unsigned int width, unsigned int height, unsigned int depth, T* data = nullptr) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -74,7 +74,7 @@ static __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int widt unsigned int height, unsigned int depth, float offsetX, float offsetY, float offsetZ, float lod, T* data = nullptr) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -367,6 +367,11 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType", char4, uchar4, short4, ushort4, int4, uint4, float4) { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, " "hipAddressModeClamp 23, 21, 47") { @@ -417,6 +422,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloa char2, uchar2, short2, ushort2, char4, uchar4, short4, ushort4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, " "hipAddressModeClamp 23, 21, 67") { @@ -489,6 +500,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloa TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType float only", "", float, float1, float2, float4) { CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION( "Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, " "hipAddressModeClamp 23, 17, 301, 0., 0., 0.") { diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc index 425ce46415..bdbf22728e 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc @@ -31,7 +31,7 @@ THE SOFTWARE. template __global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject, int width, float offsetX) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; outputData[x] = tex1D(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX); #endif @@ -122,6 +122,11 @@ static void runTest(const int width, const float offsetX) { TEST_CASE("Unit_hipTextureObj1DCheckModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { runTest(256, -3); runTest(256, 4); diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckSRGBModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckSRGBModes.cc index 9b5ffa560f..e821c5b380 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckSRGBModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckSRGBModes.cc @@ -25,7 +25,7 @@ THE SOFTWARE. template __global__ void tex1DRGBAKernel(float4 *outputData, hipTextureObject_t textureObject, int width, float offsetX) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; outputData[x] = tex1D(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX); @@ -33,7 +33,7 @@ __global__ void tex1DRGBAKernel(float4 *outputData, hipTextureObject_t textureOb } __global__ void tex1DRGBAKernelFetch(float4 *outputData, hipTextureObject_t textureObject, float offsetX) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; outputData[x] = tex1Dfetch(textureObject, int(x + offsetX)); #endif @@ -177,6 +177,11 @@ line1: TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - array") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("RGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, regularCoords") { runTest(255, -3.9); runTest(255, 4.4); @@ -228,6 +233,11 @@ TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - array") { TEST_CASE("Unit_hipTextureObj1DCheckSRGBAModes - array") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("SRGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, regularCoords") { runTest(255, -3.9); runTest(255, 4.4); @@ -277,6 +287,11 @@ TEST_CASE("Unit_hipTextureObj1DCheckSRGBAModes - array") { TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - buffer") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("RGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeLinear, regularCoords") { runTest(255); } @@ -289,6 +304,11 @@ TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - buffer") { TEST_CASE("Unit_hipTextureObj1DCheckSRGBAModes - buffer") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("SRGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeLinear, regularCoords") { runTest(255); } diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc index c267bec29d..eb3bcb5373 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc @@ -23,7 +23,7 @@ THE SOFTWARE. #define N 512 static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < N) { val[k] = tex1Dfetch(obj, k); @@ -35,6 +35,11 @@ static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + // Allocating the required buffer on gpu device float *texBuf, *texBufOut; float val[N], output[N]; diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc index 4346ef3094..4459724dcd 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc @@ -28,7 +28,7 @@ THE SOFTWARE. __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__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, x, y); @@ -50,6 +50,11 @@ __global__ void tex2DKernel(float* outputData, TEST_CASE("Unit_hipTextureObj2D_Check") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + constexpr int SIZE = 256; constexpr unsigned int width = SIZE; constexpr unsigned int height = SIZE; diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc index 8a26f975fa..1041cf1bf4 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc @@ -32,7 +32,7 @@ template __global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject, int width, int height, float offsetX, float offsetY) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__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, @@ -134,6 +134,11 @@ line1: TEST_CASE("Unit_hipTextureObj2DCheckModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") { runTest(256, 256, -3.9, 6.1); runTest(256, 256, 4.4, -7.0); diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc index 5c16490177..47b18e37b3 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc @@ -26,7 +26,7 @@ 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 +#if !__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, @@ -157,6 +157,11 @@ line1: TEST_CASE("Unit_hipTextureObj2DCheckRGBAModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("RGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") { runTest(256, 256, -3.9, 6.1); runTest(256, 256, 4.4, -7.0); @@ -202,6 +207,11 @@ TEST_CASE("Unit_hipTextureObj2DCheckRGBAModes") { TEST_CASE("Unit_hipTextureObj2DCheckSRGBAModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") { runTest(256, 256, -3.9, 6.1); runTest(256, 256, 4.4, -7.0); diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc b/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc index f4c7d7927b..208e37ef00 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc @@ -34,7 +34,7 @@ bool LinearFilter3D = false; template __global__ void tex3DKernel(float* outputData, hipTextureObject_t textureObject, int width, int height, int depth, float offsetX, float offsetY, float offsetZ) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -173,6 +173,11 @@ line1: TEST_CASE("Unit_hipTextureObj3DCheckModes") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + int device = 0; hipDeviceProp_t props; HIPCHECK(hipGetDeviceProperties(&props, device)); diff --git a/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc b/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc index b12f0e7da0..c00fa0daaf 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc @@ -27,7 +27,7 @@ THE SOFTWARE. template __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < N) { val[k] = tex1Dfetch(obj, k); @@ -143,6 +143,11 @@ bool runTest() { TEST_CASE("Unit_hipTextureFetch_vector") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + // test for char runTest(); runTest(); diff --git a/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc b/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc index 434a86a71f..d5cf387246 100644 --- a/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc +++ b/projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc @@ -26,7 +26,7 @@ THE SOFTWARE. texture tex; __global__ void tex2DKernel(float* outputData, int width) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +#if !__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(tex, x, y); @@ -36,6 +36,11 @@ __global__ void tex2DKernel(float* outputData, int width) { TEST_CASE("Unit_hipTextureRef2D_Check") { CHECK_IMAGE_SUPPORT +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + constexpr int SIZE = 256; constexpr unsigned int width = SIZE; constexpr unsigned int height = SIZE;