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/texture_fetch_functions.h>
31#include <hip/amd_detail/ockl_image.h>
35#define TEXTURE_OBJECT_PARAMETERS_INIT \
36 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
37 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
41 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
42static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject,
int x)
44 TEXTURE_OBJECT_PARAMETERS_INIT
45 auto tmp = __ockl_image_load_1Db(i, x);
46 return __hipMapFrom<T>(tmp);
51 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
52static __device__ __hip_img_chk__
void tex1Dfetch(T *ptr, hipTextureObject_t textureObject,
int x)
54 *ptr = tex1Dfetch<T>(textureObject, x);
59 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
60static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject,
float x)
62 TEXTURE_OBJECT_PARAMETERS_INIT
63 auto tmp = __ockl_image_sample_1D(i, s, x);
64 return __hipMapFrom<T>(tmp);
69 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
70static __device__ __hip_img_chk__
void tex1D(T *ptr, hipTextureObject_t textureObject,
float x)
72 *ptr = tex1D<T>(textureObject, x);
77 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
78static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject,
float x,
float y)
80 TEXTURE_OBJECT_PARAMETERS_INIT
81 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
82 return __hipMapFrom<T>(tmp);
87 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
88static __device__ __hip_img_chk__
void tex2D(T *ptr, hipTextureObject_t textureObject,
float x,
float y)
90 *ptr = tex2D<T>(textureObject, x, y);
95 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
96static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject,
float x,
float y,
float z)
98 TEXTURE_OBJECT_PARAMETERS_INIT
99 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
100 return __hipMapFrom<T>(tmp);
105 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
106static __device__ __hip_img_chk__
void tex3D(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
108 *ptr = tex3D<T>(textureObject, x, y, z);
113 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
114static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject,
float x,
int layer)
116 TEXTURE_OBJECT_PARAMETERS_INIT
117 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
118 return __hipMapFrom<T>(tmp);
123 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
124static __device__ __hip_img_chk__
void tex1DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
int layer)
126 *ptr = tex1DLayered<T>(textureObject, x, layer);
131 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
132static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject,
float x,
float y,
int layer)
134 TEXTURE_OBJECT_PARAMETERS_INIT
135 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
136 return __hipMapFrom<T>(tmp);
141 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
142static __device__ __hip_img_chk__
void tex2DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer)
144 *ptr = tex1DLayered<T>(textureObject, x, y, layer);
149 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
150static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject,
float x,
float y,
float z)
152 TEXTURE_OBJECT_PARAMETERS_INIT
153 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
154 return __hipMapFrom<T>(tmp);
159 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
160static __device__ __hip_img_chk__
void texCubemap(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
162 *ptr = texCubemap<T>(textureObject, x, y, z);
167 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
168static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
170 TEXTURE_OBJECT_PARAMETERS_INIT
171 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
172 return __hipMapFrom<T>(tmp);
177 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
178static __device__ __hip_img_chk__
void texCubemapLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
180 *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
185 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
186static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
188 TEXTURE_OBJECT_PARAMETERS_INIT
191 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
192 return __hipMapFrom<T>(tmp);
196 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
197 return __hipMapFrom<T>(tmp);
201 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
202 return __hipMapFrom<T>(tmp);
206 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
207 return __hipMapFrom<T>(tmp);
216 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
217static __device__ __hip_img_chk__
void tex2Dgather(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
219 *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
224 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
225static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject,
float x,
float level)
227 TEXTURE_OBJECT_PARAMETERS_INIT
228 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
229 return __hipMapFrom<T>(tmp);
234 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
235static __device__ __hip_img_chk__
void tex1DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float level)
237 *ptr = tex1DLod<T>(textureObject, x, level);
242 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
243static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject,
float x,
float y,
float level)
245 TEXTURE_OBJECT_PARAMETERS_INIT
246 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
247 return __hipMapFrom<T>(tmp);
252 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
253static __device__ __hip_img_chk__
void tex2DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float level)
255 *ptr = tex2DLod<T>(textureObject, x, y, level);
260 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
261static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
263 TEXTURE_OBJECT_PARAMETERS_INIT
264 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
265 return __hipMapFrom<T>(tmp);
270 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
271static __device__ __hip_img_chk__
void tex3DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
273 *ptr = tex3DLod<T>(textureObject, x, y, z, level);
278 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
279static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject,
float x,
int layer,
float level)
281 TEXTURE_OBJECT_PARAMETERS_INIT
282 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
283 return __hipMapFrom<T>(tmp);
288 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
289static __device__ __hip_img_chk__
void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float level)
291 *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
296 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
297static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
299 TEXTURE_OBJECT_PARAMETERS_INIT
300 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
301 return __hipMapFrom<T>(tmp);
306 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
307static __device__ __hip_img_chk__
void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
309 *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
314 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
315static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
317 TEXTURE_OBJECT_PARAMETERS_INIT
318 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
319 return __hipMapFrom<T>(tmp);
324 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
325static __device__ __hip_img_chk__
void texCubemapLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
327 *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
332 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
333static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
335 TEXTURE_OBJECT_PARAMETERS_INIT
344 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
345static __device__ __hip_img_chk__
void texCubemapGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
347 *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
352 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
353static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
355 TEXTURE_OBJECT_PARAMETERS_INIT
356 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
357 return __hipMapFrom<T>(tmp);
362 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
363static __device__ __hip_img_chk__
void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
365 *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
370 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
371static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
373 TEXTURE_OBJECT_PARAMETERS_INIT
374 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
375 return __hipMapFrom<T>(tmp);
380 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
381static __device__ __hip_img_chk__
void tex1DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
383 *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
388 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
389static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
391 TEXTURE_OBJECT_PARAMETERS_INIT
392 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);
393 return __hipMapFrom<T>(tmp);
398 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
399static __device__ __hip_img_chk__
void tex2DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
401 *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
406 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
407static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
409 TEXTURE_OBJECT_PARAMETERS_INIT
410 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);
411 return __hipMapFrom<T>(tmp);
416 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
417static __device__ __hip_img_chk__
void tex3DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
419 *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
424 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
425static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
427 TEXTURE_OBJECT_PARAMETERS_INIT
428 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
429 return __hipMapFrom<T>(tmp);
434 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
435static __device__ __hip_img_chk__
void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
437 *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
442 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
443static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
445 TEXTURE_OBJECT_PARAMETERS_INIT
446 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);
447 return __hipMapFrom<T>(tmp);
452 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
453static __device__ __hip_img_chk__
void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
455 *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
460 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
461static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
463 TEXTURE_OBJECT_PARAMETERS_INIT
472 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
473static __device__ __hip_img_chk__
void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
475 *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:1993