From 10ea7449965dbb0be6b4692c24c58d78c312e245 Mon Sep 17 00:00:00 2001 From: "Sang, Tao" Date: Tue, 10 Jun 2025 22:30:42 -0400 Subject: [PATCH] SWDEV-511020 - fix bugs in surface function header (#449) --- .../hip/amd_detail/amd_surface_functions.h | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) 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); }