diff --git a/bin/hipcc.pl b/bin/hipcc.pl index 1ef3a90e30..a2ebb9f330 100755 --- a/bin/hipcc.pl +++ b/bin/hipcc.pl @@ -152,7 +152,10 @@ if ($HIP_PLATFORM eq "amd") { $HIPCC="$HIP_CLANG_PATH/clang" . $execExtension; $HIPLDFLAGS = "--driver-mode=g++"; } - + # to avoid using dk linker or MSVC linker + if($isWindows) { + $HIPLDFLAGS .= " -fuse-ld=lld"; + } $HIP_CLANG_VERSION = `$HIPCC --version`; $HIP_CLANG_VERSION=~/.*clang version (\S+).*/; $HIP_CLANG_VERSION=$1; diff --git a/docs/markdown/hip_debugging.md b/docs/markdown/hip_debugging.md index 9127085a4d..b26915c50e 100644 --- a/docs/markdown/hip_debugging.md +++ b/docs/markdown/hip_debugging.md @@ -261,7 +261,7 @@ The following is the summary of the most useful environment variables in HIP. | AMD_SERIALIZE_KERNEL
Serialize kernel enqueue. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | AMD_SERIALIZE_COPY
Serialize copies. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | HIP_HOST_COHERENT
Coherent memory in hipHostMalloc. | 0 | 0: memory is not coherent between host and GPU.
1: memory is coherent with host. | -| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 0 | 0: Disable.
1: Enable. | +| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 1 | 0: Disable.
1: Enable. | ## General Debugging Tips diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index a26597ca67..71dd1084ca 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -57,7 +57,6 @@ The HIP API documentation describes each API and its limitations, if any, compar At a high-level, the following features are not supported: - Textures (partial support available) - Dynamic parallelism (CUDA 5.0) -- Managed memory (CUDA 6.5) - Graphics interoperability with OpenGL or Direct3D - CUDA IPC Functions (Under Development) - CUDA array, mipmappedArray and pitched memory @@ -70,8 +69,7 @@ See the [API Support Table](CUDA_Runtime_API_functions_supported_by_HIP.md) for - Virtual functions, indirect functions and try/catch (CUDA 4.0) - `__prof_trigger` - PTX assembly (CUDA 4.0). HIP-Clang supports inline GCN assembly. -- Several kernel features are under development. See the [HIP Kernel Language](hip_kernel_language.md) for more information. These include: - - printf +- Several kernel features are under development. See the [HIP Kernel Language](hip_kernel_language.md) for more information. ### Is HIP a drop-in replacement for CUDA? @@ -233,6 +231,7 @@ See the [HIP Logging](hip_logging.md) for more information. ### What is maximum limit of kernel launching parameter? Product of block.x, block.y, and block.z should be less than 1024. +Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32, so gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32. ### Are __shfl_*_sync functions supported on HIP platform? __shfl_*_sync is not supported on HIP but for nvcc path CUDA 9.0 and above all shuffle calls get redirected to it's sync version. diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 9af6a8ec3a..81b3e076be 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -184,7 +184,7 @@ The `__restrict__` keyword tells the compiler that the associated memory pointer ### Coordinate Built-Ins Built-ins determine the coordinate of the active work item in the execution grid. They are defined in amd_hip_runtime.h (rather than being implicitly defined by the compiler). In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance: -threadIdx.x, blockIdx.y, gridDim.y, etc. +threadIdx.x, blockIdx.y, gridDim.y, etc. The products gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32. ### warpSize @@ -696,7 +696,7 @@ void assert(int input) ``` There are two kinds of implementations for assert functions depending on the use sceneries, -- One is for the host version of assert, which is defined in assert.h, +- One is for the host version of assert, which is defined in assert.h, - Another is the device version of assert, which is implemented in hip/hip_runtime.h. Users need to include assert.h to use assert. For assert to work in both device and host functions, users need to include "hip/hip_runtime.h". diff --git a/docs/markdown/hip_programming_guide.md b/docs/markdown/hip_programming_guide.md index d5549659da..e9d25b06bc 100644 --- a/docs/markdown/hip_programming_guide.md +++ b/docs/markdown/hip_programming_guide.md @@ -60,9 +60,10 @@ HIP supports Stream Memory Operations to enable direct synchronization between N hipStreamWriteValue64 Note, CPU access to the semaphore's memory requires volatile keyword to disable CPU compiler's optimizations on memory access. - For more details, please check the documentation HIP-API.pdf. +Please note, HIP stream does not gurantee concurrency on AMD hardware for the case of multiple (at least 6) long running streams executing concurrently, using hipStreamSynchronize(nullptr) for synchronization. + ### Coherency Controls ROCm defines two coherency options for host memory: - Coherent memory : Supports fine-grain synchronization while the kernel is running.  For example, a kernel can perform atomic operations that are visible to the host CPU or to other (peer) GPUs.  Synchronization instructions include threadfence_system and C++11-style atomic operations. @@ -130,7 +131,10 @@ The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hi ## Device-Side Malloc -HIP-Clang currently doesn't supports device-side malloc and free. +HIP-Clang now supports device-side malloc and free. +This implementation does not require the use of `hipDeviceSetLimit(hipLimitMallocHeapSize,value)` nor respects any setting. The heap is fully dynamic and can grow until the available free memory on the device is consumed. + +The test codes in the link (https://github.com/ROCm-Developer-Tools/HIP/blob/develop/tests/src/deviceLib/hipDeviceMalloc.cpp) show how to implement application using malloc and free functions in device kernels. ## Use of Long Double Type diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 8462ceea40..930e8d1854 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -6177,12 +6177,15 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph * Memory allocation properties */ typedef struct hipMemAllocationProp { - unsigned char compressionType; ///< Compression type - hipMemLocation location; ///< Memory location - hipMemAllocationHandleType requestedHandleType; ///< Requested handle type hipMemAllocationType type; ///< Memory allocation type - unsigned short usage; ///< Usage + hipMemAllocationHandleType requestedHandleType; ///< Requested handle type + hipMemLocation location; ///< Memory location void* win32HandleMetaData; ///< Metadata for Win32 handles + struct { + unsigned char compressionType; ///< Compression type + unsigned char gpuDirectRDMACapable; ///< RDMA capable + unsigned short usage; ///< Usage + } allocFlags; } hipMemAllocationProp; /** 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/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 55af77e7df..a5b6edab67 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -65,7 +65,8 @@ "Unit_hipHostMalloc_NonCoherent", "Unit_hipHostMalloc_Coherent", "Unit_hipHostMalloc_Default", - "Unit_hipStreamCreate_MultistreamBasicFunctionalities" + "Unit_hipStreamCreate_MultistreamBasicFunctionalities", + "Unit_hipEventIpc" ] } 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/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 84be08bc33..89027dcbf5 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -10,17 +10,11 @@ set(TEST_SRC popc.cc ldg.cc threadfence_system.cc + syncthreadsand.cc + syncthreadscount.cc + syncthreadsor.cc ) -# skipped for windows compiler issue - Illegal instruction detected -if(UNIX) - set(TEST_SRC ${TEST_SRC} - syncthreadsand.cc - syncthreadscount.cc - syncthreadsor.cc) -endif() - - # AMD only tests set(AMD_TEST_SRC unsafeAtomicAdd.cc diff --git a/tests/catch/unit/event/CMakeLists.txt b/tests/catch/unit/event/CMakeLists.txt index 6b6ef7c3c4..fae1b4264c 100644 --- a/tests/catch/unit/event/CMakeLists.txt +++ b/tests/catch/unit/event/CMakeLists.txt @@ -3,15 +3,10 @@ set(TEST_SRC Unit_hipEvent_Negative.cc Unit_hipEvent.cc Unit_hipEventElapsedTime.cc + Unit_hipEventRecord.cc + Unit_hipEventIpc.cc ) -# skipped for windows due to duplicate symbols -if(UNIX) - set(TEST_SRC ${TEST_SRC} - Unit_hipEventRecord.cc - Unit_hipEventIpc.cc) -endif() - hip_add_exe_to_target(NAME EventTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index a9d3c35d10..cf316734dd 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -10,15 +10,9 @@ set(TEST_SRC hipStreamGetCUMask.cc hipAPIStreamDisable.cc streamCommon.cc -) - -#skipped in windows - duplicate HipTest::vector_square sym (compiler issue) -if(UNIX) - set(TEST_SRC ${TEST_SRC} - hipStreamWithCUMask.cc - hipStreamACb_MultiThread.cc) -endif() - + hipStreamWithCUMask.cc + hipStreamACb_MultiThread.cc + ) else() set(TEST_SRC hipStreamCreate.cc @@ -31,7 +25,7 @@ set(TEST_SRC hipStreamCreateWithPriority.cc hipAPIStreamDisable.cc streamCommon.cc -) + ) endif() hip_add_exe_to_target(NAME StreamTest 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