From 9eb45dd1e01c9f15246e8a6fce346dc6c8b7ad91 Mon Sep 17 00:00:00 2001 From: Ajay Date: Wed, 6 Jul 2022 22:03:12 +0000 Subject: [PATCH] SWDEV-345188 - hipamd: warnings treated as errors for catch tests Change-Id: I8981a530da2e0213c60f7a81b09d9c36e27e9025 --- .../hip/amd_detail/amd_device_functions.h | 14 +++---- .../hip/amd_detail/amd_surface_functions.h | 37 +++++++------------ 2 files changed, 18 insertions(+), 33 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h index 8383a2d720..44610d4679 100644 --- a/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -101,7 +101,7 @@ __device__ static int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset) int32_t total = 0; for (int i = 0x20; i > 0; i >>= 1) { uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); - uint32_t pcnt = __builtin_popcountll(temp_mask_lo); + int32_t pcnt = __builtin_popcountll(temp_mask_lo); if (pcnt < temp_offset) { temp_mask = temp_mask >> i; temp_offset -= pcnt; @@ -135,7 +135,7 @@ __device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset) int32_t total = 0; for (int i = 0x20; i > 0; i >>= 1) { uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); - uint32_t pcnt = __builtin_popcountll(temp_mask_lo); + int32_t pcnt = __builtin_popcountll(temp_mask_lo); if (pcnt < temp_offset) { temp_mask = temp_mask >> i; temp_offset -= pcnt; @@ -653,7 +653,7 @@ __device__ long long int __clock(); __device__ long long int clock64(); __device__ long long int clock(); // hip.amdgcn.bc - named sync -__device__ void __named_sync(int a, int b); +__device__ void __named_sync(); #ifdef __HIP_DEVICE_COMPILE__ @@ -700,7 +700,7 @@ long long int clock() { return __clock(); } // hip.amdgcn.bc - named sync __device__ inline -void __named_sync(int a, int b) { __builtin_amdgcn_s_barrier(); } +void __named_sync() { __builtin_amdgcn_s_barrier(); } #endif // __HIP_DEVICE_COMPILE__ @@ -871,11 +871,7 @@ void __assert_fail(const char *assertion, } extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) -void __assertfail(const char *assertion, - const char *file, - unsigned int line, - const char *function, - size_t charsize) +void __assertfail() { // ignore all the args for now. __builtin_trap(); diff --git a/hipamd/include/hip/amd_detail/amd_surface_functions.h b/hipamd/include/hip/amd_detail/amd_surface_functions.h index 4a8a6a63c8..bae218ee14 100644 --- a/hipamd/include/hip/amd_detail/amd_surface_functions.h +++ b/hipamd/include/hip/amd_detail/amd_surface_functions.h @@ -229,8 +229,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); auto tmp = __hipMapToNativeFloat4(data); @@ -240,8 +239,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __ockl_image_load_2D(i, int2(x, y).data); @@ -251,8 +249,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __hipMapToNativeFloat4(data); @@ -262,8 +259,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data); @@ -273,8 +269,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); auto tmp = __hipMapToNativeFloat4(data); @@ -284,8 +279,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) { __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_lod_1D(i, x, layer); @@ -295,8 +289,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); auto tmp = __hipMapToNativeFloat4(data); @@ -306,8 +299,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer); @@ -317,8 +309,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __hipMapToNativeFloat4(data); @@ -328,8 +319,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face); @@ -339,8 +329,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> -static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face, - int boundaryMode = hipBoundaryModeZero) { +static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __hipMapToNativeFloat4(data); @@ -351,7 +340,7 @@ template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, - int layer, int boundaryMode = hipBoundaryModeZero) { + int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer); @@ -362,7 +351,7 @@ template < typename T, typename std::enable_if<__hip_is_isurf_channel_type::value>::type* = nullptr> static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face, - int layer, int boundaryMode = hipBoundaryModeZero) { + int layer) { __HIP_SURFACE_OBJECT_PARAMETERS_INIT x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); auto tmp = __hipMapToNativeFloat4(data);