AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

SWDEV-303204 - Disable texture/image Apis on some devices

Kaynağa Gözat 
If the target device doesn't support image,
compiler will send out unavailable error if device texture apis
are called; and host texture apis will return hipErrorNotSupported.
Remove useless include/hip/amd_detail/texture_functions.h.
Add hipDeviceAttributeImageSupport attribute to check image
support in runtime.
Add __HIP_NO_IMAGE_SUPPORT(=1) macro in device code to indicate
texture apis not supported.
Change-Id: Ia891493e681a61980a2444859a90e046a9ebea39
Bu işleme şunda yer alıyor:
Tao Sang
2021-09-24 19:19:25 -04:00
işlemeyi yapan: Tao Sang
ebeveyn 69117705c8
işleme 5acc8b5102
7 değiştirilmiş dosya ile 329 ekleme ve 11204 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
hipamd/include/hip/amd_detail/host_defines.h
+7
Dosyayı Görüntüle
@@ -50,6 +50,12 @@ THE SOFTWARE.
#define __noinline__ __attribute__((noinline))
#define __forceinline__ inline __attribute__((always_inline))
#if __HIP_NO_IMAGE_SUPPORT
#define __hip_img_chk__ __attribute__((unavailable("The image/texture API not supported on the device")))
#else
#define __hip_img_chk__
#endif
#else
// Non-HCC compiler
@@ -67,6 +73,7 @@ THE SOFTWARE.
#define __shared__
#define __constant__
#define __hip_img_chk__
#endif
#endif
hipamd/include/hip/amd_detail/texture_fetch_functions.h
+23 -23
Dosyayı Görüntüle
@@ -138,7 +138,7 @@ struct __hip_tex_ret<
};
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t, int x)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t, int x)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_load_1Db(i, x);
@@ -146,7 +146,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(textur
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t, float x)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t, float x)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_1D(i, s, x);
@@ -154,7 +154,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1D(texture<T,
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t, float x, float y)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t, float x, float y)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
@@ -162,7 +162,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2D(texture<T,
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
@@ -170,7 +170,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayered(text
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
@@ -178,7 +178,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayered(text
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t, float x, float y, float z)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t, float x, float y, float z)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
@@ -186,7 +186,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3D(texture<T,
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
@@ -194,7 +194,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemap(textur
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t, float x, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t, float x, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
@@ -202,7 +202,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t, float x, float y, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t, float x, float y, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
@@ -210,7 +210,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_1Da(i, s, float2(x, layer).data, level);
@@ -218,7 +218,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(t
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_2Da(i, s, float4(x, y, layer, 0.0f).data, level);
@@ -226,7 +226,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(t
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
@@ -234,7 +234,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
@@ -242,7 +242,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLod(tex
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
@@ -250,7 +250,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredLod(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float level)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayeredLod(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
@@ -258,7 +258,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
// TODO missing in device libs.
@@ -268,7 +268,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapGrad(te
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
// TODO missing in device libs.
@@ -278,7 +278,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t, float x, float dPdx, float dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t, float x, float dPdx, float dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
@@ -286,7 +286,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t, float x, float y, float2 dPdx, float2 dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_2D(i, s, float2(x, y).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
@@ -294,7 +294,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float dPdx, float dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
@@ -302,7 +302,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float2 dPdx, float2 dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_2Da(i, s, float4(x, y, layer, 0.0f).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
@@ -310,7 +310,7 @@ static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_3D(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
@@ -361,7 +361,7 @@ struct __hip_tex2dgather_ret<
};
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t, float x, float y, int comp=0)
static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t, float x, float y, int comp=0)
{
TEXTURE_PARAMETERS_INIT;
switch (comp) {
hipamd/include/hip/amd_detail/texture_functions.h
-11102
Dosyayı Görüntüle
Dosya farkı çok büyük olduğundan ihmal edildi Fark Yükle
hipamd/include/hip/amd_detail/texture_indirect_functions.h
+46 -46
Dosyayı Görüntüle
@@ -64,7 +64,7 @@ struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_load_1Db(i, x);
@@ -74,7 +74,7 @@ static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
{
*ptr = tex1Dfetch<T>(textureObject, x);
}
@@ -82,7 +82,7 @@ static __device__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1D(hipTextureObject_t textureObject, float x)
static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, float x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1D(i, s, x);
@@ -92,7 +92,7 @@ static __device__ T tex1D(hipTextureObject_t textureObject, float x)
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
{
*ptr = tex1D<T>(textureObject, x);
}
@@ -100,7 +100,7 @@ static __device__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
@@ -110,7 +110,7 @@ static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
{
*ptr = tex2D<T>(textureObject, x, y);
}
@@ -118,7 +118,7 @@ static __device__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
@@ -128,7 +128,7 @@ static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, fl
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = tex3D<T>(textureObject, x, y, z);
}
@@ -136,7 +136,7 @@ static __device__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
@@ -146,7 +146,7 @@ static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, layer);
}
@@ -154,7 +154,7 @@ static __device__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, fl
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
@@ -164,7 +164,7 @@ static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, floa
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, y, layer);
}
@@ -172,7 +172,7 @@ static __device__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, fl
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
@@ -182,7 +182,7 @@ static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = texCubemap<T>(textureObject, x, y, z);
}
@@ -190,7 +190,7 @@ static __device__ void texCubemap(T *ptr, hipTextureObject_t textureObject, floa
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
@@ -200,7 +200,7 @@ static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
}
@@ -208,7 +208,7 @@ static __device__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObjec
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
TEXTURE_OBJECT_PARAMETERS_INIT
switch (comp) {
@@ -239,7 +239,7 @@ static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, comp);
}
@@ -247,7 +247,7 @@ static __device__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, flo
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
@@ -257,7 +257,7 @@ static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float le
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
{
*ptr = tex1DLod<T>(textureObject, x, level);
}
@@ -265,7 +265,7 @@ static __device__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
@@ -275,7 +275,7 @@ static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
{
*ptr = tex2DLod<T>(textureObject, x, y, level);
}
@@ -283,7 +283,7 @@ static __device__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
@@ -293,7 +293,7 @@ static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = tex3DLod<T>(textureObject, x, y, z, level);
}
@@ -301,7 +301,7 @@ static __device__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
@@ -311,7 +311,7 @@ static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, i
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
{
*ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
}
@@ -319,7 +319,7 @@ static __device__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
@@ -329,7 +329,7 @@ static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
*ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
}
@@ -337,7 +337,7 @@ static __device__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
@@ -347,7 +347,7 @@ static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, flo
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = texCubemapLod<T>(textureObject, x, y, z, level);
}
@@ -355,7 +355,7 @@ static __device__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, f
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
@@ -367,7 +367,7 @@ static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, fl
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
}
@@ -375,7 +375,7 @@ static __device__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
@@ -385,7 +385,7 @@ static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
*ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
}
@@ -393,7 +393,7 @@ static __device__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureOb
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
@@ -403,7 +403,7 @@ static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float d
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
*ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
}
@@ -411,7 +411,7 @@ static __device__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_2D(i, s, float2(x, y).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
@@ -421,7 +421,7 @@ static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
*ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
}
@@ -429,7 +429,7 @@ static __device__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_3D(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
@@ -439,7 +439,7 @@ static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
}
@@ -447,7 +447,7 @@ static __device__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
@@ -457,7 +457,7 @@ static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
*ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
}
@@ -465,7 +465,7 @@ static __device__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_2Da(i, s, float4(x, y, layer, 0.0f).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
@@ -475,7 +475,7 @@ static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
*ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
}
@@ -483,7 +483,7 @@ static __device__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
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
// TODO missing in device libs.
@@ -495,7 +495,7 @@ static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, flo
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
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<T>(textureObject, x, y, z, layer, dPdx, dPdy);
}
hipamd/src/hip_device_runtime.cpp
+3
Dosyayı Görüntüle
@@ -318,6 +318,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
// hipStreamWaitValue64() and hipStreamWaitValue32() support
*pi = g_devices[device]->devices()[0]->info().aqlBarrierValue_;
break;
case hipDeviceAttributeImageSupport:
*pi = static_cast<int>(g_devices[device]->devices()[0]->info().imageSupport_);
break;
default:
HIP_RETURN(hipErrorInvalidValue);
}
hipamd/src/hip_module.cpp
+6
Dosyayı Görüntüle
@@ -663,6 +663,12 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
if ((texRef == nullptr) || (name == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
/* Get address and size for the global symbol */
if (hipSuccess != PlatformState::instance().getDynTexRef(name, hmod, texRef)) {
hipamd/src/hip_texture.cpp
+244 -33
Dosyayı Görüntüle
@@ -76,6 +76,10 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject,
const hipResourceViewDesc* pResViewDesc) {
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
return hipErrorNotSupported;
}
// Validate input params
if (pTexObject == nullptr || pResDesc == nullptr || pTexDesc == nullptr) {
@@ -313,6 +317,12 @@ hipError_t ihipDestroyTextureObject(hipTextureObject_t texObject) {
if (texObject == nullptr) {
return hipSuccess;
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
return hipErrorNotSupported;
}
const hipResourceType type = texObject->resDesc.resType;
const bool isImageFromBuffer = (type == hipResourceTypeLinear) || (type == hipResourceTypePitch2D);
@@ -341,6 +351,13 @@ hipError_t ihipUnbindTexture(textureReference* texRef) {
break;
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
hip_error = ihipDestroyTextureObject(texRef->textureObject);
if (hip_error != hipSuccess) {
break;
@@ -367,6 +384,12 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
if ((pResDesc == nullptr) || (texObject == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pResDesc = texObject->resDesc;
@@ -380,6 +403,12 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc
if ((pResViewDesc == nullptr) || (texObject == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pResViewDesc = texObject->resViewDesc;
@@ -393,16 +422,26 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
if ((pTexDesc == nullptr) || (texObject == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pTexDesc = texObject->texDesc;
HIP_RETURN(hipSuccess);
}
inline bool ihipGetTextureAlignmentOffset(size_t* offset,
inline hipError_t ihipGetTextureAlignmentOffset(size_t* offset,
const void* devPtr) {
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
return hipErrorNotSupported;
}
const char* alignedDevPtr = amd::alignUp(static_cast<const char*>(devPtr), info.imageBaseAddressAlignment_);
const size_t alignedOffset = alignedDevPtr - static_cast<const char*>(devPtr);
@@ -411,14 +450,14 @@ inline bool ihipGetTextureAlignmentOffset(size_t* offset,
// the offset is guaranteed to be 0 and NULL may be passed as the offset parameter.
if ((alignedOffset != 0) && (offset == nullptr)) {
LogPrintfError("Texture object not aligned with offset %u \n", alignedOffset);
return false;
return hipErrorInvalidValue;
}
if (offset != nullptr) {
*offset = alignedOffset;
}
return true;
return hipSuccess;
}
hipError_t ihipBindTexture(size_t* offset,
@@ -435,21 +474,23 @@ hipError_t ihipBindTexture(size_t* offset,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texref->textureObject);
hipError_t err = ihipDestroyTextureObject(texref->textureObject);
if (err != hipSuccess) {
return err;
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeLinear;
resDesc.res.linear.devPtr = const_cast<void*>(devPtr);
resDesc.res.linear.desc = *desc;
resDesc.res.linear.sizeInBytes = size;
if (ihipGetTextureAlignmentOffset(offset, devPtr)) {
// Align the user ptr to HW requirments.
resDesc.res.linear.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
} else {
return hipErrorInvalidValue;
err = ihipGetTextureAlignmentOffset(offset, devPtr);
if (err != hipSuccess) {
return err;
}
// Align the user ptr to HW requirments.
resDesc.res.linear.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
hipTextureDesc texDesc = hip::getTextureDesc(texref);
return ihipCreateTextureObject(const_cast<hipTextureObject_t*>(&texref->textureObject), &resDesc, &texDesc, nullptr);
@@ -471,7 +512,10 @@ hipError_t ihipBindTexture2D(size_t* offset,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texref->textureObject);
hipError_t err = ihipDestroyTextureObject(texref->textureObject);
if (err != hipSuccess) {
return err;
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypePitch2D;
@@ -480,14 +524,13 @@ hipError_t ihipBindTexture2D(size_t* offset,
resDesc.res.pitch2D.width = width;
resDesc.res.pitch2D.height = height;
resDesc.res.pitch2D.pitchInBytes = pitch;
if (ihipGetTextureAlignmentOffset(offset, devPtr)) {
// Align the user ptr to HW requirments.
resDesc.res.pitch2D.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
} else {
return hipErrorInvalidValue;
err = ihipGetTextureAlignmentOffset(offset, devPtr);
if (err != hipSuccess) {
return err;
}
// Align the user ptr to HW requirments.
resDesc.res.pitch2D.devPtr = static_cast<char*>(const_cast<void*>(devPtr)) - *offset;
hipTextureDesc texDesc = hip::getTextureDesc(texref);
return ihipCreateTextureObject(const_cast<hipTextureObject_t*>(&texref->textureObject), &resDesc, &texDesc, nullptr);
@@ -530,7 +573,10 @@ hipError_t ihipBindTextureToArray(const textureReference* texref,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texref->textureObject);
hipError_t err = ihipDestroyTextureObject(texref->textureObject);
if (err != hipSuccess) {
return err;
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeArray;
@@ -576,7 +622,10 @@ hipError_t ihipBindTextureToMipmappedArray(const textureReference* texref,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texref->textureObject);
hipError_t err = ihipDestroyTextureObject(texref->textureObject);
if (err != hipSuccess) {
return err;
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeMipmappedArray;
@@ -645,6 +694,12 @@ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc,
if ((desc == nullptr) || (array == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
// It is UB to call hipGetChannelDesc() on an array created via hipArrayCreate()/hipArray3DCreate().
// This is due to hip not differentiating between runtime and driver types.
@@ -660,6 +715,12 @@ hipError_t hipGetTextureAlignmentOffset(size_t* offset,
if ((offset == nullptr) || (texref == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
// TODO enforce alignment on devPtr.
*offset = 0;
@@ -673,6 +734,13 @@ hipError_t hipGetTextureReference(const textureReference** texref, const void* s
if (texref == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*texref = reinterpret_cast<const textureReference *>(symbol);
HIP_RETURN(hipSuccess);
@@ -686,6 +754,12 @@ hipError_t hipTexRefSetFormat(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->format = fmt;
texRef->numChannels = NumPackedComponents;
@@ -700,6 +774,12 @@ hipError_t hipTexRefSetFlags(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->readMode = hipReadModeNormalizedFloat;
texRef->normalized = 0;
@@ -727,6 +807,12 @@ hipError_t hipTexRefSetFilterMode(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->filterMode = fm;
@@ -742,6 +828,12 @@ hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* pam,
if ((pam == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
// Currently, the only valid value for dim are 0 and 1.
if ((dim != 0) && (dim != 1)) {
@@ -773,6 +865,12 @@ hipError_t hipTexRefSetAddressMode(textureReference* texRef,
dim);
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->addressMode[dim] = am;
@@ -831,7 +929,10 @@ hipError_t hipTexRefSetArray(textureReference* texRef,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texRef->textureObject);
hipError_t err = ihipDestroyTextureObject(texRef->textureObject);
if (err != hipSuccess) {
HIP_RETURN(err);
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeArray;
@@ -842,7 +943,7 @@ hipError_t hipTexRefSetArray(textureReference* texRef,
hipResourceViewFormat format = hip::getResourceViewFormat(hip::getChannelFormatDesc(texRef->numChannels, texRef->format));
hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(array, format);
hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc);
err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc);
if (err != hipSuccess) {
HIP_RETURN(err);
}
@@ -906,24 +1007,26 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texRef->textureObject);
hipError_t err = ihipDestroyTextureObject(texRef->textureObject);
if (err != hipSuccess) {
HIP_RETURN(err);
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeLinear;
resDesc.res.linear.devPtr = dptr;
resDesc.res.linear.desc = hip::getChannelFormatDesc(texRef->numChannels, texRef->format);
resDesc.res.linear.sizeInBytes = bytes;
if (ihipGetTextureAlignmentOffset(ByteOffset, dptr)) {
// Align the user ptr to HW requirments.
resDesc.res.linear.devPtr = static_cast<char*>(dptr) - *ByteOffset;
} else {
HIP_RETURN(hipErrorInvalidValue);
err = ihipGetTextureAlignmentOffset(ByteOffset, dptr);
if (err != hipSuccess) {
HIP_RETURN(err);
}
// Align the user ptr to HW requirments.
resDesc.res.linear.devPtr = static_cast<char*>(dptr) - *ByteOffset;
hipTextureDesc texDesc = hip::getTextureDesc(texRef);
hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr);
err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr);
if (err != hipSuccess) {
HIP_RETURN(err);
}
@@ -950,7 +1053,10 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texRef->textureObject);
hipError_t err = ihipDestroyTextureObject(texRef->textureObject);
if (err != hipSuccess) {
HIP_RETURN(err);
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypePitch2D;
@@ -962,7 +1068,7 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef,
hipTextureDesc texDesc = hip::getTextureDesc(texRef);
hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr);
err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, nullptr);
if (err != hipSuccess) {
HIP_RETURN(err);
}
@@ -983,6 +1089,12 @@ hipError_t hipTexRefGetBorderColor(float* pBorderColor,
if ((pBorderColor == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
// TODO add textureReference::borderColor.
assert(false && "textureReference::borderColor is missing in header");
@@ -999,6 +1111,12 @@ hipError_t hipTexRefGetFilterMode(hipTextureFilterMode* pfm,
if ((pfm == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pfm = texRef->filterMode;
@@ -1013,6 +1131,12 @@ hipError_t hipTexRefGetFlags(unsigned int* pFlags,
if ((pFlags == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pFlags = 0;
@@ -1041,6 +1165,12 @@ hipError_t hipTexRefGetFormat(hipArray_Format* pFormat,
(texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pFormat = texRef->format;
*pNumChannels = texRef->numChannels;
@@ -1056,6 +1186,12 @@ hipError_t hipTexRefGetMaxAnisotropy(int* pmaxAnsio,
if ((pmaxAnsio == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pmaxAnsio = texRef->maxAnisotropy;
@@ -1070,6 +1206,12 @@ hipError_t hipTexRefGetMipmapFilterMode(hipTextureFilterMode* pfm,
if ((pfm == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pfm = texRef->mipmapFilterMode;
@@ -1084,6 +1226,12 @@ hipError_t hipTexRefGetMipmapLevelBias(float* pbias,
if ((pbias == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pbias = texRef->mipmapLevelBias;
@@ -1100,6 +1248,12 @@ hipError_t hipTexRefGetMipmapLevelClamp(float* pminMipmapLevelClamp,
(texRef == nullptr)){
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pminMipmapLevelClamp = texRef->minMipmapLevelClamp;
*pmaxMipmapLevelClamp = texRef->maxMipmapLevelClamp;
@@ -1115,6 +1269,12 @@ hipError_t hipTexRefGetMipmappedArray(hipMipmappedArray_t* pArray,
if ((pArray == nullptr) || (texRef == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
hipResourceDesc resDesc = {};
// TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace.
@@ -1144,6 +1304,12 @@ hipError_t hipTexRefSetBorderColor(textureReference* texRef,
if ((texRef == nullptr) || (pBorderColor == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
// TODO add textureReference::borderColor.
assert(false && "textureReference::borderColor is missing in header");
@@ -1159,6 +1325,12 @@ hipError_t hipTexRefSetMaxAnisotropy(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->maxAnisotropy = maxAniso;
@@ -1172,6 +1344,12 @@ hipError_t hipTexRefSetMipmapFilterMode(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->mipmapFilterMode = fm;
@@ -1185,6 +1363,12 @@ hipError_t hipTexRefSetMipmapLevelBias(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->mipmapLevelBias = bias;
@@ -1199,6 +1383,12 @@ hipError_t hipTexRefSetMipmapLevelClamp(textureReference* texRef,
if (texRef == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
texRef->minMipmapLevelClamp = minMipMapLevelClamp;
texRef->maxMipmapLevelClamp = maxMipMapLevelClamp;
@@ -1227,7 +1417,10 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef,
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
// Any memory previously bound to hTexRef is unbound.
// No need to check for errors.
ihipDestroyTextureObject(texRef->textureObject);
hipError_t err = ihipDestroyTextureObject(texRef->textureObject);
if (err != hipSuccess) {
HIP_RETURN(err);
}
hipResourceDesc resDesc = {};
resDesc.resType = hipResourceTypeMipmappedArray;
@@ -1238,7 +1431,7 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef,
hipResourceViewFormat format = hip::getResourceViewFormat(hip::getChannelFormatDesc(texRef->numChannels, texRef->format));
hipResourceViewDesc resViewDesc = hip::getResourceViewDesc(mipmappedArray, format);
hipError_t err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc);
err = ihipCreateTextureObject(&texRef->textureObject, &resDesc, &texDesc, &resViewDesc);
if (err != hipSuccess) {
HIP_RETURN(err);
}
@@ -1281,6 +1474,12 @@ hipError_t hipTexObjectGetResourceDesc(HIP_RESOURCE_DESC* pResDesc,
if ((pResDesc == nullptr) || (texObject == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pResDesc = hip::getResourceDesc(texObject->resDesc);
@@ -1294,6 +1493,12 @@ hipError_t hipTexObjectGetResourceViewDesc(HIP_RESOURCE_VIEW_DESC* pResViewDesc,
if ((pResViewDesc == nullptr) || (texObject == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pResViewDesc = hip::getResourceViewDesc(texObject->resViewDesc);
@@ -1307,6 +1512,12 @@ hipError_t hipTexObjectGetTextureDesc(HIP_TEXTURE_DESC* pTexDesc,
if ((pTexDesc == nullptr) || (texObject == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
const device::Info& info = device->info();
if (!info.imageSupport_) {
LogPrintfError("Texture not supported on the device %s", info.name_);
HIP_RETURN(hipErrorNotSupported);
}
*pTexDesc = hip::getTextureDesc(texObject->texDesc);