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