HIP: Heterogenous-computing Interface for Portability
texture_fetch_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/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 
38 template<typename T>
39 struct __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 
51 template<typename T>
52 struct __hip_is_tex_surf_channel_type
53 {
54  static constexpr bool value =
55  __hip_is_tex_surf_scalar_channel_type<T>::value;
56 };
57 
58 template<
59  typename T,
60  unsigned int rank>
61 struct __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 
70 template<typename T>
71 struct __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 
80 template<
81  typename T,
82  unsigned int rank>
83 struct __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 
92 template <
93  typename T,
94  hipTextureReadMode readMode,
95  typename Enable = void>
96 struct __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  */
104 template<typename T, typename U>
105 __forceinline__ __device__
106 typename 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  */
127 template<typename T, typename U>
128 __forceinline__ __device__
129 typename 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  */
150 template<typename U, typename T>
151 __forceinline__ __device__
152 typename 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  */
175 template<typename U, typename T>
176 __forceinline__ __device__
177 typename 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 
197 template <
198  typename T,
199  hipTextureReadMode readMode>
200 using __hip_tex_ret_t = typename __hip_tex_ret<T, readMode, bool>::type;
201 
202 template <typename T>
203 struct __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 
211 template<
212  typename T,
213  unsigned int rank>
214 struct __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 
222 template<typename T>
223 struct __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 
231 template<
232  typename T,
233  unsigned int rank>
234 struct __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 
243 template <typename T, hipTextureReadMode readMode>
244 static __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 
251 template <typename T, hipTextureReadMode readMode>
252 static __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 
259 template <typename T, hipTextureReadMode readMode>
260 static __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 
267 template <typename T, hipTextureReadMode readMode>
268 static __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 
275 template <typename T, hipTextureReadMode readMode>
276 static __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 
283 template <typename T, hipTextureReadMode readMode>
284 static __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 
291 template <typename T, hipTextureReadMode readMode>
292 static __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 
299 template <typename T, hipTextureReadMode readMode>
300 static __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 
307 template <typename T, hipTextureReadMode readMode>
308 static __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 
315 template <typename T, hipTextureReadMode readMode>
316 static __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 
323 template <typename T, hipTextureReadMode readMode>
324 static __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 
331 template <typename T, hipTextureReadMode readMode>
332 static __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 
339 template <typename T, hipTextureReadMode readMode>
340 static __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 
347 template <typename T, hipTextureReadMode readMode>
348 static __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 
355 template <typename T, hipTextureReadMode readMode>
356 static __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 
363 template <typename T, hipTextureReadMode readMode>
364 static __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 
373 template <typename T, hipTextureReadMode readMode>
374 static __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 
383 template <typename T, hipTextureReadMode readMode>
384 static __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 
391 template <typename T, hipTextureReadMode readMode>
392 static __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 
399 template <typename T, hipTextureReadMode readMode>
400 static __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 
407 template <typename T, hipTextureReadMode readMode>
408 static __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 
415 template <typename T, hipTextureReadMode readMode>
416 static __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 
423 template <
424  typename T,
425  hipTextureReadMode readMode,
426  typename Enable = void>
427 struct __hip_tex2dgather_ret
428 {
429  static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
430 };
431 
432 template <
433  typename T,
434  hipTextureReadMode readMode>
435 using __hip_tex2dgather_ret_t = typename __hip_tex2dgather_ret<T, readMode, bool>::type;
436 
437 template <typename T>
438 struct __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 
446 template<
447  typename T,
448  unsigned int rank>
449 struct __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 
457 template <typename T>
458 struct __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 
466 template <typename T, hipTextureReadMode readMode>
467 static __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:1820
Definition: amd_hip_vector_types.h:2035
Definition: amd_hip_vector_types.h:2042