HIP: Heterogenous-computing Interface for Portability
amd_surface_functions.h
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 
99 template <
100  typename T,
101  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
102 static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
103  int boundaryMode = hipBoundaryModeZero) {
104  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
105  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
106  auto tmp = __ockl_image_load_1D(i, x);
107  *data = __hipMapFrom<T>(tmp);
108 }
109 
110 template <
111  typename T,
112  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
113 static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
114  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
115  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
116  auto tmp = __hipMapTo<float4::Native_vec_>(data);
117  __ockl_image_store_1D(i, x, tmp);
118 }
119 
120 template <
121  typename T,
122  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
123 static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
124  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
125  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
126  auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
127  *data = __hipMapFrom<T>(tmp);
128 }
129 
130 template <
131  typename T,
132  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
133 static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
134  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
135  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
136  auto tmp = __hipMapTo<float4::Native_vec_>(data);
137  __ockl_image_store_2D(i, int2(x, y).data, tmp);
138 }
139 
140 template <
141  typename T,
142  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
143 static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
144  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
145  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
146  auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
147  *data = __hipMapFrom<T>(tmp);
148 }
149 
150 template <
151  typename T,
152  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
153 static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
154  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
155  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
156  auto tmp = __hipMapTo<float4::Native_vec_>(data);
157  __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
158 }
159 
160 template <
161  typename T,
162  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
163 static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
164  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
165  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
166  auto tmp = __ockl_image_load_lod_1D(i, x, layer);
167  *data = __hipMapFrom<T>(tmp);
168 }
169 
170 template <
171  typename T,
172  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
173 static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
174  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
175  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
176  auto tmp = __hipMapTo<float4::Native_vec_>(data);
177  __ockl_image_store_lod_1D(i, x, layer, tmp);
178 }
179 
180 template <
181  typename T,
182  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
183 static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
184  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
185  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
186  auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
187  *data = __hipMapFrom<T>(tmp);
188 }
189 
190 template <
191  typename T,
192  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
193 static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
194  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
195  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
196  auto tmp = __hipMapTo<float4::Native_vec_>(data);
197  __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
198 }
199 
200 template <
201  typename T,
202  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
203 static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
204  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
205  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
206  auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
207  *data = __hipMapFrom<T>(tmp);
208 }
209 
210 template <
211  typename T,
212  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
213 static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
214  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
215  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
216  auto tmp = __hipMapTo<float4::Native_vec_>(data);
217  __ockl_image_store_CM(i, int2(x, y).data, face, tmp);
218 }
219 
220 template <
221  typename T,
222  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
223 static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
224  int layer) {
225  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
226  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
227  auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
228  *data = __hipMapFrom<T>(tmp);
229 }
230 
231 template <
232  typename T,
233  typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
234 static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
235  int layer) {
236  __HIP_SURFACE_OBJECT_PARAMETERS_INIT
237  x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
238  auto tmp = __hipMapTo<float4::Native_vec_>(data);
239  __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
240 }
241 
242 #endif
243 
244 #endif
Definition: amd_hip_vector_types.h:1813
Definition: amd_hip_vector_types.h:1820