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