25#if defined(__cplusplus)
27#include <hip/hip_vector_types.h>
28#include <hip/hip_texture_types.h>
29#include <hip/amd_detail/texture_fetch_functions.h>
30#include <hip/amd_detail/ockl_image.h>
32#if !defined(__HIPCC_RTC__)
36#define TEXTURE_OBJECT_PARAMETERS_INIT \
37 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
38 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
42 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
43static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
int x)
45 TEXTURE_OBJECT_PARAMETERS_INIT
46 auto tmp = __ockl_image_load_1Db(i, x);
47 return __hipMapFrom<T>(tmp);
52 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
53static __device__ __hip_img_chk__
void tex1Dfetch(T *ptr, hipTextureObject_t textureObject,
int x)
55 *ptr = tex1Dfetch<T>(textureObject, x);
60 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
61static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject,
float x)
63 TEXTURE_OBJECT_PARAMETERS_INIT
64 auto tmp = __ockl_image_sample_1D(i, s, x);
65 return __hipMapFrom<T>(tmp);
70 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
71static __device__ __hip_img_chk__
void tex1D(T *ptr, hipTextureObject_t textureObject,
float x)
73 *ptr = tex1D<T>(textureObject, x);
78 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
79static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject,
float x,
float y)
81 TEXTURE_OBJECT_PARAMETERS_INIT
82 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
83 return __hipMapFrom<T>(tmp);
88 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
89static __device__ __hip_img_chk__
void tex2D(T *ptr, hipTextureObject_t textureObject,
float x,
float y)
91 *ptr = tex2D<T>(textureObject, x, y);
96 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
97static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject,
float x,
float y,
float z)
99 TEXTURE_OBJECT_PARAMETERS_INIT
100 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
101 return __hipMapFrom<T>(tmp);
106 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
107static __device__ __hip_img_chk__
void tex3D(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
109 *ptr = tex3D<T>(textureObject, x, y, z);
114 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
115static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject,
float x,
int layer)
117 TEXTURE_OBJECT_PARAMETERS_INIT
118 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
119 return __hipMapFrom<T>(tmp);
124 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
125static __device__ __hip_img_chk__
void tex1DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
int layer)
127 *ptr = tex1DLayered<T>(textureObject, x, layer);
132 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
133static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject,
float x,
float y,
int layer)
135 TEXTURE_OBJECT_PARAMETERS_INIT
136 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
137 return __hipMapFrom<T>(tmp);
142 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
143static __device__ __hip_img_chk__
void tex2DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer)
145 *ptr = tex1DLayered<T>(textureObject, x, y, layer);
150 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
151static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject,
float x,
float y,
float z)
153 TEXTURE_OBJECT_PARAMETERS_INIT
154 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
155 return __hipMapFrom<T>(tmp);
160 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
161static __device__ __hip_img_chk__
void texCubemap(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
163 *ptr = texCubemap<T>(textureObject, x, y, z);
168 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
169static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
171 TEXTURE_OBJECT_PARAMETERS_INIT
172 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
173 return __hipMapFrom<T>(tmp);
178 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
179static __device__ __hip_img_chk__
void texCubemapLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
181 *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
186 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
187static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
189 TEXTURE_OBJECT_PARAMETERS_INIT
192 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
193 return __hipMapFrom<T>(tmp);
197 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
198 return __hipMapFrom<T>(tmp);
202 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
203 return __hipMapFrom<T>(tmp);
207 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
208 return __hipMapFrom<T>(tmp);
217 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
218static __device__ __hip_img_chk__
void tex2Dgather(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
220 *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
225 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
226static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject,
float x,
float level)
228 TEXTURE_OBJECT_PARAMETERS_INIT
229 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
230 return __hipMapFrom<T>(tmp);
235 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
236static __device__ __hip_img_chk__
void tex1DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float level)
238 *ptr = tex1DLod<T>(textureObject, x, level);
243 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
244static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject,
float x,
float y,
float level)
246 TEXTURE_OBJECT_PARAMETERS_INIT
247 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
248 return __hipMapFrom<T>(tmp);
253 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
254static __device__ __hip_img_chk__
void tex2DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float level)
256 *ptr = tex2DLod<T>(textureObject, x, y, level);
261 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
262static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
264 TEXTURE_OBJECT_PARAMETERS_INIT
265 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
266 return __hipMapFrom<T>(tmp);
271 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
272static __device__ __hip_img_chk__
void tex3DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
274 *ptr = tex3DLod<T>(textureObject, x, y, z, level);
279 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
280static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject,
float x,
int layer,
float level)
282 TEXTURE_OBJECT_PARAMETERS_INIT
283 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
284 return __hipMapFrom<T>(tmp);
289 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
290static __device__ __hip_img_chk__
void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float level)
292 *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
297 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
298static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
300 TEXTURE_OBJECT_PARAMETERS_INIT
301 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
302 return __hipMapFrom<T>(tmp);
307 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
308static __device__ __hip_img_chk__
void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
310 *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
315 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
316static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
318 TEXTURE_OBJECT_PARAMETERS_INIT
319 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
320 return __hipMapFrom<T>(tmp);
325 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
326static __device__ __hip_img_chk__
void texCubemapLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
328 *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
333 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
334static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
336 TEXTURE_OBJECT_PARAMETERS_INIT
345 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
346static __device__ __hip_img_chk__
void texCubemapGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
348 *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
353 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
354static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
356 TEXTURE_OBJECT_PARAMETERS_INIT
357 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
358 return __hipMapFrom<T>(tmp);
363 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
364static __device__ __hip_img_chk__
void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
366 *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
371 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
372static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
374 TEXTURE_OBJECT_PARAMETERS_INIT
375 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
376 return __hipMapFrom<T>(tmp);
381 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
382static __device__ __hip_img_chk__
void tex1DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
384 *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
389 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
390static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
392 TEXTURE_OBJECT_PARAMETERS_INIT
393 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);
394 return __hipMapFrom<T>(tmp);
399 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
400static __device__ __hip_img_chk__
void tex2DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
402 *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
407 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
408static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
410 TEXTURE_OBJECT_PARAMETERS_INIT
411 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);
412 return __hipMapFrom<T>(tmp);
417 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
418static __device__ __hip_img_chk__
void tex3DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
420 *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
425 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
426static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
428 TEXTURE_OBJECT_PARAMETERS_INIT
429 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
430 return __hipMapFrom<T>(tmp);
435 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
436static __device__ __hip_img_chk__
void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
438 *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
443 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
444static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
446 TEXTURE_OBJECT_PARAMETERS_INIT
447 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);
448 return __hipMapFrom<T>(tmp);
453 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
454static __device__ __hip_img_chk__
void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
456 *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
461 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
462static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
464 TEXTURE_OBJECT_PARAMETERS_INIT
473 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
474static __device__ __hip_img_chk__
void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
476 *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:1993