HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
texture_indirect_functions.h
1/*
2Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
3
4Permission is hereby granted, free of charge, to any person obtaining a copy
5of this software and associated documentation files (the "Software"), to deal
6in the Software without restriction, including without limitation the rights
7to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8copies of the Software, and to permit persons to whom the Software is
9furnished to do so, subject to the following conditions:
10
11The above copyright notice and this permission notice shall be included in
12all copies or substantial portions of the Software.
13
14THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20THE 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
39template <
40 typename T,
41 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
42static __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
49template <
50 typename T,
51 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
52static __device__ __hip_img_chk__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
53{
54 *ptr = tex1Dfetch<T>(textureObject, x);
55}
56
57template <
58 typename T,
59 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
60static __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
67template <
68 typename T,
69 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
70static __device__ __hip_img_chk__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
71{
72 *ptr = tex1D<T>(textureObject, x);
73}
74
75template <
76 typename T,
77 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
78static __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
85template <
86 typename T,
87 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
88static __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
93template <
94 typename T,
95 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
96static __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
103template <
104 typename T,
105 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
106static __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
111template <
112 typename T,
113 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
114static __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
121template <
122 typename T,
123 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
124static __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
129template <
130 typename T,
131 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
132static __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
139template <
140 typename T,
141 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
142static __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
147template <
148 typename T,
149 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
150static __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
157template <
158 typename T,
159 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
160static __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
165template <
166 typename T,
167 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
168static __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
175template <
176 typename T,
177 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
178static __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
183template <
184 typename T,
185 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
186static __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
214template <
215 typename T,
216 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
217static __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
222template <
223 typename T,
224 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
225static __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
232template <
233 typename T,
234 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
235static __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
240template <
241 typename T,
242 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
243static __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
250template <
251 typename T,
252 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
253static __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
258template <
259 typename T,
260 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
261static __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
268template <
269 typename T,
270 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
271static __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
276template <
277 typename T,
278 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
279static __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
286template <
287 typename T,
288 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
289static __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
294template <
295 typename T,
296 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
297static __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
304template <
305 typename T,
306 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
307static __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
312template <
313 typename T,
314 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
315static __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
322template <
323 typename T,
324 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
325static __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
330template <
331 typename T,
332 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
333static __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
342template <
343 typename T,
344 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
345static __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
350template <
351 typename T,
352 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
353static __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
360template <
361 typename T,
362 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
363static __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
368template <
369 typename T,
370 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
371static __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
378template <
379 typename T,
380 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
381static __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
386template <
387 typename T,
388 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
389static __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
396template <
397 typename T,
398 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
399static __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
404template <
405 typename T,
406 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
407static __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
414template <
415 typename T,
416 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
417static __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
422template <
423 typename T,
424 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
425static __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
432template <
433 typename T,
434 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
435static __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
440template <
441 typename T,
442 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
443static __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
450template <
451 typename T,
452 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
453static __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
458template <
459 typename T,
460 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
461static __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
470template <
471 typename T,
472 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
473static __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
float2
Definition amd_hip_vector_types.h:2035
float4
Definition amd_hip_vector_types.h:2042
Generated on Wed Mar 12 2025 01:08:27 for HIP: Heterogenous-computing Interface for Portability by doxygen 1.13.1