/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h Source File#

HIP Runtime API Reference: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h Source File
amd_surface_functions.h
Go to the documentation of this file.
1 /*
2 Copyright (c) 2018 - 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 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
24 #define HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
25 
26 #if defined(__cplusplus)
27 
28 #if !defined(__HIPCC_RTC__)
29 #include <hip/surface_types.h>
30 #include <hip/hip_vector_types.h>
31 #include <hip/amd_detail/texture_fetch_functions.h>
32 #include <hip/amd_detail/ockl_image.h>
33 #endif
34 
35 #if defined(__HIPCC_RTC__)
36 #define __HOST_DEVICE__ __device__
37 #else
38 #define __HOST_DEVICE__ __host__ __device__
39 #endif
40 
41 #define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
42  unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
43 
44 // CUDA is using byte address, need map to pixel address for HIP
45 static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) {
46  /*
47  * use below format index to generate format LUT
48  typedef enum {
49  HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0,
50  HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1,
51  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2,
52  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3,
53  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4,
54  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5,
55  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6,
56  HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7,
57  HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8,
58  HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9,
59  HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10,
60  HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11,
61  HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12,
62  HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13,
63  HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14,
64  HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15
65  } hsa_ext_image_channel_type_t;
66  */
67  static const int FormatLUT[] = { 0, 1, 0, 1, 3, 1, 1, 1, 0, 1, 2, 0, 1, 2, 1, 2 };
68  x = FormatLUT[format] == 3 ? x / FormatLUT[format] : x >> FormatLUT[format];
69 
70  /*
71  * use below order index to generate order LUT
72  typedef enum {
73  HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0,
74  HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1,
75  HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2,
76  HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3,
77  HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4,
78  HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5,
79  HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6,
80  HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7,
81  HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8,
82  HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9,
83  HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10,
84  HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11,
85  HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12,
86  HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13,
87  HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14,
88  HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15,
89  HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16,
90  HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17,
91  HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18,
92  HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19
93  } hsa_ext_image_channel_order_t;
94  */
95  static const int OrderLUT[] = { 0, 0, 1, 1, 3, 1, 3, 2, 2, 2, 2, 2, 3, 2, 2, 2, 0, 0, 0, 0 };
96  return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order];
97 }
98 
107 template <
108  typename T,
109  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
110 static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
111  int boundaryMode = hipBoundaryModeZero) {
112  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
113  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
114  auto tmp = __ockl_image_load_1D(i, x);
115  *data = __hipMapFrom<T>(tmp);
116 }
117 
125 template <
126  typename T,
127  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
128 static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
129  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
130  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
131  auto tmp = __hipMapTo<float4::Native_vec_>(data);
132  __ockl_image_store_1D(i, x, tmp);
133 }
134 
135 
144 template <
145  typename T,
146  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
147 static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
148  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
149  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
150  auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
151  *data = __hipMapFrom<T>(tmp);
152 }
153 
163 template <
164  typename T,
165  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
166 static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
167  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
168  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
169  auto tmp = __hipMapTo<float4::Native_vec_>(data);
170  __ockl_image_store_2D(i, int2(x, y).data, tmp);
171 }
172 
183 template <
184  typename T,
185  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
186 static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
187  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
188  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
189  auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
190  *data = __hipMapFrom<T>(tmp);
191 }
192 
203 template <
204  typename T,
205  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
206 static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
207  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
208  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
209  auto tmp = __hipMapTo<float4::Native_vec_>(data);
210  __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
211 }
212 
222 template <
223  typename T,
224  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
225 static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
226  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
227  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
228  auto tmp = __ockl_image_load_lod_1D(i, x, layer);
229  *data = __hipMapFrom<T>(tmp);
230 }
231 
241 template <
242  typename T,
243  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
244 static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
245  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
246  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
247  auto tmp = __hipMapTo<float4::Native_vec_>(data);
248  __ockl_image_store_lod_1D(i, x, layer, tmp);
249 }
250 
261 template <
262  typename T,
263  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
264 static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
265  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
266  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
267  auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
268  *data = __hipMapFrom<T>(tmp);
269 }
270 
281 template <
282  typename T,
283  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
284 static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
285  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
286  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
287  auto tmp = __hipMapTo<float4::Native_vec_>(data);
288  __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
289 }
290 
301 template <
302  typename T,
303  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
304 static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
305  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
306  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
307  auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
308  *data = __hipMapFrom<T>(tmp);
309 }
310 
321 template <
322  typename T,
323  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
324 static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
325  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
326  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
327  auto tmp = __hipMapTo<float4::Native_vec_>(data);
328  __ockl_image_store_CM(i, int2(x, y).data, face, tmp);
329 }
330 
342 template <
343  typename T,
344  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
345 static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
346  int layer) {
347  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
348  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
349  auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
350  *data = __hipMapFrom<T>(tmp);
351 }
352 
364 template <
365  typename T,
366  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
367 static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
368  int layer) {
369  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
370  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
371  auto tmp = __hipMapTo<float4::Native_vec_>(data);
372  __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
373 }
374 
375 #endif
376 
377 #endif
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT
Definition: amd_surface_functions.h:41
__HOST_DEVICE__
#define __HOST_DEVICE__
Definition: amd_surface_functions.h:38
hip_vector_types.h
surface_types.h
Defines surface types for HIP runtime.
hipBoundaryModeZero
@ hipBoundaryModeZero
Definition: surface_types.h:54
hipSurfaceObject_t
struct __hip_surface * hipSurfaceObject_t
Definition: surface_types.h:41