From 117f0ab1020eb7ea14c6366ce528d7eaed5166b0 Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Tue, 17 Mar 2020 13:09:33 -0400 Subject: [PATCH] Add constraints to texture indirect functions Similar to the previous patch, this change adds type constraints to texture indirect functions. Since we don't have to deduce the return type for these, we simply just have to check if the user provided a valid channel type. Change-Id: Ia094bd6126e01df2ea90902c9aa59cb6cfe85773 --- .../hip/hcc_detail/texture_fetch_functions.h | 26 +-- .../hcc_detail/texture_indirect_functions.h | 211 ++++++++++++++---- 2 files changed, 178 insertions(+), 59 deletions(-) diff --git a/include/hip/hcc_detail/texture_fetch_functions.h b/include/hip/hcc_detail/texture_fetch_functions.h index 39ac3ecd89..23a850c52c 100644 --- a/include/hip/hcc_detail/texture_fetch_functions.h +++ b/include/hip/hcc_detail/texture_fetch_functions.h @@ -35,7 +35,7 @@ THE SOFTWARE. unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; template -struct __hip_is_channel_type +struct __hip_is_tex_channel_type { static constexpr bool value = std::is_same::value || @@ -50,17 +50,17 @@ struct __hip_is_channel_type template< typename T, unsigned int rank> -struct __hip_is_channel_type> +struct __hip_is_tex_channel_type> { static constexpr bool value = - __hip_is_channel_type::value && + __hip_is_tex_channel_type::value && ((rank == 1) || (rank == 2) || (rank == 4)); }; template -struct __hip_is_normalized_channel_type +struct __hip_is_tex_normalized_channel_type { static constexpr bool value = std::is_same::value || @@ -72,10 +72,10 @@ struct __hip_is_normalized_channel_type template< typename T, unsigned int rank> -struct __hip_is_normalized_channel_type> +struct __hip_is_tex_normalized_channel_type> { static constexpr bool value = - __hip_is_normalized_channel_type::value && + __hip_is_tex_normalized_channel_type::value && ((rank == 1) || (rank == 2) || (rank == 4)); @@ -99,7 +99,7 @@ template struct __hip_tex_ret< T, hipReadModeElementType, - typename std::enable_if<__hip_is_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_channel_type::value, bool>::type> { using type = T; }; @@ -110,7 +110,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeElementType, - typename std::enable_if<__hip_is_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -119,7 +119,7 @@ template struct __hip_tex_ret< T, hipReadModeNormalizedFloat, - typename std::enable_if<__hip_is_normalized_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> { using type = float; }; @@ -130,7 +130,7 @@ template< struct __hip_tex_ret< HIP_vector_type, hipReadModeNormalizedFloat, - typename std::enable_if<__hip_is_normalized_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_normalized_channel_type>::value, bool>::type> { using type = HIP_vector_type<__hip_tex_ret_t, rank>; }; @@ -333,7 +333,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeElementType, - typename std::enable_if<__hip_is_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_channel_type::value, bool>::type> { using type = HIP_vector_type; }; @@ -344,7 +344,7 @@ template< struct __hip_tex2dgather_ret< HIP_vector_type, hipReadModeElementType, - typename std::enable_if<__hip_is_channel_type>::value, bool>::type> + typename std::enable_if<__hip_is_tex_channel_type>::value, bool>::type> { using type = HIP_vector_type; }; @@ -353,7 +353,7 @@ template struct __hip_tex2dgather_ret< T, hipReadModeNormalizedFloat, - typename std::enable_if<__hip_is_normalized_channel_type::value, bool>::type> + typename std::enable_if<__hip_is_tex_normalized_channel_type::value, bool>::type> { using type = float4; }; diff --git a/include/hip/hcc_detail/texture_indirect_functions.h b/include/hip/hcc_detail/texture_indirect_functions.h index 56784f433a..4facce436e 100644 --- a/include/hip/hcc_detail/texture_indirect_functions.h +++ b/include/hip/hcc_detail/texture_indirect_functions.h @@ -28,11 +28,40 @@ THE SOFTWARE. #include #include +#include + #define TEXTURE_OBJECT_PARAMETERS_INIT \ 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 +template +struct __hip_is_itex_channel_type +{ + static constexpr bool value = + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value; +}; + +template< + typename T, + unsigned int rank> +struct __hip_is_itex_channel_type> +{ + static constexpr bool value = + __hip_is_itex_channel_type::value && + ((rank == 1) || + (rank == 2) || + (rank == 4)); +}; + +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -40,13 +69,17 @@ static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x) return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x) { *ptr = tex1Dfetch(textureObject, x); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1D(hipTextureObject_t textureObject, float x) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -54,13 +87,17 @@ static __device__ T tex1D(hipTextureObject_t textureObject, float x) return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x) { *ptr = tex1D(textureObject, x); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -68,13 +105,17 @@ static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y) return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y) { *ptr = tex2D(textureObject, x, y); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -82,13 +123,17 @@ static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, fl return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z) { *ptr = tex3D(textureObject, x, y, z); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -96,13 +141,17 @@ static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer) { *ptr = tex1DLayered(textureObject, x, layer); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -110,13 +159,17 @@ static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, floa return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer) { *ptr = tex1DLayered(textureObject, x, y, layer); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -124,13 +177,17 @@ static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z) { *ptr = texCubemap(textureObject, x, y, z); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -138,13 +195,17 @@ static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x, return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer) { *ptr = texCubemapLayered(textureObject, x, y, z, layer); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -173,13 +234,17 @@ static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float return {}; } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0) { *ptr = texCubemapLayered(textureObject, x, y, comp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -187,13 +252,17 @@ static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float le return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level) { *ptr = tex1DLod(textureObject, x, level); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -201,13 +270,17 @@ static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level) { *ptr = tex2DLod(textureObject, x, y, level); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -215,13 +288,17 @@ static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level) { *ptr = tex3DLod(textureObject, x, y, z, level); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -229,13 +306,17 @@ static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, i return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level) { *ptr = tex1DLayeredLod(textureObject, x, layer, level); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -243,13 +324,17 @@ static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level) { *ptr = tex2DLayeredLod(textureObject, x, y, layer, level); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -257,13 +342,17 @@ static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, flo return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level) { *ptr = texCubemapLod(textureObject, x, y, z, level); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -273,13 +362,17 @@ static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, fl return {}; } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ 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); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -287,13 +380,17 @@ static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ 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); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -301,13 +398,17 @@ static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float d return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy) { *ptr = tex1DGrad(textureObject, x, dPdx, dPdy); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -315,13 +416,17 @@ static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy) { *ptr = tex2DGrad(textureObject, x, y, dPdx, dPdy); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -329,13 +434,17 @@ static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ 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); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -343,13 +452,17 @@ static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy) { *ptr = tex1DLayeredGrad(textureObject, x, layer, dPdx, dPdy); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -357,13 +470,17 @@ static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, return *reinterpret_cast(&tmp); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ 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); } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy) { TEXTURE_OBJECT_PARAMETERS_INIT @@ -373,7 +490,9 @@ static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, flo return {}; } -template +template < + typename T, + typename std::enable_if<__hip_is_itex_channel_type::value>::type* = nullptr> static __device__ 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);