diff --git a/hipamd/include/hip/amd_detail/amd_surface_functions.h b/hipamd/include/hip/amd_detail/amd_surface_functions.h index 544a249dca..bdba2ef16e 100644 --- a/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -33,9 +33,6 @@ THE SOFTWARE. #define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \ unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj; -template -using __hip_is_surf_channel_type = __hip_is_tex_channel_type; - // CUDA is using byte address, need map to pixel address for HIP static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) { /* @@ -93,7 +90,7 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x, int boundaryMode = hipBoundaryModeZero) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT @@ -104,7 +101,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); @@ -114,7 +111,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); @@ -124,7 +121,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); @@ -134,7 +131,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); @@ -144,7 +141,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); @@ -154,7 +151,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); @@ -164,7 +161,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); @@ -174,7 +171,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); @@ -184,7 +181,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); @@ -194,7 +191,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); @@ -204,7 +201,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); @@ -214,7 +211,7 @@ static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT @@ -225,7 +222,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac template < typename T, - typename std::enable_if<__hip_is_surf_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT diff --git a/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/hipamd/include/hip/amd_detail/texture_fetch_functions.h index a252c0efde..f1cd70ab05 100644 --- a/hipamd/include/hip/amd_detail/texture_fetch_functions.h +++ b/hipamd/include/hip/amd_detail/texture_fetch_functions.h @@ -37,7 +37,7 @@ THE SOFTWARE. unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; template -struct __hip_is_tex_channel_type +struct __hip_is_tex_surf_scalar_channel_type { static constexpr bool value = std::is_same::value || @@ -49,13 +49,20 @@ struct __hip_is_tex_channel_type std::is_same::value; }; +template +struct __hip_is_tex_surf_channel_type +{ + static constexpr bool value = + __hip_is_tex_surf_scalar_channel_type::value; +}; + template< typename T, unsigned int rank> -struct __hip_is_tex_channel_type> +struct __hip_is_tex_surf_channel_type> { static constexpr bool value = - __hip_is_tex_channel_type::value && + __hip_is_tex_surf_scalar_channel_type::value && ((rank == 1) || (rank == 2) || (rank == 4)); @@ -98,7 +105,7 @@ struct __hip_tex_ret template __forceinline__ __device__ typename std::enable_if< - __hip_is_tex_channel_type::value && std::is_scalar::value, const T>::type + __hip_is_tex_surf_scalar_channel_type::value, const T>::type __hipMapFrom(const U &u) { if constexpr (sizeof(T) < sizeof(float)) { union { @@ -120,7 +127,8 @@ __hipMapFrom(const U &u) { */ template __forceinline__ __device__ -typename std::enable_if<__hip_is_tex_channel_type::value, const T>::type +typename std::enable_if< + __hip_is_tex_surf_scalar_channel_type::value, const T>::type __hipMapFrom(const U &u) { if constexpr (sizeof(typename T::value_type) < sizeof(float)) { union { @@ -143,7 +151,7 @@ __hipMapFrom(const U &u) { template __forceinline__ __device__ typename std::enable_if< - __hip_is_tex_channel_type::value && std::is_scalar::value, const U>::type +__hip_is_tex_surf_scalar_channel_type::value, const U>::type __hipMapTo(const T &t) { if constexpr (sizeof(T) < sizeof(float)) { union { @@ -167,7 +175,8 @@ __hipMapTo(const T &t) { */ template __forceinline__ __device__ -typename std::enable_if<__hip_is_tex_channel_type::value, const U>::type +typename std::enable_if< + __hip_is_tex_surf_scalar_channel_type::value, const U>::type __hipMapTo(const T &t) { if constexpr (sizeof(typename T::value_type) < sizeof(float)) { union { @@ -195,7 +204,7 @@ template struct __hip_tex_ret< T, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> { using type = T; }; @@ -206,7 +215,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -430,7 +439,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type::value, bool>::type> { using type = HIP_vector_type; }; @@ -441,7 +450,7 @@ template< struct __hip_tex2dgather_ret< HIP_vector_type, hipReadModeElementType, - typename std::enable_if<__hip_is_tex_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_surf_channel_type>::value, bool>::type> { using type = HIP_vector_type; }; diff --git a/hipamd/include/hip/amd_detail/texture_indirect_functions.h b/hipamd/include/hip/amd_detail/texture_indirect_functions.h index af0bfa4cec..85ebfa862f 100644 --- a/hipamd/include/hip/amd_detail/texture_indirect_functions.h +++ b/hipamd/include/hip/amd_detail/texture_indirect_functions.h @@ -37,12 +37,9 @@ THE SOFTWARE. unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \ unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; -template -using __hip_is_itex_channel_type = __hip_is_tex_channel_type; - template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, int x) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -52,7 +49,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x) { *ptr = tex1Dfetch(textureObject, x); @@ -60,7 +57,7 @@ static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t tex template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, float x) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -70,7 +67,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x) { *ptr = tex1D(textureObject, x); @@ -78,7 +75,7 @@ static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -88,7 +85,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y) { *ptr = tex2D(textureObject, x, y); @@ -96,7 +93,7 @@ static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -106,7 +103,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z) { *ptr = tex3D(textureObject, x, y, z); @@ -114,7 +111,7 @@ static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -124,7 +121,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer) { *ptr = tex1DLayered(textureObject, x, layer); @@ -132,7 +129,7 @@ static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t t template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -142,7 +139,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer) { *ptr = tex1DLayered(textureObject, x, y, layer); @@ -150,7 +147,7 @@ static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t t template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -160,7 +157,7 @@ static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z) { *ptr = texCubemap(textureObject, x, y, z); @@ -168,7 +165,7 @@ static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t tex template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -178,7 +175,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer) { *ptr = texCubemapLayered(textureObject, x, y, z, layer); @@ -186,7 +183,7 @@ static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObjec template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -217,7 +214,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0) { *ptr = texCubemapLayered(textureObject, x, y, comp); @@ -225,7 +222,7 @@ static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t te template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -235,7 +232,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level) { *ptr = tex1DLod(textureObject, x, level); @@ -243,7 +240,7 @@ static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -253,7 +250,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level) { *ptr = tex2DLod(textureObject, x, y, level); @@ -261,7 +258,7 @@ static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -271,7 +268,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level) { *ptr = tex3DLod(textureObject, x, y, z, level); @@ -279,7 +276,7 @@ static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textu template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -289,7 +286,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level) { *ptr = tex1DLayeredLod(textureObject, x, layer, level); @@ -297,7 +294,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_ template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -307,7 +304,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level) { *ptr = tex2DLayeredLod(textureObject, x, y, layer, level); @@ -315,7 +312,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_ template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -325,7 +322,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level) { *ptr = texCubemapLod(textureObject, x, y, z, level); @@ -333,7 +330,7 @@ static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -345,7 +342,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { *ptr = texCubemapGrad(textureObject, x, y, z, dPdx, dPdy); @@ -353,7 +350,7 @@ static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -363,7 +360,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level) { *ptr = texCubemapLayeredLod(textureObject, x, y, z, layer, level); @@ -371,7 +368,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureOb template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -381,7 +378,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy) { *ptr = tex1DGrad(textureObject, x, dPdx, dPdy); @@ -389,7 +386,7 @@ static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -399,7 +396,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy) { *ptr = tex2DGrad(textureObject, x, y, dPdx, dPdy); @@ -407,7 +404,7 @@ static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -417,7 +414,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { *ptr = tex3DGrad(textureObject, x, y, z, dPdx, dPdy); @@ -425,7 +422,7 @@ static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t text template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -435,7 +432,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy) { *ptr = tex1DLayeredGrad(textureObject, x, layer, dPdx, dPdy); @@ -443,7 +440,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -453,7 +450,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy) { *ptr = tex2DLayeredGrad(textureObject, x, y, layer, dPdx, dPdy); @@ -461,7 +458,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -473,7 +470,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t te template < typename T, - typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> + typename std::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy) { *ptr = texCubemapLayeredGrad(textureObject, x, y, z, layer, dPdx, dPdy);