AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-396574 – Refactor texture channel type checking

Просмотр исходного кода 
Create __hip_is_tex_surf_channel_type to replace
__hip_is_tex_channel_type
__hip_is_itex_channel_type
__hip_is_surf_channel_type .

Change-Id: I1692b92d417bad742d562679f218ebf8ca532e31
Этот коммит содержится в:
taosang2
2023-05-01 16:27:03 -04:00
коммит произвёл Tao Sang
родитель 31b362bf6e
Коммит 381e6520d0
3 изменённых файлов: 80 добавлений и 77 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
hipamd/include/hip/amd_detail/amd_surface_functions.h
+14 -17
Просмотреть файл
@@ -33,9 +33,6 @@ THE SOFTWARE.
#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
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) {
/*
@@ -93,7 +90,7 @@ static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format,
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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
@@ -104,7 +101,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -114,7 +111,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -124,7 +121,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -134,7 +131,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -144,7 +141,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -154,7 +151,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -164,7 +161,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -174,7 +171,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -184,7 +181,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -194,7 +191,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -204,7 +201,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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));
@@ -214,7 +211,7 @@ static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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
@@ -225,7 +222,7 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfac
template <
typename T,
typename std::enable_if<__hip_is_surf_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_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
hipamd/include/hip/amd_detail/texture_fetch_functions.h
+20 -11
Просмотреть файл
@@ -37,7 +37,7 @@ THE SOFTWARE.
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
template<typename T>
struct __hip_is_tex_channel_type
struct __hip_is_tex_surf_scalar_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
@@ -49,13 +49,20 @@ struct __hip_is_tex_channel_type
std::is_same<T, float>::value;
};
template<typename T>
struct __hip_is_tex_surf_channel_type
{
static constexpr bool value =
__hip_is_tex_surf_scalar_channel_type<T>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
struct __hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_tex_channel_type<T>::value &&
__hip_is_tex_surf_scalar_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
@@ -98,7 +105,7 @@ struct __hip_tex_ret
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
__hip_is_tex_surf_scalar_channel_type<T>::value, const T>::type
__hipMapFrom(const U &u) {
if constexpr (sizeof(T) < sizeof(float)) {
union {
@@ -120,7 +127,8 @@ __hipMapFrom(const U &u) {
*/
template<typename T, typename U>
__forceinline__ __device__
typename std::enable_if<__hip_is_tex_channel_type<typename T::value_type>::value, const T>::type
typename std::enable_if<
__hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value, const T>::type
__hipMapFrom(const U &u) {
if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
union {
@@ -143,7 +151,7 @@ __hipMapFrom(const U &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
__hip_is_tex_surf_scalar_channel_type<T>::value, const U>::type
__hipMapTo(const T &t) {
if constexpr (sizeof(T) < sizeof(float)) {
union {
@@ -167,7 +175,8 @@ __hipMapTo(const T &t) {
*/
template<typename U, typename T>
__forceinline__ __device__
typename std::enable_if<__hip_is_tex_channel_type<typename T::value_type>::value, const U>::type
typename std::enable_if<
__hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value, const U>::type
__hipMapTo(const T &t) {
if constexpr (sizeof(typename T::value_type) < sizeof(float)) {
union {
@@ -195,7 +204,7 @@ template <typename T>
struct __hip_tex_ret<
T,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
{
using type = T;
};
@@ -206,7 +215,7 @@ template<
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
};
@@ -430,7 +439,7 @@ template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
@@ -441,7 +450,7 @@ template<
struct __hip_tex2dgather_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
hipamd/include/hip/amd_detail/texture_indirect_functions.h
+46 -49
Просмотреть файл
@@ -37,12 +37,9 @@ 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>
using __hip_is_itex_channel_type = __hip_is_tex_channel_type<T>;
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -52,7 +49,7 @@ static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
{
*ptr = tex1Dfetch<T>(textureObject, x);
@@ -60,7 +57,7 @@ static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t tex
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, float x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -70,7 +67,7 @@ static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, floa
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
{
*ptr = tex1D<T>(textureObject, x);
@@ -78,7 +75,7 @@ static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureO
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -88,7 +85,7 @@ static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, floa
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
{
*ptr = tex2D<T>(textureObject, x, y);
@@ -96,7 +93,7 @@ static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureO
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -106,7 +103,7 @@ static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, floa
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -114,7 +111,7 @@ static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureO
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -124,7 +121,7 @@ static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObjec
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, layer);
@@ -132,7 +129,7 @@ static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t t
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -142,7 +139,7 @@ static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObjec
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -150,7 +147,7 @@ static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t t
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -160,7 +157,7 @@ static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -168,7 +165,7 @@ static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t tex
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -178,7 +175,7 @@ static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t texture
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -186,7 +183,7 @@ static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObjec
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -217,7 +214,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -225,7 +222,7 @@ static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t te
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -235,7 +232,7 @@ static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, f
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
{
*ptr = tex1DLod<T>(textureObject, x, level);
@@ -243,7 +240,7 @@ static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textu
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -253,7 +250,7 @@ static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, f
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -261,7 +258,7 @@ static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textu
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -271,7 +268,7 @@ static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, f
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -279,7 +276,7 @@ static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textu
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -289,7 +286,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureOb
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -297,7 +294,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -307,7 +304,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureO
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -315,7 +312,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -325,7 +322,7 @@ static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObje
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -333,7 +330,7 @@ static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -345,7 +342,7 @@ static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObj
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -353,7 +350,7 @@ static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -363,7 +360,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t text
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -371,7 +368,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureOb
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -381,7 +378,7 @@ static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -389,7 +386,7 @@ static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t text
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -399,7 +396,7 @@ static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -407,7 +404,7 @@ static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t text
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -417,7 +414,7 @@ static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject,
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -425,7 +422,7 @@ static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t text
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -435,7 +432,7 @@ static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureO
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -443,7 +440,7 @@ static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
@@ -453,7 +450,7 @@ static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureO
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);
@@ -461,7 +458,7 @@ static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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
@@ -473,7 +470,7 @@ static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t te
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
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);