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/hcc_detail/ockl_image.h>
30 
31 #include <type_traits>
32 
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;
36 
37 template<typename T>
38 struct __hip_is_itex_channel_type
39 {
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;
48 };
49 
50 template<
51  typename T,
52  unsigned int rank>
53 struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
54 {
55  static constexpr bool value =
56  __hip_is_itex_channel_type<T>::value &&
57  ((rank == 1) ||
58  (rank == 2) ||
59  (rank == 4));
60 };
61 
62 template <
63  typename T,
64  typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
65 static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
66 {
67  TEXTURE_OBJECT_PARAMETERS_INIT
68  auto tmp = __ockl_image_load_1Db(i, x);
69  return *reinterpret_cast<T*>(&tmp);
70 }
71 
72 template <
73  typename T,
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)
76 {
77  *ptr = tex1Dfetch<T>(textureObject, x);
78 }
79 
80 template <
81  typename T,
82  typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
83 static __device__ T tex1D(hipTextureObject_t textureObject, float x)
84 {
85  TEXTURE_OBJECT_PARAMETERS_INIT
86  auto tmp = __ockl_image_sample_1D(i, s, x);
87  return *reinterpret_cast<T*>(&tmp);
88 }
89 
90 template <
91  typename T,
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)
94 {
95  *ptr = tex1D<T>(textureObject, x);
96 }
97 
98 template <
99  typename T,
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)
102 {
103  TEXTURE_OBJECT_PARAMETERS_INIT
104  auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
105  return *reinterpret_cast<T*>(&tmp);
106 }
107 
108 template <
109  typename T,
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)
112 {
113  *ptr = tex2D<T>(textureObject, x, y);
114 }
115 
116 template <
117  typename T,
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)
120 {
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);
124 }
125 
126 template <
127  typename T,
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)
130 {
131  *ptr = tex3D<T>(textureObject, x, y, z);
132 }
133 
134 template <
135  typename T,
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)
138 {
139  TEXTURE_OBJECT_PARAMETERS_INIT
140  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
141  return *reinterpret_cast<T*>(&tmp);
142 }
143 
144 template <
145  typename T,
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)
148 {
149  *ptr = tex1DLayered<T>(textureObject, x, layer);
150 }
151 
152 template <
153  typename T,
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)
156 {
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);
160 }
161 
162 template <
163  typename T,
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)
166 {
167  *ptr = tex1DLayered<T>(textureObject, x, y, layer);
168 }
169 
170 template <
171  typename T,
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)
174 {
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);
178 }
179 
180 template <
181  typename T,
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)
184 {
185  *ptr = texCubemap<T>(textureObject, x, y, z);
186 }
187 
188 template <
189  typename T,
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)
192 {
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);
196 }
197 
198 template <
199  typename T,
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)
202 {
203  *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
204 }
205 
206 template <
207  typename T,
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)
210 {
211  TEXTURE_OBJECT_PARAMETERS_INIT
212  switch (comp) {
213  case 1: {
214  auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
215  return *reinterpret_cast<T*>(&tmp);
216  break;
217  }
218  case 2: {
219  auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
220  return *reinterpret_cast<T*>(&tmp);
221  break;
222  }
223  case 3: {
224  auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
225  return *reinterpret_cast<T*>(&tmp);
226  break;
227  }
228  default: {
229  auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
230  return *reinterpret_cast<T*>(&tmp);
231  break;
232  }
233  };
234  return {};
235 }
236 
237 template <
238  typename T,
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)
241 {
242  *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
243 }
244 
245 template <
246  typename T,
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)
249 {
250  TEXTURE_OBJECT_PARAMETERS_INIT
251  auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
252  return *reinterpret_cast<T*>(&tmp);
253 }
254 
255 template <
256  typename T,
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)
259 {
260  *ptr = tex1DLod<T>(textureObject, x, level);
261 }
262 
263 template <
264  typename T,
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)
267 {
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);
271 }
272 
273 template <
274  typename T,
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)
277 {
278  *ptr = tex2DLod<T>(textureObject, x, y, level);
279 }
280 
281 template <
282  typename T,
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)
285 {
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);
289 }
290 
291 template <
292  typename T,
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)
295 {
296  *ptr = tex3DLod<T>(textureObject, x, y, z, level);
297 }
298 
299 template <
300  typename T,
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)
303 {
304  TEXTURE_OBJECT_PARAMETERS_INIT
305  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
306  return *reinterpret_cast<T*>(&tmp);
307 }
308 
309 template <
310  typename T,
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)
313 {
314  *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
315 }
316 
317 template <
318  typename T,
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)
321 {
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);
325 }
326 
327 template <
328  typename T,
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)
331 {
332  *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
333 }
334 
335 template <
336  typename T,
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)
339 {
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);
343 }
344 
345 template <
346  typename T,
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)
349 {
350  *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
351 }
352 
353 template <
354  typename T,
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)
357 {
358  TEXTURE_OBJECT_PARAMETERS_INIT
359  // TODO missing in device libs.
360  // 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);
361  // return *reinterpret_cast<T*>(&tmp);
362  return {};
363 }
364 
365 template <
366  typename T,
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)
369 {
370  *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
371 }
372 
373 template <
374  typename T,
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)
377 {
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);
381 }
382 
383 template <
384  typename T,
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)
387 {
388  *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
389 }
390 
391 template <
392  typename T,
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)
395 {
396  TEXTURE_OBJECT_PARAMETERS_INIT
397  auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
398  return *reinterpret_cast<T*>(&tmp);
399 }
400 
401 template <
402  typename T,
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)
405 {
406  *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
407 }
408 
409 template <
410  typename T,
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)
413 {
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);
417 }
418 
419 template <
420  typename T,
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)
423 {
424  *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
425 }
426 
427 template <
428  typename T,
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)
431 {
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);
435 }
436 
437 template <
438  typename T,
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)
441 {
442  *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
443 }
444 
445 template <
446  typename T,
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)
449 {
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);
453 }
454 
455 template <
456  typename T,
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)
459 {
460  *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
461 }
462 
463 template <
464  typename T,
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)
467 {
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);
471 }
472 
473 template <
474  typename T,
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)
477 {
478  *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
479 }
480 
481 template <
482  typename T,
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)
485 {
486  TEXTURE_OBJECT_PARAMETERS_INIT
487  // TODO missing in device libs.
488  // 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);
489  // return *reinterpret_cast<T*>(&tmp);
490  return {};
491 }
492 
493 template <
494  typename T,
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)
497 {
498  *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
499 }
500 
501 #endif
float4
Definition: hip_vector_types.h:1579
float2
Definition: hip_vector_types.h:1578