diff --git a/hipamd/include/hip/amd_detail/amd_surface_functions.h b/hipamd/include/hip/amd_detail/amd_surface_functions.h index 3f62011945..0a56f96410 100644 --- a/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -151,8 +151,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __ockl_image_load_2D(i, get_native_vector(coords)); *data = __hipMapFrom(tmp); } @@ -171,8 +171,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __hipMapTo(data); __ockl_image_store_2D(i, get_native_vector(coords), tmp); } @@ -192,8 +192,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int4 coords{x, y, z, 0}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); + int4 coords{x, y, z, 0}; auto tmp = __ockl_image_load_3D(i, get_native_vector(coords)); *data = __hipMapFrom(tmp); } @@ -213,8 +213,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int4 coords{x, y, z, 0}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); + int4 coords{x, y, z, 0}; auto tmp = __hipMapTo(data); __ockl_image_store_3D(i, get_native_vector(coords), tmp); } @@ -272,8 +272,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __ockl_image_load_lod_2D(i, get_native_vector(coords), layer); *data = __hipMapFrom(tmp); } @@ -293,8 +293,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __hipMapTo(data); __ockl_image_store_lod_2D(i, get_native_vector(coords), layer, tmp); } @@ -314,8 +314,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __ockl_image_load_CM(i, get_native_vector(coords), face); *data = __hipMapFrom(tmp); } @@ -335,8 +335,8 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __hipMapTo(data); __ockl_image_store_CM(i, get_native_vector(coords), face, tmp); } @@ -358,8 +358,8 @@ template < static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __ockl_image_load_lod_CM(i, get_native_vector(coords), face, layer); *data = __hipMapFrom(tmp); } @@ -381,8 +381,8 @@ template < static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT - int2 coords{x, y}; x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); + int2 coords{x, y}; auto tmp = __hipMapTo(data); __ockl_image_store_lod_CM(i, get_native_vector(coords), face, layer, tmp); }