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>
31 #include <type_traits>
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;
39 struct __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;
52 struct __hip_is_tex_surf_channel_type
54 static constexpr
bool value =
55 __hip_is_tex_surf_scalar_channel_type<T>::value;
61 struct __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 &&
71 struct __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;
83 struct __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!");
104 template<
typename T,
typename U>
105 __forceinline__ __device__
106 typename 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);
127 template<
typename T,
typename U>
128 __forceinline__ __device__
129 typename 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);
150 template<
typename U,
typename T>
151 __forceinline__ __device__
152 typename 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);
175 template<
typename U,
typename T>
176 __forceinline__ __device__
177 typename 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>
200 using __hip_tex_ret_t =
typename __hip_tex_ret<T, readMode, bool>::type;
202 template <
typename T>
203 struct __hip_tex_ret<
205 hipReadModeElementType,
206 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value, bool>::type>
214 struct __hip_tex_ret<
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>;
223 struct __hip_tex_ret<
225 hipReadModeNormalizedFloat,
226 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
234 struct __hip_tex_ret<
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>;
243 template <
typename T, hipTextureReadMode readMode>
244 static __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);
251 template <
typename T, hipTextureReadMode readMode>
252 static __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);
259 template <
typename T, hipTextureReadMode readMode>
260 static __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);
267 template <
typename T, hipTextureReadMode readMode>
268 static __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);
275 template <
typename T, hipTextureReadMode readMode>
276 static __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);
283 template <
typename T, hipTextureReadMode readMode>
284 static __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);
291 template <
typename T, hipTextureReadMode readMode>
292 static __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);
299 template <
typename T, hipTextureReadMode readMode>
300 static __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);
307 template <
typename T, hipTextureReadMode readMode>
308 static __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);
315 template <
typename T, hipTextureReadMode readMode>
316 static __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);
323 template <
typename T, hipTextureReadMode readMode>
324 static __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);
331 template <
typename T, hipTextureReadMode readMode>
332 static __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);
339 template <
typename T, hipTextureReadMode readMode>
340 static __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);
347 template <
typename T, hipTextureReadMode readMode>
348 static __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);
355 template <
typename T, hipTextureReadMode readMode>
356 static __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);
363 template <
typename T, hipTextureReadMode readMode>
364 static __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;
373 template <
typename T, hipTextureReadMode readMode>
374 static __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;
383 template <
typename T, hipTextureReadMode readMode>
384 static __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);
391 template <
typename T, hipTextureReadMode readMode>
392 static __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);
399 template <
typename T, hipTextureReadMode readMode>
400 static __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);
407 template <
typename T, hipTextureReadMode readMode>
408 static __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);
415 template <
typename T, hipTextureReadMode readMode>
416 static __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>
427 struct __hip_tex2dgather_ret
429 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
434 hipTextureReadMode readMode>
435 using __hip_tex2dgather_ret_t =
typename __hip_tex2dgather_ret<T, readMode, bool>::type;
437 template <
typename T>
438 struct __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>;
449 struct __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>;
457 template <
typename T>
458 struct __hip_tex2dgather_ret<
460 hipReadModeNormalizedFloat,
461 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
466 template <
typename T, hipTextureReadMode readMode>
467 static __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:1820
Definition: amd_hip_vector_types.h:2035
Definition: amd_hip_vector_types.h:2042