25#if defined(__cplusplus)
27#include <hip/hip_vector_types.h>
28#include <hip/hip_texture_types.h>
29#include <hip/amd_detail/ockl_image.h>
31#if !defined(__HIPCC_RTC__)
35#define TEXTURE_PARAMETERS_INIT \
36 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
37 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
40struct __hip_is_tex_surf_scalar_channel_type
42 static constexpr bool value =
43 std::is_same<T, char>::value ||
44 std::is_same<T, unsigned char>::value ||
45 std::is_same<T, short>::value ||
46 std::is_same<T, unsigned short>::value ||
47 std::is_same<T, int>::value ||
48 std::is_same<T, unsigned int>::value ||
49 std::is_same<T, float>::value;
53struct __hip_is_tex_surf_channel_type
55 static constexpr bool value =
56 __hip_is_tex_surf_scalar_channel_type<T>::value;
62struct __hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>
64 static constexpr bool value =
65 __hip_is_tex_surf_scalar_channel_type<T>::value &&
72struct __hip_is_tex_normalized_channel_type
74 static constexpr bool value =
75 std::is_same<T, char>::value ||
76 std::is_same<T, unsigned char>::value ||
77 std::is_same<T, short>::value ||
78 std::is_same<T, unsigned short>::value;
84struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
86 static constexpr bool value =
87 __hip_is_tex_normalized_channel_type<T>::value &&
95 hipTextureReadMode readMode,
96 typename Enable =
void>
99 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
105template<
typename T,
typename U>
106__forceinline__ __device__
107typename std::enable_if<
108 __hip_is_tex_surf_scalar_channel_type<T>::value,
const T>::type
109__hipMapFrom(
const U &u) {
110 if constexpr (
sizeof(T) <
sizeof(
float)) {
115 return static_cast<T
>(d.i);
128template<
typename T,
typename U>
129__forceinline__ __device__
130typename std::enable_if<
131 __hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value,
const T>::type
132__hipMapFrom(
const U &u) {
133 if constexpr (
sizeof(
typename T::value_type) <
sizeof(
float)) {
138 return __hipMapVector<
typename T::value_type,
sizeof(T)/
sizeof(
typename T::value_type)>(d.i4);
151template<
typename U,
typename T>
152__forceinline__ __device__
153typename std::enable_if<
154__hip_is_tex_surf_scalar_channel_type<T>::value,
const U>::type
155__hipMapTo(
const T &t) {
156 if constexpr (
sizeof(T) <
sizeof(
float)) {
161 d.i =
static_cast<int>(t);
176template<
typename U,
typename T>
177__forceinline__ __device__
178typename std::enable_if<
179 __hip_is_tex_surf_scalar_channel_type<typename T::value_type>::value,
const U>::type
180__hipMapTo(
const T &t) {
181 if constexpr (
sizeof(
typename T::value_type) <
sizeof(
float)) {
186 d.i4 = __hipMapVector<int, 4>(t);
200 hipTextureReadMode readMode>
201using __hip_tex_ret_t =
typename __hip_tex_ret<T, readMode, bool>::type;
206 hipReadModeElementType,
207 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
216 HIP_vector_type<T, rank>,
217 hipReadModeElementType,
218 typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
220 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
226 hipReadModeNormalizedFloat,
227 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
236 HIP_vector_type<T, rank>,
237 hipReadModeNormalizedFloat,
238 typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
240 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
244template <
typename T, hipTextureReadMode readMode>
245static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t,
int x)
247 TEXTURE_PARAMETERS_INIT;
248 auto tmp = __ockl_image_load_1Db(i, x);
249 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
252template <
typename T, hipTextureReadMode readMode>
253static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t,
float x)
255 TEXTURE_PARAMETERS_INIT;
256 auto tmp = __ockl_image_sample_1D(i, s, x);
257 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
260template <
typename T, hipTextureReadMode readMode>
261static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t,
float x,
float y)
263 TEXTURE_PARAMETERS_INIT;
264 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
265 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
268template <
typename T, hipTextureReadMode readMode>
269static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer)
271 TEXTURE_PARAMETERS_INIT;
272 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
273 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
276template <
typename T, hipTextureReadMode readMode>
277static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer)
279 TEXTURE_PARAMETERS_INIT;
280 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
281 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
284template <
typename T, hipTextureReadMode readMode>
285static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z)
287 TEXTURE_PARAMETERS_INIT;
288 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
289 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
292template <
typename T, hipTextureReadMode readMode>
293static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z)
295 TEXTURE_PARAMETERS_INIT;
296 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
297 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
300template <
typename T, hipTextureReadMode readMode>
301static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t,
float x,
float level)
303 TEXTURE_PARAMETERS_INIT;
304 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
305 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
308template <
typename T, hipTextureReadMode readMode>
309static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float level)
311 TEXTURE_PARAMETERS_INIT;
312 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
313 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
316template <
typename T, hipTextureReadMode readMode>
317static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float level)
319 TEXTURE_PARAMETERS_INIT;
320 auto tmp = __ockl_image_sample_lod_1Da(i, s,
float2(x, layer).data, level);
321 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
324template <
typename T, hipTextureReadMode readMode>
325static __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)
327 TEXTURE_PARAMETERS_INIT;
328 auto tmp = __ockl_image_sample_lod_2Da(i, s,
float4(x, y, layer, 0.0f).data, level);
329 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
332template <
typename T, hipTextureReadMode readMode>
333static __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)
335 TEXTURE_PARAMETERS_INIT;
336 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
337 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
340template <
typename T, hipTextureReadMode readMode>
341static __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)
343 TEXTURE_PARAMETERS_INIT;
344 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
345 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
348template <
typename T, hipTextureReadMode readMode>
349static __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)
351 TEXTURE_PARAMETERS_INIT;
352 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
353 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
356template <
typename T, hipTextureReadMode readMode>
357static __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)
359 TEXTURE_PARAMETERS_INIT;
360 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
361 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
364template <
typename T, hipTextureReadMode readMode>
365static __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)
367 TEXTURE_PARAMETERS_INIT;
374template <
typename T, hipTextureReadMode readMode>
375static __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)
377 TEXTURE_PARAMETERS_INIT;
384template <
typename T, hipTextureReadMode readMode>
385static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t,
float x,
float dPdx,
float dPdy)
387 TEXTURE_PARAMETERS_INIT;
388 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
389 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
392template <
typename T, hipTextureReadMode readMode>
393static __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)
395 TEXTURE_PARAMETERS_INIT;
396 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);
397 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
400template <
typename T, hipTextureReadMode readMode>
401static __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)
403 TEXTURE_PARAMETERS_INIT;
404 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
405 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
408template <
typename T, hipTextureReadMode readMode>
409static __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)
411 TEXTURE_PARAMETERS_INIT;
412 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);
413 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
416template <
typename T, hipTextureReadMode readMode>
417static __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)
419 TEXTURE_PARAMETERS_INIT;
420 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);
421 return __hipMapFrom<__hip_tex_ret_t<T, readMode>>(tmp);
426 hipTextureReadMode readMode,
427 typename Enable =
void>
428struct __hip_tex2dgather_ret
430 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
435 hipTextureReadMode readMode>
436using __hip_tex2dgather_ret_t =
typename __hip_tex2dgather_ret<T, readMode, bool>::type;
439struct __hip_tex2dgather_ret<
441 hipReadModeElementType,
442 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
444 using type = HIP_vector_type<T, 4>;
450struct __hip_tex2dgather_ret<
451 HIP_vector_type<T, rank>,
452 hipReadModeElementType,
453 typename std::enable_if<__hip_is_tex_surf_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
455 using type = HIP_vector_type<T, 4>;
459struct __hip_tex2dgather_ret<
461 hipReadModeNormalizedFloat,
462 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
467template <
typename T, hipTextureReadMode readMode>
468static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
int comp=0)
470 TEXTURE_PARAMETERS_INIT;
473 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
474 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
477 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
478 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
481 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
482 return __hipMapFrom<__hip_tex2dgather_ret_t<T, readMode>>(tmp);
485 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
486 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