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__)
32 #include <type_traits>
33 #endif // !defined(__HIPCC_RTC__)
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;
40 struct __hip_is_itex_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;
55 struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
57 static constexpr
bool value =
58 __hip_is_itex_channel_type<T>::value &&
66 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
67 static __device__ T tex1Dfetch(hipTextureObject_t textureObject,
int x)
69 TEXTURE_OBJECT_PARAMETERS_INIT
70 auto tmp = __ockl_image_load_1Db(i, x);
71 return *
reinterpret_cast<T*
>(&tmp);
76 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
77 static __device__
void tex1Dfetch(T *ptr, hipTextureObject_t textureObject,
int x)
79 *ptr = tex1Dfetch<T>(textureObject, x);
84 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
85 static __device__ T tex1D(hipTextureObject_t textureObject,
float x)
87 TEXTURE_OBJECT_PARAMETERS_INIT
88 auto tmp = __ockl_image_sample_1D(i, s, x);
89 return *
reinterpret_cast<T*
>(&tmp);
94 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
95 static __device__
void tex1D(T *ptr, hipTextureObject_t textureObject,
float x)
97 *ptr = tex1D<T>(textureObject, x);
102 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
103 static __device__ T tex2D(hipTextureObject_t textureObject,
float x,
float y)
105 TEXTURE_OBJECT_PARAMETERS_INIT
106 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
107 return *
reinterpret_cast<T*
>(&tmp);
112 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
113 static __device__
void tex2D(T *ptr, hipTextureObject_t textureObject,
float x,
float y)
115 *ptr = tex2D<T>(textureObject, x, y);
120 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
121 static __device__ T tex3D(hipTextureObject_t textureObject,
float x,
float y,
float z)
123 TEXTURE_OBJECT_PARAMETERS_INIT
124 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
125 return *
reinterpret_cast<T*
>(&tmp);
130 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
131 static __device__
void tex3D(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
133 *ptr = tex3D<T>(textureObject, x, y, z);
138 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
139 static __device__ T tex1DLayered(hipTextureObject_t textureObject,
float x,
int layer)
141 TEXTURE_OBJECT_PARAMETERS_INIT
142 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
143 return *
reinterpret_cast<T*
>(&tmp);
148 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
149 static __device__
void tex1DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
int layer)
151 *ptr = tex1DLayered<T>(textureObject, x, layer);
156 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
157 static __device__ T tex2DLayered(hipTextureObject_t textureObject,
float x,
float y,
int layer)
159 TEXTURE_OBJECT_PARAMETERS_INIT
160 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
161 return *
reinterpret_cast<T*
>(&tmp);
166 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
167 static __device__
void tex2DLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer)
169 *ptr = tex1DLayered<T>(textureObject, x, y, layer);
174 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
175 static __device__ T texCubemap(hipTextureObject_t textureObject,
float x,
float y,
float z)
177 TEXTURE_OBJECT_PARAMETERS_INIT
178 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
179 return *
reinterpret_cast<T*
>(&tmp);
184 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
185 static __device__
void texCubemap(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z)
187 *ptr = texCubemap<T>(textureObject, x, y, z);
192 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
193 static __device__ T texCubemapLayered(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
195 TEXTURE_OBJECT_PARAMETERS_INIT
196 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
197 return *
reinterpret_cast<T*
>(&tmp);
202 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
203 static __device__
void texCubemapLayered(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer)
205 *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
210 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
211 static __device__ T tex2Dgather(hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
213 TEXTURE_OBJECT_PARAMETERS_INIT
216 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
217 return *
reinterpret_cast<T*
>(&tmp);
221 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
222 return *
reinterpret_cast<T*
>(&tmp);
226 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
227 return *
reinterpret_cast<T*
>(&tmp);
231 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
232 return *
reinterpret_cast<T*
>(&tmp);
241 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
242 static __device__
void tex2Dgather(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int comp = 0)
244 *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
249 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
250 static __device__ T tex1DLod(hipTextureObject_t textureObject,
float x,
float level)
252 TEXTURE_OBJECT_PARAMETERS_INIT
253 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
254 return *
reinterpret_cast<T*
>(&tmp);
259 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
260 static __device__
void tex1DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float level)
262 *ptr = tex1DLod<T>(textureObject, x, level);
267 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
268 static __device__ T tex2DLod(hipTextureObject_t textureObject,
float x,
float y,
float level)
270 TEXTURE_OBJECT_PARAMETERS_INIT
271 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
272 return *
reinterpret_cast<T*
>(&tmp);
277 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
278 static __device__
void tex2DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float level)
280 *ptr = tex2DLod<T>(textureObject, x, y, level);
285 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
286 static __device__ T tex3DLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
288 TEXTURE_OBJECT_PARAMETERS_INIT
289 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
290 return *
reinterpret_cast<T*
>(&tmp);
295 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
296 static __device__
void tex3DLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
298 *ptr = tex3DLod<T>(textureObject, x, y, z, level);
303 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
304 static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject,
float x,
int layer,
float level)
306 TEXTURE_OBJECT_PARAMETERS_INIT
307 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
308 return *
reinterpret_cast<T*
>(&tmp);
313 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
314 static __device__
void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float level)
316 *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
321 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
322 static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
324 TEXTURE_OBJECT_PARAMETERS_INIT
325 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
326 return *
reinterpret_cast<T*
>(&tmp);
331 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
332 static __device__
void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float level)
334 *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
339 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
340 static __device__ T texCubemapLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
342 TEXTURE_OBJECT_PARAMETERS_INIT
343 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
344 return *
reinterpret_cast<T*
>(&tmp);
349 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
350 static __device__
void texCubemapLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float level)
352 *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
357 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
358 static __device__ T texCubemapGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
360 TEXTURE_OBJECT_PARAMETERS_INIT
369 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
370 static __device__
void texCubemapGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
372 *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
377 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
378 static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
380 TEXTURE_OBJECT_PARAMETERS_INIT
381 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
382 return *
reinterpret_cast<T*
>(&tmp);
387 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
388 static __device__
void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float level)
390 *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
395 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
396 static __device__ T tex1DGrad(hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
398 TEXTURE_OBJECT_PARAMETERS_INIT
399 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
400 return *
reinterpret_cast<T*
>(&tmp);
405 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
406 static __device__
void tex1DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float dPdx,
float dPdy)
408 *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
413 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
414 static __device__ T tex2DGrad(hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
416 TEXTURE_OBJECT_PARAMETERS_INIT
417 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);
418 return *
reinterpret_cast<T*
>(&tmp);
423 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
424 static __device__
void tex2DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float2 dPdx,
float2 dPdy)
426 *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
431 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
432 static __device__ T tex3DGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
434 TEXTURE_OBJECT_PARAMETERS_INIT
435 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);
436 return *
reinterpret_cast<T*
>(&tmp);
441 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
442 static __device__
void tex3DGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
444 *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
449 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
450 static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
452 TEXTURE_OBJECT_PARAMETERS_INIT
453 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
454 return *
reinterpret_cast<T*
>(&tmp);
459 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
460 static __device__
void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
int layer,
float dPdx,
float dPdy)
462 *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
467 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
468 static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
470 TEXTURE_OBJECT_PARAMETERS_INIT
471 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);
472 return *
reinterpret_cast<T*
>(&tmp);
477 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
478 static __device__
void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
480 *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
485 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
486 static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
488 TEXTURE_OBJECT_PARAMETERS_INIT
497 typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* =
nullptr>
498 static __device__
void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
500 *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);