HIP: Heterogenous-computing Interface for Portability
texture_fetch_functions.h
1 /*
2 Copyright (c) 2015 - present 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 #include <hip/hip_vector_types.h>
28 #include <hip/texture_types.h>
29 #include <hip/amd_detail/ockl_image.h>
30 
31 #include <type_traits>
32 
33 #define TEXTURE_PARAMETERS_INIT \
34  unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
35  unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
36 
37 template<typename T>
38 struct __hip_is_tex_channel_type
39 {
40  static constexpr bool value =
41  std::is_same<T, char>::value ||
42  std::is_same<T, unsigned char>::value ||
43  std::is_same<T, short>::value ||
44  std::is_same<T, unsigned short>::value ||
45  std::is_same<T, int>::value ||
46  std::is_same<T, unsigned int>::value ||
47  std::is_same<T, float>::value;
48 };
49 
50 template<
51  typename T,
52  unsigned int rank>
53 struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
54 {
55  static constexpr bool value =
56  __hip_is_tex_channel_type<T>::value &&
57  ((rank == 1) ||
58  (rank == 2) ||
59  (rank == 4));
60 };
61 
62 template<typename T>
63 struct __hip_is_tex_normalized_channel_type
64 {
65  static constexpr bool value =
66  std::is_same<T, char>::value ||
67  std::is_same<T, unsigned char>::value ||
68  std::is_same<T, short>::value ||
69  std::is_same<T, unsigned short>::value;
70 };
71 
72 template<
73  typename T,
74  unsigned int rank>
75 struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
76 {
77  static constexpr bool value =
78  __hip_is_tex_normalized_channel_type<T>::value &&
79  ((rank == 1) ||
80  (rank == 2) ||
81  (rank == 4));
82 };
83 
84 template <
85  typename T,
86  hipTextureReadMode readMode,
87  typename Enable = void>
88 struct __hip_tex_ret
89 {
90  static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
91 };
92 
93 template <
94  typename T,
95  hipTextureReadMode readMode>
96 using __hip_tex_ret_t = typename __hip_tex_ret<T, readMode, bool>::type;
97 
98 template <typename T>
99 struct __hip_tex_ret<
100  T,
101  hipReadModeElementType,
102  typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
103 {
104  using type = T;
105 };
106 
107 template<
108  typename T,
109  unsigned int rank>
110 struct __hip_tex_ret<
111  HIP_vector_type<T, rank>,
112  hipReadModeElementType,
113  typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
114 {
115  using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
116 };
117 
118 template<typename T>
119 struct __hip_tex_ret<
120  T,
121  hipReadModeNormalizedFloat,
122  typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
123 {
124  using type = float;
125 };
126 
127 template<
128  typename T,
129  unsigned int rank>
130 struct __hip_tex_ret<
131  HIP_vector_type<T, rank>,
132  hipReadModeNormalizedFloat,
133  typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
134 {
135  using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
136 };
137 
138 template <typename T, hipTextureReadMode readMode>
139 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t, int x)
140 {
141  TEXTURE_PARAMETERS_INIT;
142  auto tmp = __ockl_image_load_1Db(i, x);
143  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
144 }
145 
146 template <typename T, hipTextureReadMode readMode>
147 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t, float x)
148 {
149  TEXTURE_PARAMETERS_INIT;
150  auto tmp = __ockl_image_sample_1D(i, s, x);
151  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
152 }
153 
154 template <typename T, hipTextureReadMode readMode>
155 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t, float x, float y)
156 {
157  TEXTURE_PARAMETERS_INIT;
158  auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
159  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
160 }
161 
162 template <typename T, hipTextureReadMode readMode>
163 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer)
164 {
165  TEXTURE_PARAMETERS_INIT;
166  auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
167  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
168 }
169 
170 template <typename T, hipTextureReadMode readMode>
171 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer)
172 {
173  TEXTURE_PARAMETERS_INIT;
174  auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
175  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
176 }
177 
178 template <typename T, hipTextureReadMode readMode>
179 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t, float x, float y, float z)
180 {
181  TEXTURE_PARAMETERS_INIT;
182  auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
183  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
184 }
185 
186 template <typename T, hipTextureReadMode readMode>
187 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z)
188 {
189  TEXTURE_PARAMETERS_INIT;
190  auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
191  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
192 }
193 
194 template <typename T, hipTextureReadMode readMode>
195 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t, float x, float level)
196 {
197  TEXTURE_PARAMETERS_INIT;
198  auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
199  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
200 }
201 
202 template <typename T, hipTextureReadMode readMode>
203 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t, float x, float y, float level)
204 {
205  TEXTURE_PARAMETERS_INIT;
206  auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
207  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
208 }
209 
210 template <typename T, hipTextureReadMode readMode>
211 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float level)
212 {
213  TEXTURE_PARAMETERS_INIT;
214  auto tmp = __ockl_image_sample_lod_1Da(i, s, float2(x, layer).data, level);
215  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
216 }
217 
218 template <typename T, hipTextureReadMode readMode>
219 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float level)
220 {
221  TEXTURE_PARAMETERS_INIT;
222  auto tmp = __ockl_image_sample_lod_2Da(i, s, float4(x, y, layer, 0.0f).data, level);
223  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
224 }
225 
226 template <typename T, hipTextureReadMode readMode>
227 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float level)
228 {
229  TEXTURE_PARAMETERS_INIT;
230  auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
231  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
232 }
233 
234 template <typename T, hipTextureReadMode readMode>
235 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float level)
236 {
237  TEXTURE_PARAMETERS_INIT;
238  auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
239  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
240 }
241 
242 template <typename T, hipTextureReadMode readMode>
243 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer)
244 {
245  TEXTURE_PARAMETERS_INIT;
246  auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
247  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
248 }
249 
250 template <typename T, hipTextureReadMode readMode>
251 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredLod(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float level)
252 {
253  TEXTURE_PARAMETERS_INIT;
254  auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
255  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
256 }
257 
258 template <typename T, hipTextureReadMode readMode>
259 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
260 {
261  TEXTURE_PARAMETERS_INIT;
262  // TODO missing in device libs.
263  // 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);
264  // return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
265  return {};
266 }
267 
268 template <typename T, hipTextureReadMode readMode>
269 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
270 {
271  TEXTURE_PARAMETERS_INIT;
272  // TODO missing in device libs.
273  // 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);
274  // return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
275  return {};
276 }
277 
278 template <typename T, hipTextureReadMode readMode>
279 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t, float x, float dPdx, float dPdy)
280 {
281  TEXTURE_PARAMETERS_INIT;
282  auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
283  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
284 }
285 
286 template <typename T, hipTextureReadMode readMode>
287 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t, float x, float y, float2 dPdx, float2 dPdy)
288 {
289  TEXTURE_PARAMETERS_INIT;
290  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);
291  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
292 }
293 
294 template <typename T, hipTextureReadMode readMode>
295 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float dPdx, float dPdy)
296 {
297  TEXTURE_PARAMETERS_INIT;
298  auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
299  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
300 }
301 
302 template <typename T, hipTextureReadMode readMode>
303 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float2 dPdx, float2 dPdy)
304 {
305  TEXTURE_PARAMETERS_INIT;
306  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);
307  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
308 }
309 
310 template <typename T, hipTextureReadMode readMode>
311 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
312 {
313  TEXTURE_PARAMETERS_INIT;
314  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);
315  return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
316 }
317 
318 template <
319  typename T,
320  hipTextureReadMode readMode,
321  typename Enable = void>
322 struct __hip_tex2dgather_ret
323 {
324  static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
325 };
326 
327 template <
328  typename T,
329  hipTextureReadMode readMode>
330 using __hip_tex2dgather_ret_t = typename __hip_tex2dgather_ret<T, readMode, bool>::type;
331 
332 template <typename T>
333 struct __hip_tex2dgather_ret<
334  T,
335  hipReadModeElementType,
336  typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
337 {
338  using type = HIP_vector_type<T, 4>;
339 };
340 
341 template<
342  typename T,
343  unsigned int rank>
344 struct __hip_tex2dgather_ret<
345  HIP_vector_type<T, rank>,
346  hipReadModeElementType,
347  typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
348 {
349  using type = HIP_vector_type<T, 4>;
350 };
351 
352 template <typename T>
353 struct __hip_tex2dgather_ret<
354  T,
355  hipReadModeNormalizedFloat,
356  typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
357 {
358  using type = float4;
359 };
360 
361 template <typename T, hipTextureReadMode readMode>
362 static __forceinline__ __device__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t, float x, float y, int comp=0)
363 {
364  TEXTURE_PARAMETERS_INIT;
365  switch (comp) {
366  case 1: {
367  auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
368  return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
369  }
370  case 2: {
371  auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
372  return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
373  }
374  case 3: {
375  auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
376  return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
377  }
378  default: {
379  auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
380  return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
381  }
382  }
383  return {};
384 }
385 
386 #endif
float4
Definition: hip_vector_types.h:1575
float2
Definition: hip_vector_types.h:1574