From 9ee9f04cdd617bf25d848529a075276a4e9eb5b9 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Wed, 30 Aug 2023 13:16:38 +0000 Subject: [PATCH] SWDEV-407358 - Address MIGraphx -Werror issue with unsafe buffer access Change-Id: If4189dfbf12b670af60782f09ff8cf9b9867c299 [ROCm/clr commit: 16fb3ca923255e325edf2f922df2236510cb32a9] --- .../hip/amd_detail/amd_hip_cooperative_groups.h | 10 +++++----- .../hipamd/include/hip/amd_detail/amd_warp_functions.h | 1 + 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 0c3d9fdf9f..589c8c950f 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -236,7 +236,7 @@ class thread_block : public thread_group { const bool pow2 = ((tile_size & (tile_size - 1)) == 0); // Invalid tile size, assert if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) { - __hip_assert(false && "invalid tile size"); + __hip_assert(false && "invalid tile size") } thread_group tiledGroup = thread_group(internal::cg_tiled_group, tile_size); @@ -291,7 +291,7 @@ class tiled_group : public thread_group { const bool pow2 = ((tile_size & (tile_size - 1)) == 0); if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) { - __hip_assert(false && "invalid tile size"); + __hip_assert(false && "invalid tile size") } if (size() <= tile_size) { @@ -515,7 +515,7 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { return (static_cast(this)->thread_rank()); } default: { - __hip_assert(false && "invalid cooperative group type"); + __hip_assert(false && "invalid cooperative group type") return -1; } } @@ -543,7 +543,7 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const { return (static_cast(this)->is_valid()); } default: { - __hip_assert(false && "invalid cooperative group type"); + __hip_assert(false && "invalid cooperative group type") return false; } } @@ -576,7 +576,7 @@ __CG_QUALIFIER__ void thread_group::sync() const { break; } default: { - __hip_assert(false && "invalid cooperative group type"); + __hip_assert(false && "invalid cooperative group type") } } } diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h index 1e9eba8920..c43b81f969 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -31,6 +31,7 @@ THE SOFTWARE. #pragma clang diagnostic ignored "-Wold-style-cast" #pragma clang diagnostic ignored "-Wc++98-compat" #pragma clang diagnostic ignored "-Wc++98-compat-pedantic" +#pragma clang diagnostic ignored "-Wunsafe-buffer-usage" #endif __device__ static inline unsigned __hip_ds_bpermute(int index, unsigned src) {