HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_surface_functions.h
1/*
2Copyright (c) 2018 - 2023 Advanced Micro Devices, Inc. All rights reserved.
3
4Permission is hereby granted, free of charge, to any person obtaining a copy
5of this software and associated documentation files (the "Software"), to deal
6in the Software without restriction, including without limitation the rights
7to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8copies of the Software, and to permit persons to whom the Software is
9furnished to do so, subject to the following conditions:
10
11The above copyright notice and this permission notice shall be included in
12all copies or substantial portions of the Software.
13
14THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20THE 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
45static __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
99template <
100 typename T,
101 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
102static __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
110template <
111 typename T,
112 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
113static __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
120template <
121 typename T,
122 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
123static __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
130template <
131 typename T,
132 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
133static __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
140template <
141 typename T,
142 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
143static __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
150template <
151 typename T,
152 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
153static __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
160template <
161 typename T,
162 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
163static __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
170template <
171 typename T,
172 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
173static __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
180template <
181 typename T,
182 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
183static __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
190template <
191 typename T,
192 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
193static __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
200template <
201 typename T,
202 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
203static __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
210template <
211 typename T,
212 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
213static __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
220template <
221 typename T,
222 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
223static __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
231template <
232 typename T,
233 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
234static __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