25 #if defined(__cplusplus)
27 #include <hip/hip_vector_types.h>
28 #include <hip/hip_texture_types.h>
29 #include <hip/hcc_detail/ockl_image.h>
31 #include <type_traits>
33 #define TEXTURE_OBJECT_PARAMETERS_INIT \
34 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
35 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
38 struct __hip_is_itex_channel_type
40 static constexpr
bool value =
41 std::is_same<T, char>::value ||
42 std::is_same<T, unsigned char>::value ||
43 std::is_same<T, short>::value ||
44 std::is_same<T, unsigned short>::value ||
45 std::is_same<T, int>::value ||
46 std::is_same<T, unsigned int>::value ||
47 std::is_same<T, float>::value;
53 struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
55 static constexpr
bool value =
56 __hip_is_itex_channel_type<T>::value &&
64 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
65 static __device__ T tex1Dfetch(hipTextureObject_t textureObject,
int x)
67 TEXTURE_OBJECT_PARAMETERS_INIT
68 auto tmp = __ockl_image_load_1Db(i, x);
69 return *
reinterpret_cast<T*
>(&tmp);
74 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
75 static __device__
void tex1Dfetch(T *ptr, hipTextureObject_t textureObject,
int x)
77 *ptr = tex1Dfetch<T>(textureObject, x);
82 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
83 static __device__ T tex1D(hipTextureObject_t textureObject,
float x)
85 TEXTURE_OBJECT_PARAMETERS_INIT
86 auto tmp = __ockl_image_sample_1D(i, s, x);
87 return *
reinterpret_cast<T*
>(&tmp);
92 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
93 static __device__
void tex1D(T *ptr, hipTextureObject_t textureObject,
float x)
95 *ptr = tex1D<T>(textureObject, x);
100 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
101 static __device__ T tex2D(hipTextureObject_t textureObject,
float x,
float y)
103 TEXTURE_OBJECT_PARAMETERS_INIT
104 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
105 return *
reinterpret_cast<T*
>(&tmp);
110 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
111 static __device__
void tex2D(T *ptr, hipTextureObject_t textureObject,
float x,
float y)
113 *ptr = tex2D<T>(textureObject, x, y);
118 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
119 static __device__ T tex3D(hipTextureObject_t textureObject,
float x,
float y,
float z)
121 TEXTURE_OBJECT_PARAMETERS_INIT
122 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
123 return *
reinterpret_cast<T*
>(&tmp);
128 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
129 static __device__
void tex3D(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
131 *ptr = tex3D<T>(textureObject, x, y, z);
136 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
137 static __device__ T tex1DLayered(hipTextureObject_t textureObject,
float x,
int layer)
139 TEXTURE_OBJECT_PARAMETERS_INIT
140 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
141 return *
reinterpret_cast<T*
>(&tmp);
146 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
147 static __device__
void tex1DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
int layer)
149 *ptr = tex1DLayered<T>(textureObject, x, layer);
154 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
155 static __device__ T tex2DLayered(hipTextureObject_t textureObject,
float x,
float y,
int layer)
157 TEXTURE_OBJECT_PARAMETERS_INIT
158 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
159 return *
reinterpret_cast<T*
>(&tmp);
164 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
165 static __device__
void tex2DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer)
167 *ptr = tex1DLayered<T>(textureObject, x, y, layer);
172 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
173 static __device__ T texCubemap(hipTextureObject_t textureObject,
float x,
float y,
float z)
175 TEXTURE_OBJECT_PARAMETERS_INIT
176 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
177 return *
reinterpret_cast<T*
>(&tmp);
182 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
183 static __device__
void texCubemap(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
185 *ptr = texCubemap<T>(textureObject, x, y, z);
190 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
191 static __device__ T texCubemapLayered(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
193 TEXTURE_OBJECT_PARAMETERS_INIT
194 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
195 return *
reinterpret_cast<T*
>(&tmp);
200 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
201 static __device__
void texCubemapLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
203 *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
208 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
209 static __device__ T tex2Dgather(hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
211 TEXTURE_OBJECT_PARAMETERS_INIT
214 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
215 return *
reinterpret_cast<T*
>(&tmp);
219 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
220 return *
reinterpret_cast<T*
>(&tmp);
224 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
225 return *
reinterpret_cast<T*
>(&tmp);
229 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
230 return *
reinterpret_cast<T*
>(&tmp);
239 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
240 static __device__
void tex2Dgather(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
242 *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
247 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
248 static __device__ T tex1DLod(hipTextureObject_t textureObject,
float x,
float level)
250 TEXTURE_OBJECT_PARAMETERS_INIT
251 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
252 return *
reinterpret_cast<T*
>(&tmp);
257 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
258 static __device__
void tex1DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float level)
260 *ptr = tex1DLod<T>(textureObject, x, level);
265 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
266 static __device__ T tex2DLod(hipTextureObject_t textureObject,
float x,
float y,
float level)
268 TEXTURE_OBJECT_PARAMETERS_INIT
269 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
270 return *
reinterpret_cast<T*
>(&tmp);
275 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
276 static __device__
void tex2DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float level)
278 *ptr = tex2DLod<T>(textureObject, x, y, level);
283 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
284 static __device__ T tex3DLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
286 TEXTURE_OBJECT_PARAMETERS_INIT
287 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
288 return *
reinterpret_cast<T*
>(&tmp);
293 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
294 static __device__
void tex3DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
296 *ptr = tex3DLod<T>(textureObject, x, y, z, level);
301 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
302 static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject,
float x,
int layer,
float level)
304 TEXTURE_OBJECT_PARAMETERS_INIT
305 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
306 return *
reinterpret_cast<T*
>(&tmp);
311 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
312 static __device__
void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float level)
314 *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
319 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
320 static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
322 TEXTURE_OBJECT_PARAMETERS_INIT
323 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
324 return *
reinterpret_cast<T*
>(&tmp);
329 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
330 static __device__
void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
332 *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
337 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
338 static __device__ T texCubemapLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
340 TEXTURE_OBJECT_PARAMETERS_INIT
341 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
342 return *
reinterpret_cast<T*
>(&tmp);
347 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
348 static __device__
void texCubemapLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
350 *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
355 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
356 static __device__ T texCubemapGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
358 TEXTURE_OBJECT_PARAMETERS_INIT
367 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
368 static __device__
void texCubemapGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
370 *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
375 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
376 static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
378 TEXTURE_OBJECT_PARAMETERS_INIT
379 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
380 return *
reinterpret_cast<T*
>(&tmp);
385 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
386 static __device__
void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
388 *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
393 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
394 static __device__ T tex1DGrad(hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
396 TEXTURE_OBJECT_PARAMETERS_INIT
397 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
398 return *
reinterpret_cast<T*
>(&tmp);
403 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
404 static __device__
void tex1DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
406 *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
411 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
412 static __device__ T tex2DGrad(hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
414 TEXTURE_OBJECT_PARAMETERS_INIT
415 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);
416 return *
reinterpret_cast<T*
>(&tmp);
421 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
422 static __device__
void tex2DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
424 *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
429 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
430 static __device__ T tex3DGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
432 TEXTURE_OBJECT_PARAMETERS_INIT
433 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);
434 return *
reinterpret_cast<T*
>(&tmp);
439 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
440 static __device__
void tex3DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
442 *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
447 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
448 static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
450 TEXTURE_OBJECT_PARAMETERS_INIT
451 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
452 return *
reinterpret_cast<T*
>(&tmp);
457 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
458 static __device__
void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
460 *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
465 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
466 static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
468 TEXTURE_OBJECT_PARAMETERS_INIT
469 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);
470 return *
reinterpret_cast<T*
>(&tmp);
475 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
476 static __device__
void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
478 *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
483 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
484 static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
486 TEXTURE_OBJECT_PARAMETERS_INIT
495 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
496 static __device__
void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
498 *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);