From 643f8dbd36e716384ec2a1d9b43c24413155e2a3 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 9 May 2022 21:16:20 +0530 Subject: [PATCH] SWDEV-305992 - Improve texture tests (#2646) Enable more texture supports for hip-rocclr. Skip texture tests on device that doesn't support images. SWDEV-322257 - Fix issues of image tests skipped on NVidia devices. Change-Id: Ia99d06b1e97fc945f1b740e47710f4dcd70f38ca --- .../11_texture_driver/tex2dKernel.cpp | 2 ++ .../11_texture_driver/texture2dDrv.cpp | 13 +++++++++++++ tests/catch/include/hip_test_common.hh | 18 ++++++++++++++++++ .../hipCreateTextureObject_ArgValidation.cc | 2 ++ .../texture/hipCreateTextureObject_Array.cc | 4 ++++ .../texture/hipCreateTextureObject_Linear.cc | 2 ++ .../texture/hipCreateTextureObject_Pitch2D.cc | 2 ++ .../unit/texture/hipTextureObjFetchVector.cc | 4 ++++ .../module/hipModuleTexture2dDrv.cpp | 9 ++------- tests/src/test_common.h | 15 +++++++++++++++ tests/src/texture/hipBindTex2DPitch.cpp | 14 +++++++++----- tests/src/texture/hipBindTexRef1DFetch.cpp | 5 ++++- tests/src/texture/hipGetChanDesc.cpp | 11 +++-------- .../src/texture/hipNormalizedFloatValueTex.cpp | 12 ++++-------- tests/src/texture/hipTex1DFetchCheckModes.cpp | 4 ++++ tests/src/texture/hipTexObjPitch.cpp | 9 ++------- tests/src/texture/hipTextureMipmapObj2D.cpp | 9 ++------- .../src/texture/hipTextureObj1DCheckModes.cpp | 4 ++++ tests/src/texture/hipTextureObj1DFetch.cpp | 7 ++++++- tests/src/texture/hipTextureObj2D.cpp | 13 ++++--------- .../src/texture/hipTextureObj2DCheckModes.cpp | 4 ++++ .../src/texture/hipTextureObj3DCheckModes.cpp | 4 ++++ tests/src/texture/hipTextureRef2D.cpp | 10 +++++++--- tests/src/texture/simpleTexture2DLayered.cpp | 10 +++++++--- tests/src/texture/simpleTexture3D.cpp | 13 ++++--------- 25 files changed, 132 insertions(+), 68 deletions(-) diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 613db7ff01..5f2ded8518 100644 --- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -25,7 +25,9 @@ THE SOFTWARE. texture tex; extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; outputData[y * width + x] = tex2D(tex, x, y); +#endif } diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 1e27b14008..161b3ea93d 100644 --- a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -127,7 +127,20 @@ bool runTest(int argc, char** argv) { return testResult; } +inline bool isImageSupported() { + int imageSupport = 1; +#ifdef __HIP_PLATFORM_AMD__ + HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, + 0)); +#endif + return imageSupport != 0; +} + int main(int argc, char** argv) { + if (!isImageSupported()) { + printf("Texture is not support on the device. Skipped.\n"); + return 0; + } hipInit(0); testResult = runTest(argc, argv); printf("%s ...\n", testResult ? "PASSED" : "FAILED"); diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index 68817ef1c2..e767f07387 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -115,4 +115,22 @@ static inline int RAND_R(unsigned* rand_seed) return rand_r(rand_seed); #endif } + +inline bool isImageSupported() { + int imageSupport = 1; +#ifdef __HIP_PLATFORM_AMD__ + int device; + HIP_CHECK(hipGetDevice(&device)); + HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, + device)); +#endif + return imageSupport != 0; } + +} + +// This must be called in the beginning of image test app's main() to indicate whether image +// is supported. +#define checkImageSupport() \ + if (!HipTest::isImageSupported()) \ + { printf("Texture is not support on the device. Skipped.\n"); return; } diff --git a/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc b/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc index 91a0827127..c7570aea03 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc @@ -25,6 +25,8 @@ THE SOFTWARE. * Validate argument list of texture object api. */ TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { + checkImageSupport(); + float *texBuf; hipError_t ret; constexpr int xsize = 32; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_Array.cc b/tests/catch/unit/texture/hipCreateTextureObject_Array.cc index 037efea1af..da42c7878e 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_Array.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_Array.cc @@ -23,6 +23,8 @@ THE SOFTWARE. * Validates Array Resource texture object with negative/functional tests. */ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { + checkImageSupport(); + hipError_t ret; hipResourceDesc resDesc; hipTextureDesc texDesc; @@ -47,6 +49,8 @@ TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { * with negative/functional tests. */ TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") { + checkImageSupport(); + hipError_t ret; hipResourceDesc resDesc; hipTextureDesc texDesc; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc b/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc index 29bd35a821..0e2caf6a80 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc @@ -26,6 +26,8 @@ THE SOFTWARE. * Validates Linear Resource texture object with negative/functional tests. */ TEST_CASE("Unit_hipCreateTextureObject_LinearResource") { + checkImageSupport(); + float *texBuf; hipError_t ret; constexpr int xsize = 32; diff --git a/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc b/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc index 89eb936c3c..7fc593708c 100644 --- a/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc +++ b/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc @@ -29,6 +29,8 @@ THE SOFTWARE. * Validates Pitch2D Resource texture object with negative and functional tests */ TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { + checkImageSupport(); + hipError_t ret; hipResourceDesc resDesc; hipTextureDesc texDesc; diff --git a/tests/catch/unit/texture/hipTextureObjFetchVector.cc b/tests/catch/unit/texture/hipTextureObjFetchVector.cc index 8fc6c6730a..d3faa5d5b0 100644 --- a/tests/catch/unit/texture/hipTextureObjFetchVector.cc +++ b/tests/catch/unit/texture/hipTextureObjFetchVector.cc @@ -23,10 +23,12 @@ THE SOFTWARE. template __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < N) { val[k] = tex1Dfetch(obj, k); } +#endif } template @@ -203,6 +205,8 @@ bool runTest(const char *description) { } TEST_CASE("Unit_hipTextureFetch_vector") { + checkImageSupport(); + // test for char runTest("char1"); runTest("char2"); diff --git a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp index b4abf307dd..67c62f412e 100644 --- a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp +++ b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp @@ -592,13 +592,8 @@ bool testTexSingleStreamMultGPU(unsigned int numOfGPUs, int main(int argc, char** argv) { HipTest::parseStandardArguments(argc, argv, true); - int imageSupport = 0; - hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - p_gpuDevice); - if (!imageSupport) { - printf("Texture is not support on the device\n"); - passed(); - } + checkImageSupport(); + bool TestPassed = true; if (p_tests == 0x01) { TestPassed = testTexType(HIP_AD_FORMAT_FLOAT, diff --git a/tests/src/test_common.h b/tests/src/test_common.h index b6dcba4538..21e82a71ca 100644 --- a/tests/src/test_common.h +++ b/tests/src/test_common.h @@ -564,5 +564,20 @@ struct MemTraits { } }; +inline bool isImageSupported() { + int imageSupport = 1; +#ifdef __HIP_PLATFORM_AMD__ + HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, + p_gpuDevice)); +#endif + return imageSupport != 0; +} + }; // namespace HipTest + +// This must be called in the beginning of image test app's main() to indicate whether image +// is supported. +#define checkImageSupport() \ + if (!HipTest::isImageSupported()) \ + { printf("Texture is not support on the device. Skipped.\n"); passed(); } #endif //__cplusplus diff --git a/tests/src/texture/hipBindTex2DPitch.cpp b/tests/src/texture/hipBindTex2DPitch.cpp index 9de2ec52ae..41c99a8a61 100644 --- a/tests/src/texture/hipBindTex2DPitch.cpp +++ b/tests/src/texture/hipBindTex2DPitch.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 */ @@ -32,16 +32,18 @@ texture tex; // texture object is a kernel argument __global__ void texture2dCopyKernel( TYPE_t* dst) { - +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; if ( (x< SIZE_W) && (y< SIZE_H) ){ dst[SIZE_W*y+x] = tex2D(tex, x, y); } +#endif } int main (void) { + checkImageSupport(); TYPE_t* B; TYPE_t* A; TYPE_t* devPtrB; @@ -49,8 +51,8 @@ int main (void) B = new TYPE_t[SIZE_H*SIZE_W]; A = new TYPE_t[SIZE_H*SIZE_W]; - for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){ - A[i-1] = i; + for (size_t i = 0; i < (SIZE_H * SIZE_W); i++) { + A[i] = i + 1; } size_t devPitchA, tex_ofs; @@ -58,12 +60,14 @@ int main (void) HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); + tex.addressMode[0] = hipAddressModeClamp; + tex.addressMode[1] = hipAddressModeClamp; tex.normalized = false; HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, SIZE_W, SIZE_H, devPitchA)); HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; - hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB); + hipLaunchKernelGGL(texture2dCopyKernel, dim3(3, 2, 1), dim3(4, 4, 1), 0, 0, devPtrB); hipDeviceSynchronize(); HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t), SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost)); diff --git a/tests/src/texture/hipBindTexRef1DFetch.cpp b/tests/src/texture/hipBindTexRef1DFetch.cpp index 2903bf46a4..29510ac161 100644 --- a/tests/src/texture/hipBindTexRef1DFetch.cpp +++ b/tests/src/texture/hipBindTexRef1DFetch.cpp @@ -35,15 +35,19 @@ THE SOFTWARE. texture tex; __global__ void kernel(float *out) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; if(x *tex) 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)); + tex->addressMode[0] = hipAddressModeClamp; tex->normalized = true; tex->channelDesc = desc; tex->filterMode = fMode; @@ -172,15 +173,10 @@ bool runTest() { int main(int argc, char** argv) { - int imageSupport = 0; - hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - p_gpuDevice); - if (!imageSupport) { - printf("Texture is not support on the device\n"); - passed(); - } HipTest::parseStandardArguments(argc, argv, true); - int device = 0; + checkImageSupport(); + + int device = p_gpuDevice; bool status = false; HIPCHECK(hipSetDevice(device)); hipDeviceProp_t props; diff --git a/tests/src/texture/hipTex1DFetchCheckModes.cpp b/tests/src/texture/hipTex1DFetchCheckModes.cpp index 433185c124..a7fbee744a 100644 --- a/tests/src/texture/hipTex1DFetchCheckModes.cpp +++ b/tests/src/texture/hipTex1DFetchCheckModes.cpp @@ -29,14 +29,18 @@ THE SOFTWARE. #define N 16 #define offset 3 __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < N) val[k] = tex1Dfetch(obj, k+offset); +#endif } int runTest(hipTextureAddressMode, hipTextureFilterMode); int main(int argc, char **argv) { + checkImageSupport(); + int testResult = runTest(hipAddressModeClamp,hipFilterModePoint); testResult = testResult & runTest(hipAddressModeClamp,hipFilterModeLinear); testResult = testResult & runTest(hipAddressModeBorder,hipFilterModePoint); diff --git a/tests/src/texture/hipTexObjPitch.cpp b/tests/src/texture/hipTexObjPitch.cpp index 2e59d91d7c..948556bf38 100644 --- a/tests/src/texture/hipTexObjPitch.cpp +++ b/tests/src/texture/hipTexObjPitch.cpp @@ -96,13 +96,8 @@ void texture2Dtest() int main() { - int imageSupport = 0; - hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - p_gpuDevice); - if (!imageSupport) { - printf("Texture is not support on the device\n"); - passed(); - } + checkImageSupport(); + texture2Dtest(); texture2Dtest(); texture2Dtest(); diff --git a/tests/src/texture/hipTextureMipmapObj2D.cpp b/tests/src/texture/hipTextureMipmapObj2D.cpp index 3a59a9640d..d3b174be52 100644 --- a/tests/src/texture/hipTextureMipmapObj2D.cpp +++ b/tests/src/texture/hipTextureMipmapObj2D.cpp @@ -171,14 +171,9 @@ bool runTest(int argc, char** argv) { } int main(int argc, char** argv) { + checkImageSupport(); + bool testResult = true; - int imageSupport = 0; - hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - p_gpuDevice); - if (!imageSupport) { - printf("Texture is not support on the device\n"); - passed(); - } #ifdef _WIN32 testResult = runTest(argc, argv); #else diff --git a/tests/src/texture/hipTextureObj1DCheckModes.cpp b/tests/src/texture/hipTextureObj1DCheckModes.cpp index fbcf14cf62..59c20201e9 100644 --- a/tests/src/texture/hipTextureObj1DCheckModes.cpp +++ b/tests/src/texture/hipTextureObj1DCheckModes.cpp @@ -14,8 +14,10 @@ template __global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject, int width, float offsetX) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int x = blockIdx.x * blockDim.x + threadIdx.x; outputData[x] = tex1D(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX); +#endif } template @@ -88,6 +90,8 @@ bool runTest(const int width, const float offsetX) { } int main(int argc, char **argv) { + checkImageSupport(); + bool testResult = true; testResult = testResult && runTest(256, -3); testResult = testResult && runTest(256, 4); diff --git a/tests/src/texture/hipTextureObj1DFetch.cpp b/tests/src/texture/hipTextureObj1DFetch.cpp index e10781da06..95f05a0a1e 100644 --- a/tests/src/texture/hipTextureObj1DFetch.cpp +++ b/tests/src/texture/hipTextureObj1DFetch.cpp @@ -21,7 +21,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 */ @@ -32,14 +32,18 @@ THE SOFTWARE. #define N 512 __global__ void tex1dKernel(float *val, hipTextureObject_t obj) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT int k = blockIdx.x * blockDim.x + threadIdx.x; if (k < N) val[k] = tex1Dfetch(obj, k); +#endif } int runTest(void); int main(int argc, char **argv) { + checkImageSupport(); + int testResult = runTest(); if(testResult) { passed(); @@ -72,6 +76,7 @@ int runTest() { hipTextureDesc texDesc; memset(&texDesc, 0, sizeof(texDesc)); texDesc.readMode = hipReadModeElementType; + texDesc.addressMode[0]= hipAddressModeClamp; // Creating texture object hipTextureObject_t texObj = 0; diff --git a/tests/src/texture/hipTextureObj2D.cpp b/tests/src/texture/hipTextureObj2D.cpp index 18f40eab11..2641c28e17 100644 --- a/tests/src/texture/hipTextureObj2D.cpp +++ b/tests/src/texture/hipTextureObj2D.cpp @@ -22,13 +22,8 @@ __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int runTest(int argc, char** argv); int main(int argc, char** argv) { - int imageSupport = 0; - hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - p_gpuDevice); - if (!imageSupport) { - printf("Texture is not support on the device\n"); - passed(); - } + checkImageSupport(); + int testResult = runTest(argc, argv); if (testResult) { @@ -70,8 +65,8 @@ int runTest(int argc, char** argv) { // Specify texture object parameters hipTextureDesc texDesc; memset(&texDesc, 0, sizeof(texDesc)); - texDesc.addressMode[0] = hipAddressModeWrap; - texDesc.addressMode[1] = hipAddressModeWrap; + texDesc.addressMode[0] = hipAddressModeClamp; + texDesc.addressMode[1] = hipAddressModeClamp; texDesc.filterMode = hipFilterModePoint; texDesc.readMode = hipReadModeElementType; texDesc.normalizedCoords = 0; diff --git a/tests/src/texture/hipTextureObj2DCheckModes.cpp b/tests/src/texture/hipTextureObj2DCheckModes.cpp index 1c79ff59d3..f8418d3828 100644 --- a/tests/src/texture/hipTextureObj2DCheckModes.cpp +++ b/tests/src/texture/hipTextureObj2DCheckModes.cpp @@ -15,11 +15,13 @@ 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 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 @@ -100,6 +102,8 @@ line1: } int main(int argc, char **argv) { + checkImageSupport(); + bool testResult = true; testResult = testResult && runTest(256, 256, -3.9, 6.1); diff --git a/tests/src/texture/hipTextureObj3DCheckModes.cpp b/tests/src/texture/hipTextureObj3DCheckModes.cpp index 3c71ee2af8..25390f6c7e 100644 --- a/tests/src/texture/hipTextureObj3DCheckModes.cpp +++ b/tests/src/texture/hipTextureObj3DCheckModes.cpp @@ -16,6 +16,7 @@ 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 int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; @@ -23,6 +24,7 @@ __global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX, normalizedCoords ? (y + offsetY) / height : y + offsetY, normalizedCoords ? (z + offsetZ) / depth : z + offsetZ); +#endif } template @@ -119,6 +121,8 @@ line1: } int main(int argc, char **argv) { + checkImageSupport(); + bool testResult = true; testResult = testResult && runTest(256, 256, 256, -3.9, 6.1, 9.5); diff --git a/tests/src/texture/hipTextureRef2D.cpp b/tests/src/texture/hipTextureRef2D.cpp index 533dddde00..7fa5d3930d 100644 --- a/tests/src/texture/hipTextureRef2D.cpp +++ b/tests/src/texture/hipTextureRef2D.cpp @@ -1,5 +1,5 @@ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_RUNTIME rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ @@ -14,14 +14,18 @@ texture tex; __global__ void tex2DKernel(float* outputData, int width, int height) { +#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(tex, x, y); +#endif } int runTest(int argc, char** argv); int main(int argc, char** argv) { + checkImageSupport(); + int testResult = runTest(argc, argv); if (testResult) { passed(); @@ -54,8 +58,8 @@ int runTest(int argc, char** argv) { hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - tex.addressMode[0] = hipAddressModeWrap; - tex.addressMode[1] = hipAddressModeWrap; + tex.addressMode[0] = hipAddressModeClamp; + tex.addressMode[1] = hipAddressModeClamp; tex.filterMode = hipFilterModePoint; tex.normalized = 0; diff --git a/tests/src/texture/simpleTexture2DLayered.cpp b/tests/src/texture/simpleTexture2DLayered.cpp index 7709a8cbec..b73d909baf 100644 --- a/tests/src/texture/simpleTexture2DLayered.cpp +++ b/tests/src/texture/simpleTexture2DLayered.cpp @@ -21,7 +21,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 */ @@ -34,9 +34,11 @@ texture tex2DL; __global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__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(tex2DL, x, y, layer); +#endif } //////////////////////////////////////////////////////////////////////////////// @@ -67,8 +69,8 @@ void runTest(int width,int height,int num_layers,textureaddressMode[0] = hipAddressModeWrap; - tex->addressMode[1] = hipAddressModeWrap; + tex->addressMode[0] = hipAddressModeClamp; + tex->addressMode[1] = hipAddressModeClamp; tex->filterMode = hipFilterModePoint; tex->normalized = false; @@ -103,6 +105,8 @@ void runTest(int width,int height,int num_layers,textureaddressMode[0] = hipAddressModeWrap; - tex->addressMode[1] = hipAddressModeWrap; + tex->addressMode[0] = hipAddressModeClamp; + tex->addressMode[1] = hipAddressModeClamp; tex->filterMode = hipFilterModePoint; tex->normalized = false; @@ -126,13 +126,8 @@ void runTest(int width,int height,int depth,texture