AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-511020 - fix bugs in surface function header (#449)

Browse Source
This commit is contained in:
Sang, Tao
2025-06-10 22:30:42 -04:00
committed by GitHub
parent 5fabfedbd1
commit 10ea744996
1 changed files with 10 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_surface_functions.h
+10 -10
View File
@@ -151,8 +151,8 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::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<T>(tmp);
}
@@ -171,8 +171,8 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::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<float4::Native_vec_>(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<T>::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<T>(tmp);
}
@@ -213,8 +213,8 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::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<float4::Native_vec_>(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<T>::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<T>(tmp);
}
@@ -293,8 +293,8 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::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<float4::Native_vec_>(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<T>::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<T>(tmp);
}
@@ -335,8 +335,8 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::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<float4::Native_vec_>(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<T>(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<float4::Native_vec_>(data);
__ockl_image_store_lod_CM(i, get_native_vector(coords), face, layer, tmp);
}