HIP: Heterogenous-computing Interface for Portability
texture_indirect_functions.h
1 /*
2 Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
23 #pragma once
24 
25 #if defined(__cplusplus)
26 
27 #include <hip/hip_vector_types.h>
28 #include <hip/hip_texture_types.h>
29 #include <hip/amd_detail/ockl_image.h>
30 
31 #if !defined(__HIPCC_RTC__)
32 #include <type_traits>
33 #endif // !defined(__HIPCC_RTC__)
34 
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;
38 
39 template<typename T>
40 struct __hip_is_itex_channel_type
41 {
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;
50 };
51 
52 template<
53  typename T,
54  unsigned int rank>
55 struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
56 {
57  static constexpr bool value =
58  __hip_is_itex_channel_type<T>::value &&
59  ((rank == 1) ||
60  (rank == 2) ||
61  (rank == 4));
62 };
63 
64 template <
65  typename T,
66  typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
67 static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
68 {
69  TEXTURE_OBJECT_PARAMETERS_INIT
70  auto tmp = __ockl_image_load_1Db(i, x);
71  return *reinterpret_cast<T*>(&tmp);
72 }
73 
74 template <
75  typename T,
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)
78 {
79  *ptr = tex1Dfetch<T>(textureObject, x);
80 }
81 
82 template <
83  typename T,
84  typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
85 static __device__ T tex1D(hipTextureObject_t textureObject, float x)
86 {
87  TEXTURE_OBJECT_PARAMETERS_INIT
88  auto tmp = __ockl_image_sample_1D(i, s, x);
89  return *reinterpret_cast<T*>(&tmp);
90 }
91 
92 template <
93  typename T,
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)
96 {
97  *ptr = tex1D<T>(textureObject, x);
98 }
99 
100 template <
101  typename T,
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)
104 {
105  TEXTURE_OBJECT_PARAMETERS_INIT
106  auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
107  return *reinterpret_cast<T*>(&tmp);
108 }
109 
110 template <
111  typename T,
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)
114 {
115  *ptr = tex2D<T>(textureObject, x, y);
116 }
117 
118 template <
119  typename T,
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)
122 {
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);
126 }
127 
128 template <
129  typename T,
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)
132 {
133  *ptr = tex3D<T>(textureObject, x, y, z);
134 }
135 
136 template <
137  typename T,
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)
140 {
141  TEXTURE_OBJECT_PARAMETERS_INIT
142  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
143  return *reinterpret_cast<T*>(&tmp);
144 }
145 
146 template <
147  typename T,
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)
150 {
151  *ptr = tex1DLayered<T>(textureObject, x, layer);
152 }
153 
154 template <
155  typename T,
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)
158 {
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);
162 }
163 
164 template <
165  typename T,
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)
168 {
169  *ptr = tex1DLayered<T>(textureObject, x, y, layer);
170 }
171 
172 template <
173  typename T,
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)
176 {
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);
180 }
181 
182 template <
183  typename T,
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)
186 {
187  *ptr = texCubemap<T>(textureObject, x, y, z);
188 }
189 
190 template <
191  typename T,
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)
194 {
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);
198 }
199 
200 template <
201  typename T,
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)
204 {
205  *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
206 }
207 
208 template <
209  typename T,
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)
212 {
213  TEXTURE_OBJECT_PARAMETERS_INIT
214  switch (comp) {
215  case 1: {
216  auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
217  return *reinterpret_cast<T*>(&tmp);
218  break;
219  }
220  case 2: {
221  auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
222  return *reinterpret_cast<T*>(&tmp);
223  break;
224  }
225  case 3: {
226  auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
227  return *reinterpret_cast<T*>(&tmp);
228  break;
229  }
230  default: {
231  auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
232  return *reinterpret_cast<T*>(&tmp);
233  break;
234  }
235  };
236  return {};
237 }
238 
239 template <
240  typename T,
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)
243 {
244  *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
245 }
246 
247 template <
248  typename T,
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)
251 {
252  TEXTURE_OBJECT_PARAMETERS_INIT
253  auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
254  return *reinterpret_cast<T*>(&tmp);
255 }
256 
257 template <
258  typename T,
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)
261 {
262  *ptr = tex1DLod<T>(textureObject, x, level);
263 }
264 
265 template <
266  typename T,
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)
269 {
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);
273 }
274 
275 template <
276  typename T,
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)
279 {
280  *ptr = tex2DLod<T>(textureObject, x, y, level);
281 }
282 
283 template <
284  typename T,
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)
287 {
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);
291 }
292 
293 template <
294  typename T,
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)
297 {
298  *ptr = tex3DLod<T>(textureObject, x, y, z, level);
299 }
300 
301 template <
302  typename T,
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)
305 {
306  TEXTURE_OBJECT_PARAMETERS_INIT
307  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
308  return *reinterpret_cast<T*>(&tmp);
309 }
310 
311 template <
312  typename T,
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)
315 {
316  *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
317 }
318 
319 template <
320  typename T,
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)
323 {
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);
327 }
328 
329 template <
330  typename T,
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)
333 {
334  *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
335 }
336 
337 template <
338  typename T,
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)
341 {
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);
345 }
346 
347 template <
348  typename T,
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)
351 {
352  *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
353 }
354 
355 template <
356  typename T,
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)
359 {
360  TEXTURE_OBJECT_PARAMETERS_INIT
361  // TODO missing in device libs.
362  // auto tmp = __ockl_image_sample_grad_CM(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);
363  // return *reinterpret_cast<T*>(&tmp);
364  return {};
365 }
366 
367 template <
368  typename T,
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)
371 {
372  *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
373 }
374 
375 template <
376  typename T,
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)
379 {
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);
383 }
384 
385 template <
386  typename T,
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)
389 {
390  *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
391 }
392 
393 template <
394  typename T,
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)
397 {
398  TEXTURE_OBJECT_PARAMETERS_INIT
399  auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
400  return *reinterpret_cast<T*>(&tmp);
401 }
402 
403 template <
404  typename T,
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)
407 {
408  *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
409 }
410 
411 template <
412  typename T,
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)
415 {
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);
419 }
420 
421 template <
422  typename T,
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)
425 {
426  *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
427 }
428 
429 template <
430  typename T,
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)
433 {
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);
437 }
438 
439 template <
440  typename T,
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)
443 {
444  *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
445 }
446 
447 template <
448  typename T,
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)
451 {
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);
455 }
456 
457 template <
458  typename T,
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)
461 {
462  *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
463 }
464 
465 template <
466  typename T,
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)
469 {
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);
473 }
474 
475 template <
476  typename T,
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)
479 {
480  *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
481 }
482 
483 template <
484  typename T,
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)
487 {
488  TEXTURE_OBJECT_PARAMETERS_INIT
489  // TODO missing in device libs.
490  // auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
491  // return *reinterpret_cast<T*>(&tmp);
492  return {};
493 }
494 
495 template <
496  typename T,
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)
499 {
500  *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
501 }
502 
503 #endif
float4
Definition: hip_vector_types.h:1583
float2
Definition: hip_vector_types.h:1582