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