clr/hipamd/include/hip/amd_detail/amd_surface_functions.h Source File

clr/hipamd/include/hip/amd_detail/amd_surface_functions.h Source File#

HIP Runtime API Reference: clr/hipamd/include/hip/amd_detail/amd_surface_functions.h Source File
amd_surface_functions.h
Go to the documentation of this file.
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
49// CUDA is using byte address, need map to pixel address for HIP
50static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) {
51 /*
52 * use below format index to generate format LUT
53 typedef enum {
54 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0,
55 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1,
56 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2,
57 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3,
58 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4,
59 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5,
60 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6,
61 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7,
62 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8,
63 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9,
64 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10,
65 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11,
66 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12,
67 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13,
68 HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14,
69 HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15
70 } hsa_ext_image_channel_type_t;
71 */
72 static const int FormatLUT[] = { 0, 1, 0, 1, 3, 1, 1, 1, 0, 1, 2, 0, 1, 2, 1, 2 };
73 x = FormatLUT[format] == 3 ? x / FormatLUT[format] : x >> FormatLUT[format];
74
75 /*
76 * use below order index to generate order LUT
77 typedef enum {
78 HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0,
79 HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1,
80 HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2,
81 HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3,
82 HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4,
83 HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5,
84 HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6,
85 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7,
86 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8,
87 HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9,
88 HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10,
89 HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11,
90 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12,
91 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13,
92 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14,
93 HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15,
94 HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16,
95 HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17,
96 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18,
97 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19
98 } hsa_ext_image_channel_order_t;
99 */
100 static const int OrderLUT[] = { 0, 0, 1, 1, 3, 1, 3, 2, 2, 2, 2, 2, 3, 2, 2, 2, 0, 0, 0, 0 };
101 return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order];
102}
103
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 surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
116 int boundaryMode = hipBoundaryModeZero) {
117 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
118 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
119 auto tmp = __ockl_image_load_1D(i, x);
120 *data = __hipMapFrom<T>(tmp);
121}
122
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 surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
134 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
135 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
136 auto tmp = __hipMapTo<float4::Native_vec_>(data);
137 __ockl_image_store_1D(i, x, tmp);
138}
139
140
149template <
150 typename T,
151 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
152static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
153 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
154 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
155 auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
156 *data = __hipMapFrom<T>(tmp);
157}
158
168template <
169 typename T,
170 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
171static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
172 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
173 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
174 auto tmp = __hipMapTo<float4::Native_vec_>(data);
175 __ockl_image_store_2D(i, int2(x, y).data, tmp);
176}
177
188template <
189 typename T,
190 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
191static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
192 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
193 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
194 auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
195 *data = __hipMapFrom<T>(tmp);
196}
197
208template <
209 typename T,
210 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
211static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
212 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
213 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
214 auto tmp = __hipMapTo<float4::Native_vec_>(data);
215 __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);
216}
217
227template <
228 typename T,
229 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
230static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
231 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
232 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
233 auto tmp = __ockl_image_load_lod_1D(i, x, layer);
234 *data = __hipMapFrom<T>(tmp);
235}
236
246template <
247 typename T,
248 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
249static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
250 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
251 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
252 auto tmp = __hipMapTo<float4::Native_vec_>(data);
253 __ockl_image_store_lod_1D(i, x, layer, tmp);
254}
255
266template <
267 typename T,
268 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
269static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
270 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
271 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
272 auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
273 *data = __hipMapFrom<T>(tmp);
274}
275
286template <
287 typename T,
288 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
289static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
290 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
291 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
292 auto tmp = __hipMapTo<float4::Native_vec_>(data);
293 __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);
294}
295
306template <
307 typename T,
308 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
309static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
310 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
311 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
312 auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
313 *data = __hipMapFrom<T>(tmp);
314}
315
326template <
327 typename T,
328 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
329static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
330 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
331 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
332 auto tmp = __hipMapTo<float4::Native_vec_>(data);
333 __ockl_image_store_CM(i, int2(x, y).data, face, tmp);
334}
335
347template <
348 typename T,
349 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
350static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
351 int layer) {
352 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
353 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
354 auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
355 *data = __hipMapFrom<T>(tmp);
356}
357
369template <
370 typename T,
371 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
372static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
373 int layer) {
374 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
375 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
376 auto tmp = __hipMapTo<float4::Native_vec_>(data);
377 __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);
378}
379
380// Doxygen end group SurfaceAPI
385#endif
386
387#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