AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

SWDEV-396574 – Refactor device mapping functions

Explorar el Código 
Refactor mapFrom(), mapTo(), mapElem(),
__hip_is_surf_channel_type() and
__hip_is_itex_channel_type()

Change-Id: I1692b92d407bad742d562678f218ebf8ca532e91


[ROCm/clr commit: 2f2d02649f]
Este commit está contenido en:
taosang2
2023-04-25 17:58:58 -04:00
cometido por Tao Sang
padre ebcab6e910
commit 3a37f33e4e
Se han modificado 4 ficheros con 203 adiciones y 346 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
projects/clr/hipamd/include/hip/amd_detail/amd_hip_vector_types.h
+22 -217
Ver fichero
@@ -1057,6 +1057,28 @@ template <typename __T> struct is_scalar : public integral_constant<bool, __is_s
return HIP_vector_type<T, n>{x} <<= y;
}
/*
* Map HIP_vector_type<U, rankU> to HIP_vector_type<T, rankT>
*/
template<
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
__forceinline__
__HOST_DEVICE__
typename std::enable_if<
(rankT >= 1 && rankT <= 4 && rankU >= 1 && rankU <= 4),
const HIP_vector_type<T, rankT>>::type
__hipMapVector(const HIP_vector_type<U, rankU> &u) {
HIP_vector_type<T, rankT> t; // Initialized to 0
if constexpr (rankT >= 1 && rankU >= 1) t.x = static_cast<T>(u.x);
if constexpr (rankT >= 2 && rankU >= 2) t.y = static_cast<T>(u.y);
if constexpr (rankT >= 3 && rankU >= 3) t.z = static_cast<T>(u.z);
if constexpr (rankT >= 4 && rankU >= 4) t.w = static_cast<T>(u.w);
return t;
};
#define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
using CUDA_name##1 = HIP_vector_type<T, 1>;\
using CUDA_name##2 = HIP_vector_type<T, 2>;\
@@ -2030,223 +2052,6 @@ typedef union {
type r{x, y, z, w}; \
return r; \
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 1 &&
sizeof(U) / sizeof(typename U::value_type) >= 1, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 2 &&
sizeof(U) / sizeof(typename U::value_type) >= 2, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x),
static_cast<typename T::value_type>(u.y)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 3 &&
sizeof(U) / sizeof(typename U::value_type) >= 3, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x),
static_cast<typename T::value_type>(u.y),
static_cast<typename T::value_type>(u.z)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 4 &&
sizeof(U) / sizeof(typename U::value_type) == 1, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x),
static_cast<typename T::value_type>(0),
static_cast<typename T::value_type>(0),
static_cast<typename T::value_type>(0)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 4 &&
sizeof(U) / sizeof(typename U::value_type) == 2, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x),
static_cast<typename T::value_type>(u.y),
static_cast<typename T::value_type>(0),
static_cast<typename T::value_type>(0)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 4 &&
sizeof(U) / sizeof(typename U::value_type) == 3, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x),
static_cast<typename T::value_type>(u.y),
static_cast<typename T::value_type>(u.z),
static_cast<typename T::value_type>(0)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
sizeof(T) / sizeof(typename T::value_type) == 4 &&
sizeof(U) / sizeof(typename U::value_type) >= 4, T>::type
mapElem(const U &u) {
return {
static_cast<typename T::value_type>(u.x),
static_cast<typename T::value_type>(u.y),
static_cast<typename T::value_type>(u.z),
static_cast<typename T::value_type>(u.w)
};
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value, const T>::type
mapFrom(const U &u) {
union {
U u;
T t;
} d = { u };
return d.t;
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
(sizeof(T) == sizeof(typename T::value_type)) ||
std::is_same<typename T::value_type, int>::value ||
std::is_same<typename T::value_type, unsigned int>::value ||
std::is_same<typename T::value_type, float>::value, const T>::type
mapFrom(const U &u) {
union {
U u;
T t;
} d = { u };
return d.t;
}
template<typename T, typename U>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
(sizeof(T) > sizeof(typename T::value_type)) && (
std::is_same<typename T::value_type, char>::value ||
std::is_same<typename T::value_type, unsigned char>::value ||
std::is_same<typename T::value_type, short>::value ||
std::is_same<typename T::value_type, unsigned short>::value), const T>::type
mapFrom(const U &u) {
union {
U u;
int4 i4;
uint4 u4;
} d = { u };
if (std::is_signed<typename T::value_type>::value) {
return mapElem<T>(d.i4) ;
} else {
return mapElem<T>(d.u4);
}
}
template<typename U, typename T>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value, const U>::type
mapTo(const T &t) {
union {
U u;
T t;
} d = { 0 };
d.t = t;
return d.u;
}
template<typename U, typename T>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
(sizeof(T) == sizeof(typename T::value_type)) ||
std::is_same<typename T::value_type, int>::value ||
std::is_same<typename T::value_type, unsigned int>::value ||
std::is_same<typename T::value_type, float>::value, const U>::type
mapTo(const T &t) {
union {
U u;
T t;
} d = { 0 };
d.t = t;
return d.u;
}
template<typename U, typename T>
__HOST_DEVICE__
__forceinline__
typename std::enable_if<
(sizeof(T) > sizeof(typename T::value_type)) && (
std::is_same<typename T::value_type, char>::value ||
std::is_same<typename T::value_type, unsigned char>::value ||
std::is_same<typename T::value_type, short>::value ||
std::is_same<typename T::value_type, unsigned short>::value), const U>::type
mapTo(const T &t) {
union {
U u;
int4 i4;
uint4 u4;
} d = { 0 };
if (std::is_signed<typename T::value_type>::value) {
d.i4 = mapElem<int4>(t);
} else {
d.u4 = mapElem<uint4>(t);
}
return d.u;
}
#else
#define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
static inline __HOST_DEVICE__ type make_##type(comp x) { \
projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h
+31 -53
Ver fichero
@@ -27,36 +27,14 @@ THE SOFTWARE.
#include <hip/surface_types.h>
#include <hip/hip_vector_types.h>
#include <hip/amd_detail/texture_fetch_functions.h>
#include <hip/amd_detail/ockl_image.h>
#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
template<typename T>
struct __hip_is_isurf_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_isurf_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_isurf_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 3) ||
(rank == 4));
};
template <typename T>
using __hip_is_surf_channel_type = __hip_is_tex_channel_type<T>;
// CUDA is using byte address, need map to pixel address for HIP
static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) {
@@ -115,144 +93,144 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format,
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
int boundaryMode = hipBoundaryModeZero) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __ockl_image_load_1D(i, x);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_1D(i, x, tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_2D(i, int2(x, y).data, tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __ockl_image_load_lod_1D(i, x, layer);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_lod_1D(i, x, layer, tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_CM(i, int2(x, y).data, face, tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
*data = mapFrom<T>(tmp);
*data = __hipMapFrom<T>(tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = mapTo<float4::Native_vec_>(data);
auto tmp = __hipMapTo<float4::Native_vec_>(data);
__ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
}
projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
+121 -26
Ver fichero
@@ -92,6 +92,100 @@ struct __hip_tex_ret
static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
};
/*
* Map from device function return U to scalar texture type T
*/
template<typename T, typename U>
__forceinline__ __device__
typename std::enable_if<
__hip_is_tex_channel_type<T>::value && std::is_scalar<T>::value, const T>::type
__hipMapFrom(const U &u) {
if constexpr (sizeof(T) < sizeof(float)) {
union {
U u;
int i;
} d = { u };
return static_cast<T>(d.i);
} else { // sizeof(T) == sizeof(float)
union {
U u;
T t;
} d = { u };
return d.t;
}
}
/*
* Map from device function return U to vector texture type T
*/
template<typename T, typename U>
__forceinline__ __device__
typename std::enable_if<__hip_is_tex_channel_type<typename T::value_type>::value, const T>::type
__hipMapFrom(const U &u) {
if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
union {
U u;
int4 i4;
} d = { u };
return __hipMapVector<typename T::value_type, sizeof(T)/sizeof(typename T::value_type)>(d.i4);
} else { // sizeof(typename T::value_type) == sizeof(float)
union {
U u;
T t;
} d = { u };
return d.t;
}
}
/*
* Map from scalar texture type T to device function input U
*/
template<typename U, typename T>
__forceinline__ __device__
typename std::enable_if<
__hip_is_tex_channel_type<T>::value && std::is_scalar<T>::value, const U>::type
__hipMapTo(const T &t) {
if constexpr (sizeof(T) < sizeof(float)) {
union {
U u;
int i;
} d = { 0 };
d.i = static_cast<int>(t);
return d.u;
} else { // sizeof(T) == sizeof(float)
union {
U u;
T t;
} d = { 0 };
d.t = t;
return d.u;
}
}
/*
* Map from vector texture type T to device function input U
*/
template<typename U, typename T>
__forceinline__ __device__
typename std::enable_if<__hip_is_tex_channel_type<typename T::value_type>::value, const U>::type
__hipMapTo(const T &t) {
if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
union {
U u;
int4 i4;
} d = { 0 };
d.i4 = __hipMapVector<int, 4>(t);
return d.u;
} else { // sizeof(typename T::value_type) == sizeof(float)
union {
U u;
T t;
} d = { 0 };
d.t = t;
return d.u;
}
}
template <
typename T,
hipTextureReadMode readMode>
@@ -137,12 +231,13 @@ struct __hip_tex_ret<
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
};
template <typename T, hipTextureReadMode readMode>
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);
return mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -150,7 +245,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -158,7 +253,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -166,7 +261,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -174,7 +269,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -182,7 +277,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -190,7 +285,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -198,7 +293,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -206,7 +301,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -214,7 +309,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -222,7 +317,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -230,7 +325,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -238,7 +333,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -246,7 +341,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -254,7 +349,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -263,7 +358,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
// return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return {};
}
@@ -273,7 +368,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
// return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return {};
}
@@ -282,7 +377,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -290,7 +385,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -298,7 +393,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -306,7 +401,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <typename T, hipTextureReadMode readMode>
@@ -314,7 +409,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 mapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
}
template <
@@ -367,19 +462,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 mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
case 2: {
auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
case 3: {
auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
default: {
auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
return mapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
}
}
return {};
projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
+29 -50
Ver fichero
@@ -26,6 +26,7 @@ THE SOFTWARE.
#include <hip/hip_vector_types.h>
#include <hip/hip_texture_types.h>
#include <hip/amd_detail/texture_fetch_functions.h>
#include <hip/amd_detail/ockl_image.h>
#if !defined(__HIPCC_RTC__)
@@ -36,30 +37,8 @@ THE SOFTWARE.
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
template<typename T>
struct __hip_is_itex_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_itex_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
};
template <typename T>
using __hip_is_itex_channel_type = __hip_is_tex_channel_type<T>;
template <
typename T,
@@ -68,7 +47,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_load_1Db(i, x);
return mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -86,7 +65,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -104,7 +83,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -122,7 +101,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -140,7 +119,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -158,7 +137,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -176,7 +155,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -194,7 +173,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -214,22 +193,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
break;
}
case 2: {
auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
return mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
break;
}
case 3: {
auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
return mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
break;
}
default: {
auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
return mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
break;
}
};
@@ -251,7 +230,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -269,7 +248,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -287,7 +266,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -305,7 +284,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -323,7 +302,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -341,7 +320,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -360,7 +339,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 mapFrom<T>(tmp);
// return __hipMapFrom<T>(tmp);
return {};
}
@@ -379,7 +358,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -397,7 +376,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -415,7 +394,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -433,7 +412,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -451,7 +430,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -469,7 +448,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 mapFrom<T>(tmp);
return __hipMapFrom<T>(tmp);
}
template <
@@ -488,7 +467,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 mapFrom<T>(tmp);
// return __hipMapFrom<T>(tmp);
return {};
}