25#if defined(__cplusplus)
27#if !defined(__HIPCC_RTC__)
28#include <hip/hip_vector_types.h>
29#include <hip/hip_texture_types.h>
30#include <hip/amd_detail/ockl_image.h>
34#define TEXTURE_PARAMETERS_INIT \
35 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
36 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
39struct __hip_is_tex_surf_scalar_channel_type
41 static constexpr bool value =
42 std::is_same<T, char>::value ||
43 std::is_same<T, unsigned char>::value ||
44 std::is_same<T, short>::value ||
45 std::is_same<T, unsigned short>::value ||
46 std::is_same<T, int>::value ||
47 std::is_same<T, unsigned int>::value ||
48 std::is_same<T, float>::value;
52struct __hip_is_tex_surf_channel_type
54 static constexpr bool value =
55 __hip_is_tex_surf_scalar_channel_type<T>::value;
61struct __hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>
63 static constexpr bool value =
64 __hip_is_tex_surf_scalar_channel_type<T>::value &&
71struct __hip_is_tex_normalized_channel_type
73 static constexpr bool value =
74 std::is_same<T, char>::value ||
75 std::is_same<T, unsigned char>::value ||
76 std::is_same<T, short>::value ||
77 std::is_same<T, unsigned short>::value;
83struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
85 static constexpr bool value =
86 __hip_is_tex_normalized_channel_type<T>::value &&
94 hipTextureReadMode readMode,
95 typename Enable =
void>
98 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
104template<
typename T,
typename U>
105__forceinline__ __device__
106typename std::enable_if<
107 __hip_is_tex_surf_scalar_channel_type<T>::value,
const T>::type
108__hipMapFrom(
const U &u) {
109 if constexpr (
sizeof(T) <
sizeof(
float)) {
114 return static_cast<T
>(d.i);
127template<
typename T,
typename U>
128__forceinline__ __device__
129typename std::enable_if<
130 __hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value,
const T>::type
131__hipMapFrom(
const U &u) {
132 if constexpr (
sizeof(
typename T::value_type) <
sizeof(
float)) {
137 return __hipMapVector<
typename T::value_type,
sizeof(T)/
sizeof(
typename T::value_type)>(d.i4);
150template<
typename U,
typename T>
151__forceinline__ __device__
152typename std::enable_if<
153__hip_is_tex_surf_scalar_channel_type<T>::value,
const U>::type
154__hipMapTo(
const T &t) {
155 if constexpr (
sizeof(T) <
sizeof(
float)) {
160 d.i =
static_cast<int>(t);
175template<
typename U,
typename T>
176__forceinline__ __device__
177typename std::enable_if<
178 __hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value,
const U>::type
179__hipMapTo(
const T &t) {
180 if constexpr (
sizeof(
typename T::value_type) <
sizeof(
float)) {
185 d.i4 = __hipMapVector<int, 4>(t);
199 hipTextureReadMode readMode>
200using __hip_tex_ret_t =
typename __hip_tex_ret<T, readMode, bool>::type;
205 hipReadModeElementType,
206 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
215 HIP_vector_type<T, rank>,
216 hipReadModeElementType,
217 typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
219 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
225 hipReadModeNormalizedFloat,
226 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
235 HIP_vector_type<T, rank>,
236 hipReadModeNormalizedFloat,
237 typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
239 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
243template <
typename T, hipTextureReadMode readMode>
244static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t,
int x)
246 TEXTURE_PARAMETERS_INIT;
247 auto tmp = __ockl_image_load_1Db(i, x);
248 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
251template <
typename T, hipTextureReadMode readMode>
252static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t,
float x)
254 TEXTURE_PARAMETERS_INIT;
255 auto tmp = __ockl_image_sample_1D(i, s, x);
256 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
259template <
typename T, hipTextureReadMode readMode>
260static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t,
float x,
float y)
262 TEXTURE_PARAMETERS_INIT;
263 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
264 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
267template <
typename T, hipTextureReadMode readMode>
268static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer)
270 TEXTURE_PARAMETERS_INIT;
271 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
272 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
275template <
typename T, hipTextureReadMode readMode>
276static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer)
278 TEXTURE_PARAMETERS_INIT;
279 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
280 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
283template <
typename T, hipTextureReadMode readMode>
284static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z)
286 TEXTURE_PARAMETERS_INIT;
287 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
288 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
291template <
typename T, hipTextureReadMode readMode>
292static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z)
294 TEXTURE_PARAMETERS_INIT;
295 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
296 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
299template <
typename T, hipTextureReadMode readMode>
300static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t,
float x,
float level)
302 TEXTURE_PARAMETERS_INIT;
303 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
304 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
307template <
typename T, hipTextureReadMode readMode>
308static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float level)
310 TEXTURE_PARAMETERS_INIT;
311 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
312 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
315template <
typename T, hipTextureReadMode readMode>
316static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float level)
318 TEXTURE_PARAMETERS_INIT;
319 auto tmp = __ockl_image_sample_lod_1Da(i, s,
float2(x, layer).data, level);
320 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
323template <
typename T, hipTextureReadMode readMode>
324static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer,
float level)
326 TEXTURE_PARAMETERS_INIT;
327 auto tmp = __ockl_image_sample_lod_2Da(i, s,
float4(x, y, layer, 0.0f).data, level);
328 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
331template <
typename T, hipTextureReadMode readMode>
332static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z,
float level)
334 TEXTURE_PARAMETERS_INIT;
335 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
336 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
339template <
typename T, hipTextureReadMode readMode>
340static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z,
float level)
342 TEXTURE_PARAMETERS_INIT;
343 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
344 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
347template <
typename T, hipTextureReadMode readMode>
348static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer)
350 TEXTURE_PARAMETERS_INIT;
351 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
352 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
355template <
typename T, hipTextureReadMode readMode>
356static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayeredLod(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer,
float level)
358 TEXTURE_PARAMETERS_INIT;
359 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
360 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
363template <
typename T, hipTextureReadMode readMode>
364static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
366 TEXTURE_PARAMETERS_INIT;
373template <
typename T, hipTextureReadMode readMode>
374static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
376 TEXTURE_PARAMETERS_INIT;
383template <
typename T, hipTextureReadMode readMode>
384static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t,
float x,
float dPdx,
float dPdy)
386 TEXTURE_PARAMETERS_INIT;
387 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
388 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
391template <
typename T, hipTextureReadMode readMode>
392static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float2 dPdx,
float2 dPdy)
394 TEXTURE_PARAMETERS_INIT;
395 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);
396 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
399template <
typename T, hipTextureReadMode readMode>
400static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float dPdx,
float dPdy)
402 TEXTURE_PARAMETERS_INIT;
403 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
404 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
407template <
typename T, hipTextureReadMode readMode>
408static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
410 TEXTURE_PARAMETERS_INIT;
411 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);
412 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
415template <
typename T, hipTextureReadMode readMode>
416static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
418 TEXTURE_PARAMETERS_INIT;
419 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);
420 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
425 hipTextureReadMode readMode,
426 typename Enable =
void>
427struct __hip_tex2dgather_ret
429 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
434 hipTextureReadMode readMode>
435using __hip_tex2dgather_ret_t =
typename __hip_tex2dgather_ret<T, readMode, bool>::type;
438struct __hip_tex2dgather_ret<
440 hipReadModeElementType,
441 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
443 using type = HIP_vector_type<T, 4>;
449struct __hip_tex2dgather_ret<
450 HIP_vector_type<T, rank>,
451 hipReadModeElementType,
452 typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
454 using type = HIP_vector_type<T, 4>;
458struct __hip_tex2dgather_ret<
460 hipReadModeNormalizedFloat,
461 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
466template <
typename T, hipTextureReadMode readMode>
467static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
int comp=0)
469 TEXTURE_PARAMETERS_INIT;
472 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
473 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
476 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
477 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
480 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
481 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
484 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
485 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
Definition amd_hip_vector_types.h:1771
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:1993