From d96114fd22ddce03823713415b9202398bb9674e Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Tue, 20 Dec 2022 11:56:43 +0000 Subject: [PATCH] SWDEV-372396 - Address the errors in hiprtc headers with -Werror Change-Id: Ied553216dfc2ff8b5c72de617cb47b6ca8f52063 [ROCm/clr commit: 0ccd41050b0c434fcd595a5f865350781b29dda1] --- .../hip/amd_detail/amd_device_functions.h | 6 ++- .../amd_detail/amd_hip_cooperative_groups.h | 32 +++++++------ .../hip/amd_detail/amd_hip_unsafe_atomics.h | 6 ++- .../hip/amd_detail/amd_warp_functions.h | 11 ++++- .../hip_cooperative_groups_helper.h | 46 ++++++++++++------- .../clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake | 9 +++- .../clr/hipamd/src/hiprtc/hiprtcInternal.cpp | 4 +- 7 files changed, 78 insertions(+), 36 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index ce421c63d8..57403eb390 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -444,7 +444,8 @@ __device__ static inline unsigned long long int __double2ull_ru(double x) { __device__ static inline unsigned long long int __double2ull_rz(double x) { return (unsigned long long int)x; } - +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wc++98-compat-pedantic" __device__ static inline long long int __double_as_longlong(double x) { static_assert(sizeof(long long) == sizeof(double), ""); @@ -453,6 +454,7 @@ __device__ static inline long long int __double_as_longlong(double x) { return tmp; } +#pragma clang diagnostic pop /* __device__ unsigned short __float2half_rn(float x); 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 747f65a490..575a9f8ef7 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 @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -32,6 +32,13 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wc++98-compat" +#pragma clang diagnostic ignored "-Wsign-conversion" +#pragma clang diagnostic ignored "-Wunused-parameter" +#pragma clang diagnostic ignored "-Wreserved-macro-identifier" +#pragma clang diagnostic ignored "-Wpadded" + #if __cplusplus #if !defined(__HIPCC_RTC__) #include @@ -69,8 +76,8 @@ class thread_group { // only when the group is supposed to contain only the calling the thread // (throurh the API - `this_thread()`), and in all other cases, this thread // group object is a sub-object of some other derived thread group object - __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size = (uint64_t)0, - uint64_t mask = (uint64_t)0) { + __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size = static_cast(0), + uint64_t mask = static_cast(0)) { _type = type; _size = size; _mask = mask; @@ -199,7 +206,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); @@ -246,7 +253,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) { @@ -282,7 +289,6 @@ class tiled_group : public thread_group { * \details Represents a active thread group in a wavefront. * This group type also supports sub-wave level intrinsics. */ - class coalesced_group : public thread_group { private: friend __CG_QUALIFIER__ coalesced_group coalesced_threads(); @@ -300,8 +306,8 @@ class coalesced_group : public thread_group { // prepare a mask for further partitioning it so that it stays coalesced. if (coalesced_info.tiled_info.is_tiled) { unsigned int base_offset = (thread_rank() & (~(tile_size - 1))); - unsigned int masklength = min((unsigned int)size() - base_offset, tile_size); - lane_mask member_mask = (lane_mask)(-1) >> (__AMDGCN_WAVEFRONT_SIZE - masklength); + unsigned int masklength = min(static_cast(size()) - base_offset, tile_size); + lane_mask member_mask = static_cast(-1) >> (__AMDGCN_WAVEFRONT_SIZE - masklength); member_mask <<= (__lane_id() & ~(tile_size - 1)); coalesced_group coalesced_tile = coalesced_group(member_mask); @@ -358,7 +364,7 @@ class coalesced_group : public thread_group { __CG_QUALIFIER__ T shfl(T var, int srcRank) const { static_assert(is_valid_type::value, "Neither an integer or float type."); - srcRank = srcRank % size(); + srcRank = srcRank % static_cast(size()); int lane = (size() == __AMDGCN_WAVEFRONT_SIZE) ? srcRank : (__AMDGCN_WAVEFRONT_SIZE == 64) ? __fns64(coalesced_info.member_mask, 0, (srcRank + 1)) @@ -452,7 +458,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; } } @@ -476,7 +482,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; } } @@ -505,7 +511,7 @@ __CG_QUALIFIER__ void thread_group::sync() const { break; } default: { - __hip_assert(false && "invalid cooperative group type"); + __hip_assert(false && "invalid cooperative group type") } } } @@ -697,6 +703,6 @@ __CG_QUALIFIER__ thread_block_tile tiled_partition(const Paren return impl::tiled_partition_internal(g); } } // namespace cooperative_groups - +#pragma clang diagnostic pop #endif // __cplusplus #endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h index 243b5a64d9..f9e9738a8c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 - Present Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -23,6 +23,9 @@ THE SOFTWARE. #pragma once #ifdef __cplusplus +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wold-style-cast" + /** * @brief Unsafe floating point rmw atomic add. * @@ -563,4 +566,5 @@ __device__ inline double safeAtomicMin(double* addr, double val) { #endif } +#pragma clang diagnostic pop #endif 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 b18ff5f5d7..fb6065b187 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 @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -23,6 +23,14 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_WARP_FUNCTIONS_H #define HIP_INCLUDE_HIP_AMD_DETAIL_WARP_FUNCTIONS_H +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wreserved-identifier" +#pragma clang diagnostic ignored "-Wreserved-macro-identifier" +#pragma clang diagnostic ignored "-Wsign-conversion" +#pragma clang diagnostic ignored "-Wold-style-cast" +#pragma clang diagnostic ignored "-Wc++98-compat" +#pragma clang diagnostic ignored "-Wc++98-compat-pedantic" + __device__ static inline unsigned __hip_ds_bpermute(int index, unsigned src) { union { int i; unsigned u; float f; } tmp; tmp.u = src; tmp.i = __builtin_amdgcn_ds_bpermute(index, tmp.i); @@ -491,4 +499,5 @@ unsigned long long __shfl_xor(unsigned long long var, int lane_mask, int width = return tmp1; } +#pragma clang diagnostic pop #endif diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index a90f0a3a78..877c6a43b5 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -39,6 +39,12 @@ THE SOFTWARE. #define __align__(x) __attribute__((aligned(x))) #endif +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wreserved-macro-identifier" +#pragma clang diagnostic ignored "-Wc++98-compat" +#pragma clang diagnostic ignored "-Wc++98-compat-pedantic" +#pragma clang diagnostic ignored "-Wshorten-64-to-32" + #if !defined(__CG_QUALIFIER__) #define __CG_QUALIFIER__ __device__ __forceinline__ #endif @@ -92,15 +98,18 @@ typedef enum { */ namespace multi_grid { -__CG_STATIC_QUALIFIER__ uint32_t num_grids() { return (uint32_t)__ockl_multi_grid_num_grids(); } +__CG_STATIC_QUALIFIER__ uint32_t num_grids() { + return static_cast(__ockl_multi_grid_num_grids()); } -__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { return (uint32_t)__ockl_multi_grid_grid_rank(); } +__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { + return static_cast(__ockl_multi_grid_grid_rank()); } -__CG_STATIC_QUALIFIER__ uint32_t size() { return (uint32_t)__ockl_multi_grid_size(); } +__CG_STATIC_QUALIFIER__ uint32_t size() { return static_cast(__ockl_multi_grid_size()); } -__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { return (uint32_t)__ockl_multi_grid_thread_rank(); } +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { + return static_cast(__ockl_multi_grid_thread_rank()); } -__CG_STATIC_QUALIFIER__ bool is_valid() { return (bool)__ockl_multi_grid_is_valid(); } +__CG_STATIC_QUALIFIER__ bool is_valid() { return static_cast(__ockl_multi_grid_is_valid()); } __CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); } @@ -112,28 +121,28 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); } namespace grid { __CG_STATIC_QUALIFIER__ uint32_t size() { - return (uint32_t)((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) * + return static_cast((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) * (blockDim.x * gridDim.x)); } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { // Compute global id of the workgroup to which the current thread belongs to - uint32_t blkIdx = (uint32_t)((blockIdx.z * gridDim.y * gridDim.x) + + uint32_t blkIdx = static_cast((blockIdx.z * gridDim.y * gridDim.x) + (blockIdx.y * gridDim.x) + (blockIdx.x)); // Compute total number of threads being passed to reach current workgroup // within grid uint32_t num_threads_till_current_workgroup = - (uint32_t)(blkIdx * (blockDim.x * blockDim.y * blockDim.z)); + static_cast(blkIdx * (blockDim.x * blockDim.y * blockDim.z)); // Compute thread local rank within current workgroup - uint32_t local_thread_rank = (uint32_t)((threadIdx.z * blockDim.y * blockDim.x) + + uint32_t local_thread_rank = static_cast((threadIdx.z * blockDim.y * blockDim.x) + (threadIdx.y * blockDim.x) + (threadIdx.x)); return (num_threads_till_current_workgroup + local_thread_rank); } -__CG_STATIC_QUALIFIER__ bool is_valid() { return (bool)__ockl_grid_is_valid(); } +__CG_STATIC_QUALIFIER__ bool is_valid() { return static_cast(__ockl_grid_is_valid()); } __CG_STATIC_QUALIFIER__ void sync() { __ockl_grid_sync(); } @@ -146,19 +155,21 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_grid_sync(); } namespace workgroup { __CG_STATIC_QUALIFIER__ dim3 group_index() { - return (dim3((uint32_t)blockIdx.x, (uint32_t)blockIdx.y, (uint32_t)blockIdx.z)); + return (dim3(static_cast(blockIdx.x), static_cast(blockIdx.y), + static_cast(blockIdx.z))); } __CG_STATIC_QUALIFIER__ dim3 thread_index() { - return (dim3((uint32_t)threadIdx.x, (uint32_t)threadIdx.y, (uint32_t)threadIdx.z)); + return (dim3(static_cast(threadIdx.x), static_cast(threadIdx.y), + static_cast(threadIdx.z))); } __CG_STATIC_QUALIFIER__ uint32_t size() { - return ((uint32_t)(blockDim.x * blockDim.y * blockDim.z)); + return (static_cast(blockDim.x * blockDim.y * blockDim.z)); } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { - return ((uint32_t)((threadIdx.z * blockDim.y * blockDim.x) + + return (static_cast((threadIdx.z * blockDim.y * blockDim.x) + (threadIdx.y * blockDim.x) + (threadIdx.x))); } @@ -187,8 +198,8 @@ __CG_STATIC_QUALIFIER__ void sync() { __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, " // // For each thread, this function returns the number of active threads which // have i-th bit of x set and come before the current thread. -__device__ unsigned int masked_bit_count(lane_mask x, unsigned int add = 0) { - int counter=0; +__CG_STATIC_QUALIFIER__ unsigned int masked_bit_count(lane_mask x, unsigned int add = 0) { + unsigned int counter=0; #if __AMDGCN_WAVEFRONT_SIZE == 32 counter = __builtin_amdgcn_mbcnt_lo(x, add); #else @@ -206,5 +217,6 @@ __device__ unsigned int masked_bit_count(lane_mask x, unsigned int add = 0) { } // namespace cooperative_groups +#pragma clang diagnostic pop #endif // __cplusplus #endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H diff --git a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake index 19b59f3655..521b02345d 100644 --- a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -1,4 +1,4 @@ -# Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2021 - 2023 Advanced Micro Devices, Inc. All Rights Reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -29,6 +29,7 @@ function(get_hiprtc_macros HIPRTC_DEFINES) set(${HIPRTC_DEFINES} "#pragma clang diagnostic push\n\ #pragma clang diagnostic ignored \"-Wreserved-id-macro\"\n\ +#pragma clang diagnostic ignored \"-Wc++98-compat-pedantic\"\n\ #define __device__ __attribute__((device))\n\ #define __host__ __attribute__((host))\n\ #define __global__ __attribute__((global))\n\ @@ -51,7 +52,10 @@ function(get_hiprtc_macros HIPRTC_DEFINES) #pragma clang diagnostic pop\n\ #define HIP_INCLUDE_HIP_HIP_RUNTIME_H\n\ #define HIP_INCLUDE_HIP_HIP_FP16_H\n\ +#pragma clang diagnostic push\n\ +#pragma clang diagnostic ignored \"-Wreserved-macro-identifier\"\n\ #define _HIP_BFLOAT16_H_\n\ +#pragma clang diagnostic pop\n\ #define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H" PARENT_SCOPE) endfunction(get_hiprtc_macros) @@ -63,12 +67,15 @@ if(HIPRTC_ADD_MACROS) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_DEFINES}") FILE(READ "${HIPRTC_WARP_HEADER_FILE}" HIPRTC_WARP_HEADER) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_WARP_HEADER}") +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wreserved-macro-identifier" FILE(READ "${HIPRTC_COOP_HELPER_FILE}" HIPRTC_COOP_HELPER) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_COOP_HELPER}") FILE(READ "${HIPRTC_COOP_HEADER_FILE}" HIPRTC_COOP_HEADER) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_COOP_HEADER}") FILE(READ "${HIPRTC_UNSAFE_ATOMICS_FILE}" HIPRTC_UNSAFE_ATOMICS) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_UNSAFE_ATOMICS}") +#pragma clang diagnostic pop endif() macro(generate_hiprtc_header HiprtcHeader) diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp index dfff262cc9..e4e4f04460 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 - Present Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 - 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -135,6 +135,8 @@ RTCCompileProgram::RTCCompileProgram(std::string name_) : RTCProgram(name_), fgp compile_options_.push_back("hiprtc_runtime.h"); compile_options_.push_back("-std=c++14"); compile_options_.push_back("-nogpuinc"); + compile_options_.push_back("-Wno-gnu-line-marker"); + compile_options_.push_back("-Wno-missing-prototypes"); #ifdef _WIN32 compile_options_.push_back("-target"); compile_options_.push_back("x86_64-pc-windows-msvc");