/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/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 - 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

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

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

46 /*

47 * use below format index to generate format LUT

48 typedef enum {

49 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 = 0,

50 HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 = 1,

51 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 = 2,

52 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 = 3,

53 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT24 = 4,

54 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555 = 5,

55 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565 = 6,

56 HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_SHORT_101010 = 7,

57 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8 = 8,

58 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16 = 9,

59 HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32 = 10,

60 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8 = 11,

61 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16 = 12,

62 HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32 = 13,

63 HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT = 14,

64 HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT = 15

65 } hsa_ext_image_channel_type_t;

66 */

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

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

69

70 /*

71 * use below order index to generate order LUT

72 typedef enum {

73 HSA_EXT_IMAGE_CHANNEL_ORDER_A = 0,

74 HSA_EXT_IMAGE_CHANNEL_ORDER_R = 1,

75 HSA_EXT_IMAGE_CHANNEL_ORDER_RX = 2,

76 HSA_EXT_IMAGE_CHANNEL_ORDER_RG = 3,

77 HSA_EXT_IMAGE_CHANNEL_ORDER_RGX = 4,

78 HSA_EXT_IMAGE_CHANNEL_ORDER_RA = 5,

79 HSA_EXT_IMAGE_CHANNEL_ORDER_RGB = 6,

80 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBX = 7,

81 HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA = 8,

82 HSA_EXT_IMAGE_CHANNEL_ORDER_BGRA = 9,

83 HSA_EXT_IMAGE_CHANNEL_ORDER_ARGB = 10,

84 HSA_EXT_IMAGE_CHANNEL_ORDER_ABGR = 11,

85 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGB = 12,

86 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBX = 13,

87 HSA_EXT_IMAGE_CHANNEL_ORDER_SRGBA = 14,

88 HSA_EXT_IMAGE_CHANNEL_ORDER_SBGRA = 15,

89 HSA_EXT_IMAGE_CHANNEL_ORDER_INTENSITY = 16,

90 HSA_EXT_IMAGE_CHANNEL_ORDER_LUMINANCE = 17,

91 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH = 18,

92 HSA_EXT_IMAGE_CHANNEL_ORDER_DEPTH_STENCIL = 19

93 } hsa_ext_image_channel_order_t;

94 */

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

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

97}

98

107template <

108 typename T,

109 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

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

111 int boundaryMode = hipBoundaryModeZero) {

112 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

114 auto tmp = __ockl_image_load_1D(i, x);

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

116}

117

125template <

126 typename T,

127 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

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

129 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

132 __ockl_image_store_1D(i, x, tmp);

133}

134

135

144template <

145 typename T,

146 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

147static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {

148 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

150 auto tmp = __ockl_image_load_2D(i, int2(x, y).data);

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

152}

153

163template <

164 typename T,

165 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

166static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {

167 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

170 __ockl_image_store_2D(i, int2(x, y).data, tmp);

171}

172

183template <

184 typename T,

185 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

186static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {

187 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

189 auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);

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

191}

192

203template <

204 typename T,

205 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

206static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {

207 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

210 __ockl_image_store_3D(i, int4(x, y, z, 0).data, tmp);

211}

212

222template <

223 typename T,

224 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

225static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {

226 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

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

230}

231

241template <

242 typename T,

243 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

244static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {

245 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

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

249}

250

261template <

262 typename T,

263 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

264static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {

265 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

267 auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);

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

269}

270

281template <

282 typename T,

283 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

284static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {

285 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

288 __ockl_image_store_lod_2D(i, int2(x, y).data, layer, tmp);

289}

290

301template <

302 typename T,

303 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

304static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {

305 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

307 auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);

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

309}

310

321template <

322 typename T,

323 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

324static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {

325 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

328 __ockl_image_store_CM(i, int2(x, y).data, face, tmp);

329}

330

342template <

343 typename T,

344 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

345static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,

346 int layer) {

347 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

349 auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);

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

351}

352

364template <

365 typename T,

366 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>

367static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,

368 int layer) {

369 __HIP_SURFACE_OBJECT_PARAMETERS_INIT

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

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

372 __ockl_image_store_lod_CM(i, int2(x, y).data, face, layer, tmp);

373}

374

375#endif

376

377#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