HIP: Heterogenous-computing Interface for Portability
texture_indirect_functions.h
1 /*
2 Copyright (c) 2015 - 2023 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 #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>
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 <
40  typename T,
41  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
42 static __device__ __hip_img_chk__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
43 {
44  TEXTURE_OBJECT_PARAMETERS_INIT
45  auto tmp = __ockl_image_load_1Db(i, x);
46  return __hipMapFrom<T>(tmp);
47 }
48 
49 template <
50  typename T,
51  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
52 static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
53 {
54  *ptr = tex1Dfetch<T>(textureObject, x);
55 }
56 
57 template <
58  typename T,
59  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
60 static __device__ __hip_img_chk__ T tex1D(hipTextureObject_t textureObject, float x)
61 {
62  TEXTURE_OBJECT_PARAMETERS_INIT
63  auto tmp = __ockl_image_sample_1D(i, s, x);
64  return __hipMapFrom<T>(tmp);
65 }
66 
67 template <
68  typename T,
69  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
70 static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
71 {
72  *ptr = tex1D<T>(textureObject, x);
73 }
74 
75 template <
76  typename T,
77  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
78 static __device__ __hip_img_chk__ T tex2D(hipTextureObject_t textureObject, float x, float y)
79 {
80  TEXTURE_OBJECT_PARAMETERS_INIT
81  auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
82  return __hipMapFrom<T>(tmp);
83 }
84 
85 template <
86  typename T,
87  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
88 static __device__ __hip_img_chk__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
89 {
90  *ptr = tex2D<T>(textureObject, x, y);
91 }
92 
93 template <
94  typename T,
95  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
96 static __device__ __hip_img_chk__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
97 {
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);
101 }
102 
103 template <
104  typename T,
105  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
106 static __device__ __hip_img_chk__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
107 {
108  *ptr = tex3D<T>(textureObject, x, y, z);
109 }
110 
111 template <
112  typename T,
113  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
114 static __device__ __hip_img_chk__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
115 {
116  TEXTURE_OBJECT_PARAMETERS_INIT
117  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
118  return __hipMapFrom<T>(tmp);
119 }
120 
121 template <
122  typename T,
123  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
124 static __device__ __hip_img_chk__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
125 {
126  *ptr = tex1DLayered<T>(textureObject, x, layer);
127 }
128 
129 template <
130  typename T,
131  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
132 static __device__ __hip_img_chk__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
133 {
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);
137 }
138 
139 template <
140  typename T,
141  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
142 static __device__ __hip_img_chk__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
143 {
144  *ptr = tex1DLayered<T>(textureObject, x, y, layer);
145 }
146 
147 template <
148  typename T,
149  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
150 static __device__ __hip_img_chk__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
151 {
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);
155 }
156 
157 template <
158  typename T,
159  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
160 static __device__ __hip_img_chk__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
161 {
162  *ptr = texCubemap<T>(textureObject, x, y, z);
163 }
164 
165 template <
166  typename T,
167  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
168 static __device__ __hip_img_chk__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
169 {
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);
173 }
174 
175 template <
176  typename T,
177  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
178 static __device__ __hip_img_chk__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
179 {
180  *ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
181 }
182 
183 template <
184  typename T,
185  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
186 static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
187 {
188  TEXTURE_OBJECT_PARAMETERS_INIT
189  switch (comp) {
190  case 1: {
191  auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
192  return __hipMapFrom<T>(tmp);
193  break;
194  }
195  case 2: {
196  auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
197  return __hipMapFrom<T>(tmp);
198  break;
199  }
200  case 3: {
201  auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
202  return __hipMapFrom<T>(tmp);
203  break;
204  }
205  default: {
206  auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
207  return __hipMapFrom<T>(tmp);
208  break;
209  }
210  }
211  return {};
212 }
213 
214 template <
215  typename T,
216  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
217 static __device__ __hip_img_chk__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
218 {
219  *ptr = texCubemapLayered<T>(textureObject, x, y, comp);
220 }
221 
222 template <
223  typename T,
224  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
225 static __device__ __hip_img_chk__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
226 {
227  TEXTURE_OBJECT_PARAMETERS_INIT
228  auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
229  return __hipMapFrom<T>(tmp);
230 }
231 
232 template <
233  typename T,
234  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
235 static __device__ __hip_img_chk__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
236 {
237  *ptr = tex1DLod<T>(textureObject, x, level);
238 }
239 
240 template <
241  typename T,
242  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
243 static __device__ __hip_img_chk__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
244 {
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);
248 }
249 
250 template <
251  typename T,
252  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
253 static __device__ __hip_img_chk__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
254 {
255  *ptr = tex2DLod<T>(textureObject, x, y, level);
256 }
257 
258 template <
259  typename T,
260  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
261 static __device__ __hip_img_chk__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
262 {
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);
266 }
267 
268 template <
269  typename T,
270  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
271 static __device__ __hip_img_chk__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
272 {
273  *ptr = tex3DLod<T>(textureObject, x, y, z, level);
274 }
275 
276 template <
277  typename T,
278  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
279 static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
280 {
281  TEXTURE_OBJECT_PARAMETERS_INIT
282  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
283  return __hipMapFrom<T>(tmp);
284 }
285 
286 template <
287  typename T,
288  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
289 static __device__ __hip_img_chk__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
290 {
291  *ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
292 }
293 
294 template <
295  typename T,
296  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
297 static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
298 {
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);
302 }
303 
304 template <
305  typename T,
306  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
307 static __device__ __hip_img_chk__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
308 {
309  *ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
310 }
311 
312 template <
313  typename T,
314  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
315 static __device__ __hip_img_chk__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
316 {
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);
320 }
321 
322 template <
323  typename T,
324  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
325 static __device__ __hip_img_chk__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
326 {
327  *ptr = texCubemapLod<T>(textureObject, x, y, z, level);
328 }
329 
330 template <
331  typename T,
332  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
333 static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
334 {
335  TEXTURE_OBJECT_PARAMETERS_INIT
336  // TODO missing in device libs.
337  // 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);
338  // return __hipMapFrom<T>(tmp);
339  return {};
340 }
341 
342 template <
343  typename T,
344  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
345 static __device__ __hip_img_chk__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
346 {
347  *ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
348 }
349 
350 template <
351  typename T,
352  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
353 static __device__ __hip_img_chk__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
354 {
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);
358 }
359 
360 template <
361  typename T,
362  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
363 static __device__ __hip_img_chk__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
364 {
365  *ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
366 }
367 
368 template <
369  typename T,
370  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
371 static __device__ __hip_img_chk__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
372 {
373  TEXTURE_OBJECT_PARAMETERS_INIT
374  auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
375  return __hipMapFrom<T>(tmp);
376 }
377 
378 template <
379  typename T,
380  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
381 static __device__ __hip_img_chk__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
382 {
383  *ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
384 }
385 
386 template <
387  typename T,
388  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
389 static __device__ __hip_img_chk__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
390 {
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);
394 }
395 
396 template <
397  typename T,
398  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
399 static __device__ __hip_img_chk__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
400 {
401  *ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
402 }
403 
404 template <
405  typename T,
406  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
407 static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
408 {
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);
412 }
413 
414 template <
415  typename T,
416  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
417 static __device__ __hip_img_chk__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
418 {
419  *ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
420 }
421 
422 template <
423  typename T,
424  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
425 static __device__ __hip_img_chk__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
426 {
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);
430 }
431 
432 template <
433  typename T,
434  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
435 static __device__ __hip_img_chk__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
436 {
437  *ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
438 }
439 
440 template <
441  typename T,
442  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
443 static __device__ __hip_img_chk__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
444 {
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);
448 }
449 
450 template <
451  typename T,
452  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
453 static __device__ __hip_img_chk__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
454 {
455  *ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
456 }
457 
458 template <
459  typename T,
460  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
461 static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
462 {
463  TEXTURE_OBJECT_PARAMETERS_INIT
464  // TODO missing in device libs.
465  // 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);
466  // return __hipMapFrom<T>(tmp);
467  return {};
468 }
469 
470 template <
471  typename T,
472  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
473 static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
474 {
475  *ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
476 }
477 
478 #endif
Definition: amd_hip_vector_types.h:2035
Definition: amd_hip_vector_types.h:2042