diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h index f7b7e3f67b..dcf5f2fdf8 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h @@ -150,7 +150,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t t { TEXTURE_PARAMETERS_INIT; auto tmp = __ockl_image_sample_1D(i, s, x); - return *reinterpret_cast<__hip_tex_ret_t*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -158,7 +158,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t t { TEXTURE_PARAMETERS_INIT; auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data); - return *reinterpret_cast<__hip_tex_ret_t*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -166,7 +166,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t t { TEXTURE_PARAMETERS_INIT; auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data); - return *reinterpret_cast<__hip_tex_ret_t*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -174,7 +174,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -182,7 +182,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -190,7 +190,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -198,7 +198,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t t { TEXTURE_PARAMETERS_INIT; auto tmp = __ockl_image_sample_lod_1D(i, s, x, level); - return *reinterpret_cast<__hip_tex_ret_t*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -206,7 +206,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -214,7 +214,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -222,7 +222,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -230,7 +230,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -238,7 +238,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -246,7 +246,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -254,7 +254,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -263,7 +263,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + // return mapFrom<__hip_tex_ret_t>(tmp); return {}; } @@ -273,7 +273,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + // return mapFrom<__hip_tex_ret_t>(tmp); return {}; } @@ -282,7 +282,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t t { TEXTURE_PARAMETERS_INIT; auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy); - return *reinterpret_cast<__hip_tex_ret_t*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -290,7 +290,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -298,7 +298,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -306,7 +306,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template @@ -314,7 +314,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t 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*>(&tmp); + return mapFrom<__hip_tex_ret_t>(tmp); } template < @@ -367,19 +367,19 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t*>(&tmp); + return mapFrom<__hip_tex2dgather_ret_t>(tmp); } case 2: { auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data); - return *reinterpret_cast<__hip_tex2dgather_ret_t*>(&tmp); + return mapFrom<__hip_tex2dgather_ret_t>(tmp); } case 3: { auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data); - return *reinterpret_cast<__hip_tex2dgather_ret_t*>(&tmp); + return mapFrom<__hip_tex2dgather_ret_t>(tmp); } default: { auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data); - return *reinterpret_cast<__hip_tex2dgather_ret_t*>(&tmp); + return mapFrom<__hip_tex2dgather_ret_t>(tmp); } } return {}; diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h index 5bd3f6f9db..0dd04d74dd 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h @@ -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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(tmp); break; } case 2: { auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data); - return *reinterpret_cast(&tmp); + return mapFrom(tmp); break; } case 3: { auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data); - return *reinterpret_cast(&tmp); + return mapFrom(tmp); break; } default: { auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data); - return *reinterpret_cast(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + // return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + return mapFrom(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(&tmp); + // return mapFrom(tmp); return {}; }