AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

SWDEV-315590 - Fix texture vector issue

Przeglądaj źródła 
Apply mapFrom() to related texture device functions.

Change-Id: Ifed918c87373b784780b2180320ccb746845daff


[ROCm/clr commit: 78de8cd672]
This commit is contained in:
Tao Sang
2022-01-04 19:26:14 -05:00
rodzic 77bb6fd205
commit 8929cc4544
2 zmienionych plików z 50 dodań i 50 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
+25 -25
Wyświetl plik
@@ -150,7 +150,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_1D(i, s, x);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -158,7 +158,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -166,7 +166,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -174,7 +174,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -182,7 +182,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -190,7 +190,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -198,7 +198,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -206,7 +206,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -214,7 +214,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_1Da(i, s, float2(x, layer).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -222,7 +222,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_2Da(i, s, float4(x, y, layer, 0.0f).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -230,7 +230,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -238,7 +238,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -246,7 +246,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -254,7 +254,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -263,7 +263,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
TEXTURE_PARAMETERS_INIT;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(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);
// return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
// return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return {};
}
@@ -273,7 +273,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
TEXTURE_PARAMETERS_INIT;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
// return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
// return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return {};
}
@@ -282,7 +282,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -290,7 +290,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
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);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -298,7 +298,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -306,7 +306,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
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);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -314,7 +314,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> t
{
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);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <
@@ -367,19 +367,19 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, rea
switch (comp) {
case 1: {
auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
case 2: {
auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
case 3: {
auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
default: {
auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
}
return {};
projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
+25 -25
Wyświetl plik
@@ -86,7 +86,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1D(i, s, x);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -104,7 +104,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -122,7 +122,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -140,7 +140,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -158,7 +158,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -176,7 +176,7 @@ static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -194,7 +194,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -214,22 +214,22 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject
switch (comp) {
case 1: {
auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
break;
}
case 2: {
auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
break;
}
case 3: {
auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
break;
}
default: {
auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
break;
}
};
@@ -251,7 +251,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -269,7 +269,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -287,7 +287,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -305,7 +305,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -323,7 +323,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureO
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -341,7 +341,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -360,7 +360,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(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);
// return *reinterpret_cast<T*>(&tmp);
// return mapFrom<T>(tmp);
return {};
}
@@ -379,7 +379,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -397,7 +397,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject,
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -415,7 +415,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject,
{
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);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -433,7 +433,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject,
{
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);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -451,7 +451,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -469,7 +469,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO
{
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);
return *reinterpret_cast<T*>(&tmp);
return mapFrom<T>(tmp);
}
template <
@@ -488,7 +488,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t te
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
// return *reinterpret_cast<T*>(&tmp);
// return mapFrom<T>(tmp);
return {};
}