From e91332340a37b8bf788ea08285498ba614d28a4d Mon Sep 17 00:00:00 2001 From: "Arandjelovic, Marko" Date: Fri, 15 Aug 2025 22:10:23 +0200 Subject: [PATCH] SWDEV-542815 - Fix unused-parameter compiler warnings (#719) * SWDEV-542815 - Fix unused-parameter compiler warnings * SWDEV-542815 - Run clang-format * SWDEV-542815 - Run clang-format * SWDEV-542815 - Run clang-format * SWDEV-542815 - Run clang-format [ROCm/clr commit: 1385d159ba72ea0d57a92e6dcde4545115a1187d] --- .../hip/amd_detail/amd_surface_functions.h | 17 +++-- .../hip/amd_detail/texture_fetch_functions.h | 18 ++++- .../amd_detail/texture_indirect_functions.h | 75 +++++++++++-------- 3 files changed, 69 insertions(+), 41 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h index 0a56f96410..c08bb2873b 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -114,10 +114,11 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x, int boundaryMode = hipBoundaryModeZero) { - __HIP_SURFACE_OBJECT_PARAMETERS_INIT - x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); - auto tmp = __ockl_image_load_1D(i, x); - *data = __hipMapFrom(tmp); + __HIP_SURFACE_OBJECT_PARAMETERS_INIT; + (void)boundaryMode; + x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); + auto tmp = __ockl_image_load_1D(i, x); + *data = __hipMapFrom(tmp); } /** \brief Writes the value data to the one-dimensional surface at coordinate x. @@ -157,7 +158,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su *data = __hipMapFrom(tmp); } -/** \brief Writes the value data to the two-dimensional surface at coordinate +/** \brief Writes the value data to the two-dimensional surface at coordinate * x, y. * * \tparam T The data type of the surface. @@ -177,7 +178,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su __ockl_image_store_2D(i, get_native_vector(coords), tmp); } -/** \brief Reads the value from the three-dimensional surface at coordinate +/** \brief Reads the value from the three-dimensional surface at coordinate * x, y, z. * * \tparam T The data type of the surface. @@ -238,7 +239,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje *data = __hipMapFrom(tmp); } -/** \brief Writes the value data to the one-dimensional layered surface at +/** \brief Writes the value data to the one-dimensional layered surface at * coordinate x and layer index. * * \tparam T The data type of the surface. @@ -257,7 +258,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje __ockl_image_store_lod_1D(i, x, layer, tmp); } -/** \brief Reads the value from the two-dimensional layered surface at +/** \brief Reads the value from the two-dimensional layered surface at * coordinate x, y and layer index. * * \tparam T The data type of the surface. diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h index 4123ea130b..c484416a8e 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h @@ -31,9 +31,10 @@ THE SOFTWARE. #include #endif // !defined(__HIPCC_RTC__) -#define TEXTURE_PARAMETERS_INIT \ - unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \ - unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; +#define TEXTURE_PARAMETERS_INIT \ + unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \ + unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ + (void)s; template struct __hip_is_tex_surf_scalar_channel_type @@ -376,6 +377,11 @@ template static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t texCubemapGrad(texture t, float x, float y, float z, float4 dPdx, float4 dPdy) { TEXTURE_PARAMETERS_INIT; + (void)x; + (void)y; + (void)z; + (void)dPdx; + (void)dPdy; // TODO missing in device libs. // auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)), // get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x, @@ -387,6 +393,12 @@ template static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t texCubemapLayeredGrad(texture t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy) { TEXTURE_PARAMETERS_INIT; + (void)x; + (void)y; + (void)z; + (void)layer; + (void)dPdx; + (void)dPdy; // TODO missing in device libs. // auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)), // get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x, diff --git a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h index 6d7a328b6b..54399f58b3 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h @@ -32,9 +32,10 @@ THE SOFTWARE. #include #endif // !defined(__HIPCC_RTC__) -#define TEXTURE_OBJECT_PARAMETERS_INIT \ - unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \ - unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; +#define TEXTURE_OBJECT_PARAMETERS_INIT \ + unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \ + unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ + (void)s; template < typename T, @@ -287,10 +288,11 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level) { - TEXTURE_OBJECT_PARAMETERS_INIT - float2 coords{x, layer}; - auto tmp = __ockl_image_sample_1Da(i, s, get_native_vector(coords)); - return __hipMapFrom(tmp); + TEXTURE_OBJECT_PARAMETERS_INIT; + (void)level; + float2 coords{x, layer}; + auto tmp = __ockl_image_sample_1Da(i, s, get_native_vector(coords)); + return __hipMapFrom(tmp); } template < @@ -306,10 +308,11 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level) { - TEXTURE_OBJECT_PARAMETERS_INIT - float4 coords{x, y, layer, 0.0f}; - auto tmp = __ockl_image_sample_2Da(i, s, get_native_vector(coords)); - return __hipMapFrom(tmp); + TEXTURE_OBJECT_PARAMETERS_INIT; + (void)level; + float4 coords{x, y, layer, 0.0f}; + auto tmp = __ockl_image_sample_2Da(i, s, get_native_vector(coords)); + return __hipMapFrom(tmp); } template < @@ -344,12 +347,17 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { - TEXTURE_OBJECT_PARAMETERS_INIT - // TODO missing in device libs. - // auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)), - // get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x, - // dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom(tmp); - return {}; + TEXTURE_OBJECT_PARAMETERS_INIT; + (void)x; + (void)y; + (void)z; + (void)dPdx; + (void)dPdy; + // TODO missing in device libs. + // auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)), + // get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x, + // dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom(tmp); + return {}; } template < @@ -422,13 +430,14 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy) { - TEXTURE_OBJECT_PARAMETERS_INIT - float4 coords{x, y, z, 0.0f}; - float4 gradx{dPdy.x, dPdy.y, dPdy.z, 0.0f}; - float4 grady{dPdy.x, dPdy.y, dPdy.z, 0.0f}; - auto tmp = __ockl_image_sample_grad_3D(i, s, get_native_vector(coords), - get_native_vector(gradx), get_native_vector(grady)); - return __hipMapFrom(tmp); + TEXTURE_OBJECT_PARAMETERS_INIT; + (void)dPdx; + float4 coords{x, y, z, 0.0f}; + float4 gradx{dPdy.x, dPdy.y, dPdy.z, 0.0f}; + float4 grady{dPdy.x, dPdy.y, dPdy.z, 0.0f}; + auto tmp = __ockl_image_sample_grad_3D(i, s, get_native_vector(coords), get_native_vector(gradx), + get_native_vector(grady)); + return __hipMapFrom(tmp); } template < @@ -483,12 +492,18 @@ template < typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy) { - TEXTURE_OBJECT_PARAMETERS_INIT - // TODO missing in device libs. - // auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)), - // get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x, - // dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom(tmp); - return {}; + TEXTURE_OBJECT_PARAMETERS_INIT; + (void)x; + (void)y; + (void)z; + (void)layer; + (void)dPdx; + (void)dPdy; + // TODO missing in device libs. + // auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)), + // get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x, + // dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom(tmp); + return {}; } template <