23#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
24#define HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
26#if defined(__cplusplus)
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>
35#if defined(__HIPCC_RTC__)
36#define __HOST_DEVICE__ __device__
38#define __HOST_DEVICE__ __host__ __device__
41#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
42 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
45static __HOST_DEVICE__ __forceinline__
int __hipGetPixelAddr(
int x,
int format,
int order) {
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];
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];
101 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
102static __device__ __hip_img_chk__
void surf1Dread(T* data, hipSurfaceObject_t surfObj,
int x,
103 int boundaryMode = hipBoundaryModeZero) {
104 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
105 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
106 auto tmp = __ockl_image_load_1D(i, x);
107 *data = __hipMapFrom<T>(tmp);
112 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
113static __device__ __hip_img_chk__
void surf1Dwrite(T data, hipSurfaceObject_t surfObj,
int x) {
114 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
115 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
116 auto tmp = __hipMapTo<float4::Native_vec_>(data);
117 __ockl_image_store_1D(i, x, tmp);
122 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
123static __device__ __hip_img_chk__
void surf2Dread(T* data, hipSurfaceObject_t surfObj,
int x,
int y) {
124 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
125 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
126 auto tmp = __ockl_image_load_2D(i,
int2(x, y).data);
127 *data = __hipMapFrom<T>(tmp);
132 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
133static __device__ __hip_img_chk__
void surf2Dwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y) {
134 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
135 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
136 auto tmp = __hipMapTo<float4::Native_vec_>(data);
137 __ockl_image_store_2D(i,
int2(x, y).data, tmp);
142 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
143static __device__ __hip_img_chk__
void surf3Dread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int z) {
144 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
145 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
146 auto tmp = __ockl_image_load_3D(i,
int4(x, y, z, 0).data);
147 *data = __hipMapFrom<T>(tmp);
152 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
153static __device__ __hip_img_chk__
void surf3Dwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y,
int z) {
154 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
155 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
156 auto tmp = __hipMapTo<float4::Native_vec_>(data);
157 __ockl_image_store_3D(i,
int4(x, y, z, 0).data, tmp);
162 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
163static __device__ __hip_img_chk__
void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj,
int x,
int layer) {
164 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
165 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
166 auto tmp = __ockl_image_load_lod_1D(i, x, layer);
167 *data = __hipMapFrom<T>(tmp);
172 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
173static __device__ __hip_img_chk__
void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj,
int x,
int layer) {
174 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
175 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
176 auto tmp = __hipMapTo<float4::Native_vec_>(data);
177 __ockl_image_store_lod_1D(i, x, layer, tmp);
182 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
183static __device__ __hip_img_chk__
void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int layer) {
184 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
185 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
186 auto tmp = __ockl_image_load_lod_2D(i,
int2(x, y).data, layer);
187 *data = __hipMapFrom<T>(tmp);
192 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
193static __device__ __hip_img_chk__
void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y,
int layer) {
194 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
195 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
196 auto tmp = __hipMapTo<float4::Native_vec_>(data);
197 __ockl_image_store_lod_2D(i,
int2(x, y).data, layer, tmp);
202 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
203static __device__ __hip_img_chk__
void surfCubemapread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int face) {
204 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
205 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
206 auto tmp = __ockl_image_load_CM(i,
int2(x, y).data, face);
207 *data = __hipMapFrom<T>(tmp);
212 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
213static __device__ __hip_img_chk__
void surfCubemapwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y,
int face) {
214 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
215 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
216 auto tmp = __hipMapTo<float4::Native_vec_>(data);
217 __ockl_image_store_CM(i,
int2(x, y).data, face, tmp);
222 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
223static __device__ __hip_img_chk__
void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int face,
225 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
226 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
227 auto tmp = __ockl_image_load_lod_CM(i,
int2(x, y).data, face, layer);
228 *data = __hipMapFrom<T>(tmp);
233 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
234static __device__ __hip_img_chk__
void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int face,
236 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
237 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
238 auto tmp = __hipMapTo<float4::Native_vec_>(data);
239 __ockl_image_store_lod_CM(i,
int2(x, y).data, face, layer, tmp);
Definition amd_hip_vector_types.h:1764
Definition amd_hip_vector_types.h:1771