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#include <hip/surface_types.h>
29#include <hip/hip_vector_types.h>
30#include <hip/amd_detail/texture_fetch_functions.h>
31#include <hip/amd_detail/ockl_image.h>
33#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
34 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
37static __HOST_DEVICE__ __forceinline__
int __hipGetPixelAddr(
int x,
int format,
int order) {
59 static const int FormatLUT[] = { 0, 1, 0, 1, 3, 1, 1, 1, 0, 1, 2, 0, 1, 2, 1, 2 };
60 x = FormatLUT[format] == 3 ? x / FormatLUT[format] : x >> FormatLUT[format];
87 static const int OrderLUT[] = { 0, 0, 1, 1, 3, 1, 3, 2, 2, 2, 2, 2, 3, 2, 2, 2, 0, 0, 0, 0 };
88 return x = OrderLUT[order] == 3 ? x / OrderLUT[order] : x >> OrderLUT[order];
93 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
94static __device__ __hip_img_chk__
void surf1Dread(T* data, hipSurfaceObject_t surfObj,
int x,
95 int boundaryMode = hipBoundaryModeZero) {
96 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
97 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
98 auto tmp = __ockl_image_load_1D(i, x);
99 *data = __hipMapFrom<T>(tmp);
104 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
105static __device__ __hip_img_chk__
void surf1Dwrite(T data, hipSurfaceObject_t surfObj,
int x) {
106 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
107 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
108 auto tmp = __hipMapTo<float4::Native_vec_>(data);
109 __ockl_image_store_1D(i, x, tmp);
114 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
115static __device__ __hip_img_chk__
void surf2Dread(T* data, hipSurfaceObject_t surfObj,
int x,
int y) {
116 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
117 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
118 auto tmp = __ockl_image_load_2D(i,
int2(x, y).data);
119 *data = __hipMapFrom<T>(tmp);
124 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
125static __device__ __hip_img_chk__
void surf2Dwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y) {
126 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
127 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
128 auto tmp = __hipMapTo<float4::Native_vec_>(data);
129 __ockl_image_store_2D(i,
int2(x, y).data, tmp);
134 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
135static __device__ __hip_img_chk__
void surf3Dread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int z) {
136 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
137 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
138 auto tmp = __ockl_image_load_3D(i,
int4(x, y, z, 0).data);
139 *data = __hipMapFrom<T>(tmp);
144 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
145static __device__ __hip_img_chk__
void surf3Dwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y,
int z) {
146 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
147 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
148 auto tmp = __hipMapTo<float4::Native_vec_>(data);
149 __ockl_image_store_3D(i,
int4(x, y, z, 0).data, tmp);
154 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
155static __device__ __hip_img_chk__
void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj,
int x,
int layer) {
156 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
157 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
158 auto tmp = __ockl_image_load_lod_1D(i, x, layer);
159 *data = __hipMapFrom<T>(tmp);
164 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
165static __device__ __hip_img_chk__
void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj,
int x,
int layer) {
166 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
167 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
168 auto tmp = __hipMapTo<float4::Native_vec_>(data);
169 __ockl_image_store_lod_1D(i, x, layer, tmp);
174 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
175static __device__ __hip_img_chk__
void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int layer) {
176 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
177 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
178 auto tmp = __ockl_image_load_lod_2D(i,
int2(x, y).data, layer);
179 *data = __hipMapFrom<T>(tmp);
184 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
185static __device__ __hip_img_chk__
void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y,
int layer) {
186 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
187 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
188 auto tmp = __hipMapTo<float4::Native_vec_>(data);
189 __ockl_image_store_lod_2D(i,
int2(x, y).data, layer, tmp);
194 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
195static __device__ __hip_img_chk__
void surfCubemapread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int face) {
196 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
197 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
198 auto tmp = __ockl_image_load_CM(i,
int2(x, y).data, face);
199 *data = __hipMapFrom<T>(tmp);
204 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
205static __device__ __hip_img_chk__
void surfCubemapwrite(T data, hipSurfaceObject_t surfObj,
int x,
int y,
int face) {
206 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
207 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
208 auto tmp = __hipMapTo<float4::Native_vec_>(data);
209 __ockl_image_store_CM(i,
int2(x, y).data, face, tmp);
214 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
215static __device__ __hip_img_chk__
void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int face,
217 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
218 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
219 auto tmp = __ockl_image_load_lod_CM(i,
int2(x, y).data, face, layer);
220 *data = __hipMapFrom<T>(tmp);
225 typename std::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* =
nullptr>
226static __device__ __hip_img_chk__
void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj,
int x,
int y,
int face,
228 __HIP_SURFACE_OBJECT_PARAMETERS_INIT
229 x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
230 auto tmp = __hipMapTo<float4::Native_vec_>(data);
231 __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