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