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

• clr
• hipamd
• include
• hip
• amd_detail

amd_surface_functions.h

Go to the documentation of this file.

1/*

2 * Copyright (c) Advanced Micro Devices, Inc., or its affiliates.

3 *

4 * SPDX-License-Identifier: MIT

5 */

6

7#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H

8#define HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H

9

10#if defined(__cplusplus)

11

12#if !defined(__HIPCC_RTC__)

13#include <hip/surface_types.h>

14#include <hip/hip_vector_types.h>

15#include <hip/amd_detail/texture_fetch_functions.h>

16#include <hip/amd_detail/ockl_image.h>

17#endif

18

19#if defined(__HIPCC_RTC__)

20#define __HOST_DEVICE__ __device__

21#else

22#define __HOST_DEVICE__ __host__ __device__

23#endif

24

25#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \

26 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;

27

33// CUDA is using byte address, need map to pixel address for HIP

34static __HOST_DEVICE__ __forceinline__ int __hipGetPixelAddr(int x, int format, int order) {

35 /*

36 * use below format index to generate format LUT

37 typedef enum {

38 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0,

39 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1,

40 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2,

41 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3,

42 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4,

43 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5,

44 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6,

45 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7,

46 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8,

47 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9,

48 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10,

49 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11,

50 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12,

51 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13,

52 HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14,

53 HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15

54 } hsa_ext_image_channel_type_t;

55 */

56 static const int FormatLUT[] = {0, 1, 0, 1, 3, 1, 1, 1, 0, 1, 2, 0, 1, 2, 1, 2};

57 x = FormatLUT[format] == 3 ? x / FormatLUT[format] : x >> FormatLUT[format];

58

59 /*

60 * use below order index to generate order LUT

61 typedef enum {

62 HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0,

63 HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1,

64 HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2,

65 HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3,

66 HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4,

67 HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5,

68 HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6,

69 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7,

70 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8,

71 HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9,

72 HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10,

73 HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11,

74 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12,

75 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13,

76 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14,

77 HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15,

78 HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16,

79 HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17,

80 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18,

81 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19

82 } hsa_ext_image_channel_order_t;

83 */

84 static const int OrderLUT[] = {0, 0, 1, 1, 3, 1, 3, 2, 2, 2, 2, 2, 3, 2, 2, 2, 0, 0, 0, 0};

85 return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order];

86}

87

96template <typename T, typename __hip_internal::enable_if<

97 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

98static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,

99 int boundaryMode = hipBoundaryModeZero) {

100 __HIP_SURFACE_OBJECT_PARAMETERS_INIT;

101 (void)boundaryMode;

102 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));

103 auto tmp = __ockl_image_load_1D(i, x);

104 *data = __hipMapFrom<T>(tmp);

105}

106

114template <typename T, typename __hip_internal::enable_if<

115 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

116static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {

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 = __hipMapTo<float4::Native_vec_>(data);

120 __ockl_image_store_1D(i, x, tmp);

121}

122

123

132template <typename T, typename __hip_internal::enable_if<

133 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

134static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x,

135 int y) {

136 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

137 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

138 int2 coords{x, y};

139 auto tmp = __ockl_image_load_2D(i, get_native_vector(coords));

140 *data = __hipMapFrom<T>(tmp);

141}

142

152template <typename T, typename __hip_internal::enable_if<

153 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

154static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x,

155 int y) {

156 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

157 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

158 int2 coords{x, y};

159 auto tmp = __hipMapTo<float4::Native_vec_>(data);

160 __ockl_image_store_2D(i, get_native_vector(coords), tmp);

161}

162

173template <typename T, typename __hip_internal::enable_if<

174 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

175static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y,

176 int z) {

177 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

178 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));

179 int4 coords{x, y, z, 0};

180 auto tmp = __ockl_image_load_3D(i, get_native_vector(coords));

181 *data = __hipMapFrom<T>(tmp);

182}

183

194template <typename T, typename __hip_internal::enable_if<

195 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

196static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y,

197 int z) {

198 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

199 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));

200 int4 coords{x, y, z, 0};

201 auto tmp = __hipMapTo<float4::Native_vec_>(data);

202 __ockl_image_store_3D(i, get_native_vector(coords), tmp);

203}

204

214template <typename T, typename __hip_internal::enable_if<

215 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

216static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x,

217 int layer) {

218 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

219 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));

220 auto tmp = __ockl_image_load_lod_1D(i, x, layer);

221 *data = __hipMapFrom<T>(tmp);

222}

223

233template <typename T, typename __hip_internal::enable_if<

234 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

235static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x,

236 int layer) {

237 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

238 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));

239 auto tmp = __hipMapTo<float4::Native_vec_>(data);

240 __ockl_image_store_lod_1D(i, x, layer, tmp);

241}

242

253template <typename T, typename __hip_internal::enable_if<

254 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

255static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x,

256 int y, int layer) {

257 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

258 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

259 int2 coords{x, y};

260 auto tmp = __ockl_image_load_lod_2D(i, get_native_vector(coords), layer);

261 *data = __hipMapFrom<T>(tmp);

262}

263

274template <typename T, typename __hip_internal::enable_if<

275 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

276static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x,

277 int y, int layer) {

278 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

279 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

280 int2 coords{x, y};

281 auto tmp = __hipMapTo<float4::Native_vec_>(data);

282 __ockl_image_store_lod_2D(i, get_native_vector(coords), layer, tmp);

283}

284

295template <typename T, typename __hip_internal::enable_if<

296 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

297static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x,

298 int y, int face) {

299 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

300 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

301 int2 coords{x, y};

302 auto tmp = __ockl_image_load_CM(i, get_native_vector(coords), face);

303 *data = __hipMapFrom<T>(tmp);

304}

305

316template <typename T, typename __hip_internal::enable_if<

317 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

318static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x,

319 int y, int face) {

320 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

321 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

322 int2 coords{x, y};

323 auto tmp = __hipMapTo<float4::Native_vec_>(data);

324 __ockl_image_store_CM(i, get_native_vector(coords), face, tmp);

325}

326

338template <typename T, typename __hip_internal::enable_if<

339 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

340static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj,

341 int x, int y, int face, int layer) {

342 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

343 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

344 int2 coords{x, y};

345 auto tmp = __ockl_image_load_lod_CM(i, get_native_vector(coords), face, layer);

346 *data = __hipMapFrom<T>(tmp);

347}

348

360template <typename T, typename __hip_internal::enable_if<

361 __hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

362static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj,

363 int x, int y, int face, int layer) {

364 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

365 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));

366 int2 coords{x, y};

367 auto tmp = __hipMapTo<float4::Native_vec_>(data);

368 __ockl_image_store_lod_CM(i, get_native_vector(coords), face, layer, tmp);

369}

370

371// Doxygen end group SurfaceAPI

376#endif

377

378#endif

__HIP_SURFACE_OBJECT_PARAMETERS_INIT

#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT

Definition amd_surface_functions.h:25

__HOST_DEVICE__

#define __HOST_DEVICE__

Definition amd_surface_functions.h:22

hip_vector_types.h

surface_types.h

Defines surface types for HIP runtime.

hipBoundaryModeZero

@ hipBoundaryModeZero

Definition surface_types.h:41

hipSurfaceObject_t

struct __hip_surface * hipSurfaceObject_t

Definition surface_types.h:28