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) 2018 - 2025 Advanced Micro Devices, Inc. All rights reserved.

3 

4 Permission is hereby granted, free of charge, to any person obtaining a copy

5 of this software and associated documentation files (the "Software"), to deal

6 in the Software without restriction, including without limitation the rights

7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell

8 copies of the Software, and to permit persons to whom the Software is

9 furnished to do so, subject to the following conditions:

10 

11 The above copyright notice and this permission notice shall be included in

12 all copies or substantial portions of the Software.

13 

14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR

15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,

16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE

17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER

18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,

19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN

20 THE 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

50 static __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 

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

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

114 static __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 

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

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

132 static __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 

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

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

150 static __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 

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

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

170 static __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 

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

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

191 static __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 

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

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

212 static __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 

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

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

232 static __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 

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

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

251 static __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 

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

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

271 static __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 

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

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

292 static __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 

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

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

313 static __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 

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

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

334 static __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 

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

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

356 static __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 

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

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

378 static __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