diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 03bcd9810e..70a02acb29 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -494,7 +494,11 @@ Following is the list of supported floating-point intrinsics. Note that intrinsi | double __dsqrt_rn ( double x )
Compute `√x` in round-to-nearest-even mode. | ## Texture Functions -The supported Texture functions are listed in header files "texture_functions.h"(https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_functions.h) and"texture_indirect_functions.h" (https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_indirect_functions.h). +The supported Texture functions are listed in header files "texture_fetch_functions.h"(https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_fetch_functions.h) and"texture_indirect_functions.h" (https://github.com/ROCm-Developer-Tools/HIP/blob/main/include/hip/hcc_detail/texture_indirect_functions.h). + +Texture functions are not supported on some devices. +Macro __HIP_NO_IMAGE_SUPPORT == 1 can be used to check whether texture functions are not supported in device code. +Attribute hipDeviceAttributeImageSupport can be queried to check whether texture functions are supported in host runtime code. ## Surface Functions Surface functions are not supported. diff --git a/include/hip/driver_types.h b/include/hip/driver_types.h index ec15687492..fb02516354 100644 --- a/include/hip/driver_types.h +++ b/include/hip/driver_types.h @@ -436,7 +436,41 @@ typedef enum hipFunction_attribute { HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT, HIP_FUNC_ATTRIBUTE_MAX -}hipFunction_attribute; +} hipFunction_attribute; + +typedef enum hipPointer_attribute { + HIP_POINTER_ATTRIBUTE_CONTEXT = 1, ///< The context on which a pointer was allocated + ///< @warning - not supported in HIP + HIP_POINTER_ATTRIBUTE_MEMORY_TYPE, ///< memory type describing location of a pointer + HIP_POINTER_ATTRIBUTE_DEVICE_POINTER,///< address at which the pointer is allocated on device + HIP_POINTER_ATTRIBUTE_HOST_POINTER, ///< address at which the pointer is allocated on host + HIP_POINTER_ATTRIBUTE_P2P_TOKENS, ///< A pair of tokens for use with linux kernel interface + ///< @warning - not supported in HIP + HIP_POINTER_ATTRIBUTE_SYNC_MEMOPS, ///< Synchronize every synchronous memory operation + ///< initiated on this region + HIP_POINTER_ATTRIBUTE_BUFFER_ID, ///< Unique ID for an allocated memory region + HIP_POINTER_ATTRIBUTE_IS_MANAGED, ///< Indicates if the pointer points to managed memory + HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL,///< device ordinal of a device on which a pointer + ///< was allocated or registered + HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE, ///< if this pointer maps to an allocation + ///< that is suitable for hipIpcGetMemHandle + ///< @warning - not supported in HIP + HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR,///< Starting address for this requested pointer + HIP_POINTER_ATTRIBUTE_RANGE_SIZE, ///< Size of the address range for this requested pointer + HIP_POINTER_ATTRIBUTE_MAPPED, ///< tells if this pointer is in a valid address range + ///< that is mapped to a backing allocation + HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES,///< Bitmask of allowed hipmemAllocationHandleType + ///< for this allocation @warning - not supported in HIP + HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE, ///< returns if the memory referenced by + ///< this pointer can be used with the GPUDirect RDMA API + ///< @warning - not supported in HIP + HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS, ///< Returns the access flags the device associated with + ///< for the corresponding memory referenced by the ptr + HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE ///< Returns the mempool handle for the allocation if + ///< it was allocated from a mempool + ///< @warning - not supported in HIP +} hipPointer_attribute; + #endif // !defined(__HIPCC_RTC__) #else #error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__"); diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index f5311e767d..6fb56a0325 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -442,7 +442,8 @@ typedef enum hipDeviceAttribute_t { hipDeviceAttributeIsLargeBar, ///< Whether it is LargeBar hipDeviceAttributeAsicRevision, ///< Revision of the GPU in this device hipDeviceAttributeCanUseStreamWaitValue, ///< '1' if Device supports hipStreamWaitValue32() and - ///< hipStreamWaitValue64() , '0' otherwise. + ///< hipStreamWaitValue64(), '0' otherwise. + hipDeviceAttributeImageSupport, ///< '1' if Device supports image, '0' otherwise. hipDeviceAttributeAmdSpecificEnd = 19999, hipDeviceAttributeVendorSpecificBegin = 20000, @@ -1976,15 +1977,48 @@ hipError_t hipEventQuery(hipEvent_t event); /** * @brief Return attributes for the specified pointer * - * @param[out] attributes for the specified pointer - * @param[in] pointer to get attributes for + * @param [out] attributes attributes for the specified pointer + * @param [in] ptr pointer to get attributes for * * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue * - * @see hipGetDeviceCount, hipGetDevice, hipSetDevice, hipChooseDevice + * @see hipPointerGetAttribute */ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr); - +/** + * @brief Returns information about the specified pointer.[BETA] + * + * @param [in, out] data returned pointer attribute value + * @param [in] atribute attribute to query for + * @param [in] ptr pointer to get attributes for + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue + * + * @beta This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + * + * @see hipPointerGetAttributes + */ +hipError_t hipPointerGetAttribute(void* data, hipPointer_attribute attribute, + hipDeviceptr_t ptr); +/** + * @brief Returns information about the specified pointer.[BETA] + * + * @param [in] numAttributes number of attributes to query for + * @param [in] attributes attributes to query for + * @param [in, out] data a two-dimensional containing pointers to memory locations + * where the result of each attribute query will be written to + * @param [in] ptr pointer to get attributes for + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue + * + * @beta This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + * + * @see hipPointerGetAttribute + */ +hipError_t hipDrvPointerGetAttributes(unsigned int numAttributes, hipPointer_attribute* attributes, + void** data, hipDeviceptr_t ptr); /** * @brief Imports an external semaphore. * diff --git a/tests/src/runtimeApi/memory/hipMemcpy2DFromArray.cpp b/tests/src/runtimeApi/memory/hipMemcpy2DFromArray.cpp index 076cfd12d8..c8ff891eba 100644 --- a/tests/src/runtimeApi/memory/hipMemcpy2DFromArray.cpp +++ b/tests/src/runtimeApi/memory/hipMemcpy2DFromArray.cpp @@ -173,6 +173,7 @@ bool Memcpy2DFromArray::hipMemcpy2DFromArray_PeerDeviceContext() { err = hipMemcpy2DFromArray(A_h, width, A_d, 0, 0, width, NUM_H, hipMemcpyDeviceToHost); + if (err == hipSuccess) { TestPassed = ValidateResult(A_h, INITIAL_VAL); } else { @@ -281,7 +282,11 @@ int main(int argc, char **argv) { TestPassed &= Array_obj.hipMemcpy2DFromArray_SizeCheck(); } else if (p_tests == 3) { if (numDevices > 1) { +#ifndef _WIN64 TestPassed &= Array_obj.hipMemcpy2DFromArray_PeerDeviceContext(); +#else + printf("xgmi memory test not supported on windows\n"); +#endif } else { printf("skipped the testcase as noof devices <2\n"); } @@ -289,8 +294,12 @@ int main(int argc, char **argv) { TestPassed &= Array_obj.hipMemcpy2DFromArray_NegativeTests(); } else if (p_tests == 5) { if (numDevices > 1) { +#ifndef _WIN64 TestPassed &= Array_obj.hipMemcpy2DFromArray_PinnedHostMemory_SameGPU(); TestPassed &= Array_obj.hipMemcpy2DFromArray_PinnedHostMemory_PeerGPU(); +#else + printf("xgmi memory test not supported on windows\n"); +#endif } else { printf("skipped the testcases as noof devices <2\n"); } diff --git a/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp b/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp index 0e5ba8e453..08d3e6b66f 100644 --- a/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp +++ b/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp @@ -209,7 +209,9 @@ int main(int argc, char **argv) { } else if (p_tests == 2) { TestPassed = AtoH_obj.hipMemcpyAtoH_ByteCountZero(); } else if (p_tests == 3) { +#ifndef _WIN64 TestPassed = AtoH_obj.hipMemcpyAtoH_PeerDeviceContext(); +#endif } else if (p_tests == 4) { TestPassed = AtoH_obj.hipMemcpyAtoH_NegativeTests(); } else if (p_tests == 5) { diff --git a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp index 3bd8ed3f99..b4abf307dd 100644 --- a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp +++ b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp @@ -592,6 +592,13 @@ 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(); + } bool TestPassed = true; if (p_tests == 0x01) { TestPassed = testTexType(HIP_AD_FORMAT_FLOAT, diff --git a/tests/src/runtimeApi/module/tex2d_kernel.cpp b/tests/src/runtimeApi/module/tex2d_kernel.cpp index 5b558ffbf2..5ba0ebf89f 100644 --- a/tests/src/runtimeApi/module/tex2d_kernel.cpp +++ b/tests/src/runtimeApi/module/tex2d_kernel.cpp @@ -30,36 +30,44 @@ __device__ float deviceGlobalFloat; extern "C" __global__ void tex2dKernelFloat(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; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(ftex, x, y); } +#endif } extern "C" __global__ void tex2dKernelInt(int* 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; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(itex, x, y); } +#endif } extern "C" __global__ void tex2dKernelInt16(short* 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; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(stex, x, y); } +#endif } extern "C" __global__ void tex2dKernelInt8(char* 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; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(ctex, x, y); } +#endif } diff --git a/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp b/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp index 2d5b605de4..a1f2850320 100644 --- a/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp +++ b/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp @@ -42,8 +42,12 @@ int main() { HIPCHECK(hipGetDeviceCount(&count)) if (count >= 2){ +#ifdef _WIN64 + std::cout<<"XGMI p2p attribute test no supported on windows"< texus; template __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT unsigned int elementID = hipThreadIdx_x; if(elementID >= numElements) return; @@ -72,6 +73,7 @@ __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) pDst[elementID] = tex1D(texs, coord); else if(std::is_same::value) pDst[elementID] = tex1D(texus, coord); +#endif } bool textureVerifyFilterModePoint(float *hOutputData, float *expected, size_t size) { @@ -171,6 +173,13 @@ 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; bool status = false; diff --git a/tests/src/texture/hipTexObjPitch.cpp b/tests/src/texture/hipTexObjPitch.cpp index 6869081a20..2e59d91d7c 100644 --- a/tests/src/texture/hipTexObjPitch.cpp +++ b/tests/src/texture/hipTexObjPitch.cpp @@ -29,11 +29,12 @@ THE SOFTWARE. // texture object is a kernel argument template __global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) { - +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT for(int i =0;i(texObj, j, i); __syncthreads(); +#endif } template @@ -76,7 +77,7 @@ void texture2Dtest() texDescr.readMode = hipReadModeElementType; hipTextureObject_t texObj; - HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL)); + HIPCHECK(hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL)); HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; @@ -95,6 +96,13 @@ void texture2Dtest() int main() { + int imageSupport = 0; + hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, + p_gpuDevice); + if (!imageSupport) { + printf("Texture is not support on the device\n"); + passed(); + } texture2Dtest(); texture2Dtest(); texture2Dtest(); diff --git a/tests/src/texture/hipTextureMipmapObj2D.cpp b/tests/src/texture/hipTextureMipmapObj2D.cpp index 719a427f0b..3a59a9640d 100644 --- a/tests/src/texture/hipTextureMipmapObj2D.cpp +++ b/tests/src/texture/hipTextureMipmapObj2D.cpp @@ -40,10 +40,12 @@ std::vector mip_vector = {8, 4, 2, 1}; __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width, int height, float level) { #ifndef __gfx90a__ +#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] = tex2DLod(textureObject, x, y, level); #endif +#endif } bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) { @@ -148,6 +150,7 @@ bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_ hipDestroyTextureObject(textureObject); hipFree(dData); hipFreeArray(hipArray); + free(hData); return testResult; } @@ -169,7 +172,13 @@ bool runTest(int argc, char** argv) { int main(int argc, char** argv) { 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/hipTextureObj2D.cpp b/tests/src/texture/hipTextureObj2D.cpp index 649c748fd2..18f40eab11 100644 --- a/tests/src/texture/hipTextureObj2D.cpp +++ b/tests/src/texture/hipTextureObj2D.cpp @@ -12,14 +12,23 @@ __global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, 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(textureObject, x, y); +#endif } 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(); + } int testResult = runTest(argc, argv); if (testResult) { @@ -70,7 +79,6 @@ int runTest(int argc, char** argv) { // Create texture object hipTextureObject_t textureObject = 0; hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); - float* dData = NULL; hipMalloc((void**)&dData, size); @@ -104,5 +112,6 @@ int runTest(int argc, char** argv) { hipDestroyTextureObject(textureObject); hipFree(dData); hipFreeArray(hipArray); + free(hData); return testResult; } diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp index a156ff46bf..06f2a31c97 100644 --- a/tests/src/texture/simpleTexture3D.cpp +++ b/tests/src/texture/simpleTexture3D.cpp @@ -42,6 +42,7 @@ __global__ void simpleKernel3DArray(T* outputData, int width, int height,int depth) { +#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__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++) { @@ -54,6 +55,7 @@ __global__ void simpleKernel3DArray(T* outputData, } } } +#endif } //////////////////////////////////////////////////////////////////////////////// @@ -127,6 +129,13 @@ void runTest(int width,int height,int depth,texture