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/*

2Copyright (c) 2018 - 2025 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 <typename T, typename __hip_internal::enable_if<

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

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

115 int boundaryMode = hipBoundaryModeZero) {

116 __HIP_SURFACE_OBJECT_PARAMETERS_INIT;

117 (void)boundaryMode;

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 <typename T, typename __hip_internal::enable_if<

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

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

133 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

136 __ockl_image_store_1D(i, x, tmp);

137}

138

139

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

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

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

151 int y) {

152 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

154 int2 coords{x, y};

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

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

157}

158

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

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

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

171 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 int2 coords{x, y};

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

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

177}

178

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

190 __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,

192 int z) {

193 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

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

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

198}

199

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

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

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

213 int z) {

214 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

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

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

219}

220

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

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

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

233 int layer) {

234 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

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

238}

239

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

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

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

252 int layer) {

253 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

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

257}

258

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

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

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

272 int y, int layer) {

273 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

275 int2 coords{x, y};

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

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

278}

279

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

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

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

293 int y, int layer) {

294 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

296 int2 coords{x, y};

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

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

299}

300

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

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

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

314 int y, int face) {

315 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

317 int2 coords{x, y};

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

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

320}

321

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

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

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

335 int y, int face) {

336 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

338 int2 coords{x, y};

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

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

341}

342

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

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

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

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

358 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

360 int2 coords{x, y};

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

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

363}

364

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

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

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

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

380 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

382 int2 coords{x, y};

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

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

385}

386

387// Doxygen end group SurfaceAPI

392#endif

393

394#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:56

hipSurfaceObject_t

struct __hip_surface * hipSurfaceObject_t

Definition surface_types.h:43