HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_surface_functions.h
1/*
2Copyright (c) 2018 - 2022 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#include <hip/surface_types.h>
29#include <hip/hip_vector_types.h>
30#include <hip/amd_detail/texture_fetch_functions.h>
31#include <hip/amd_detail/ockl_image.h>
32
33#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
34 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
35
36// CUDA is using byte address, need map to pixel address for HIP
37static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) {
38 /*
39 * use below format index to generate format LUT
40 typedef enum {
41 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0,
42 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1,
43 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2,
44 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3,
45 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4,
46 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5,
47 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6,
48 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7,
49 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8,
50 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9,
51 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10,
52 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11,
53 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12,
54 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13,
55 HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14,
56 HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15
57 } hsa_ext_image_channel_type_t;
58 */
59 static const int FormatLUT[] = { 0, 1, 0, 1, 3, 1, 1, 1, 0, 1, 2, 0, 1, 2, 1, 2 };
60 x = FormatLUT[format] == 3 ? x / FormatLUT[format] : x >> FormatLUT[format];
61
62 /*
63 * use below order index to generate order LUT
64 typedef enum {
65 HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0,
66 HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1,
67 HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2,
68 HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3,
69 HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4,
70 HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5,
71 HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6,
72 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7,
73 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8,
74 HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9,
75 HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10,
76 HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11,
77 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12,
78 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13,
79 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14,
80 HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15,
81 HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16,
82 HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17,
83 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18,
84 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19
85 } hsa_ext_image_channel_order_t;
86 */
87 static const int OrderLUT[] = { 0, 0, 1, 1, 3, 1, 3, 2, 2, 2, 2, 2, 3, 2, 2, 2, 0, 0, 0, 0 };
88 return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order];
89}
90
91template <
92 typename T,
93 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
94static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
95 int boundaryMode = hipBoundaryModeZero) {
96 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
97 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
98 auto tmp = __ockl_image_load_1D(i, x);
99 *data = __hipMapFrom<T>(tmp);
100}
101
102template <
103 typename T,
104 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
105static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
106 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
107 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
108 auto tmp = __hipMapTo<float4::Native_vec_>(data);
109 __ockl_image_store_1D(i, x, tmp);
110}
111
112template <
113 typename T,
114 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
115static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
116 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
117 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
118 auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
119 *data = __hipMapFrom<T>(tmp);
120}
121
122template <
123 typename T,
124 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
125static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
126 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
127 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
128 auto tmp = __hipMapTo<float4::Native_vec_>(data);
129 __ockl_image_store_2D(i, int2(x, y).data, tmp);
130}
131
132template <
133 typename T,
134 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
135static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
136 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
137 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
138 auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
139 *data = __hipMapFrom<T>(tmp);
140}
141
142template <
143 typename T,
144 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
145static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
146 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
147 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
148 auto tmp = __hipMapTo<float4::Native_vec_>(data);
149 __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
150}
151
152template <
153 typename T,
154 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
155static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
156 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
157 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
158 auto tmp = __ockl_image_load_lod_1D(i, x, layer);
159 *data = __hipMapFrom<T>(tmp);
160}
161
162template <
163 typename T,
164 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
165static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
166 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
167 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
168 auto tmp = __hipMapTo<float4::Native_vec_>(data);
169 __ockl_image_store_lod_1D(i, x, layer, tmp);
170}
171
172template <
173 typename T,
174 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
175static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
176 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
177 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
178 auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
179 *data = __hipMapFrom<T>(tmp);
180}
181
182template <
183 typename T,
184 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
185static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
186 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
187 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
188 auto tmp = __hipMapTo<float4::Native_vec_>(data);
189 __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
190}
191
192template <
193 typename T,
194 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
195static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
196 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
197 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
198 auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
199 *data = __hipMapFrom<T>(tmp);
200}
201
202template <
203 typename T,
204 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
205static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
206 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
207 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
208 auto tmp = __hipMapTo<float4::Native_vec_>(data);
209 __ockl_image_store_CM(i, int2(x, y).data, face, tmp);
210}
211
212template <
213 typename T,
214 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
215static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
216 int layer) {
217 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
218 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
219 auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
220 *data = __hipMapFrom<T>(tmp);
221}
222
223template <
224 typename T,
225 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
226static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
227 int layer) {
228 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
229 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
230 auto tmp = __hipMapTo<float4::Native_vec_>(data);
231 __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
232}
233
234#endif
235#endif
Definition amd_hip_vector_types.h:1764
Definition amd_hip_vector_types.h:1771