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

SWDEV-372396 - Address the errors in hiprtc headers with -Werror

Browse Source 
Change-Id: Ied553216dfc2ff8b5c72de617cb47b6ca8f52063


[ROCm/clr commit: 0ccd41050b]
This commit is contained in:
Satyanvesh Dittakavi
2022-12-20 11:56:43 +00:00
parent 7b3cf2e9d7
commit d96114fd22
7 changed files with 78 additions and 36 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
+4 -2
View File
@@ -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);
projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
+19 -13
View File
@@ -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 <hip/amd_detail/hip_cooperative_groups_helper.h>
@@ -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<uint64_t>(0),
uint64_t mask = static_cast<uint64_t>(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<unsigned int>(size()) - base_offset, tile_size);
lane_mask member_mask = static_cast<lane_mask>(-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<T>::value, "Neither an integer or float type.");
srcRank = srcRank % size();
srcRank = srcRank % static_cast<int>(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<const coalesced_group*>(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<const coalesced_group*>(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<size, ParentCGTy> tiled_partition(const Paren
return impl::tiled_partition_internal<size, ParentCGTy>(g);
}
} // namespace cooperative_groups
#pragma clang diagnostic pop
#endif // __cplusplus
#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H
projects/clr/hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h
+5 -1
View File
@@ -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
projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h
+10 -1
View File
@@ -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
projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+29 -17
View File
@@ -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<uint32_t>(__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<uint32_t>(__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<uint32_t>(__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<uint32_t>(__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<bool>(__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<uint32_t>((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<uint32_t>((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<uint32_t>(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<uint32_t>((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<bool>(__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<uint32_t>(blockIdx.x), static_cast<uint32_t>(blockIdx.y),
static_cast<uint32_t>(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<uint32_t>(threadIdx.x), static_cast<uint32_t>(threadIdx.y),
static_cast<uint32_t>(threadIdx.z)));
}
__CG_STATIC_QUALIFIER__ uint32_t size() {
return ((uint32_t)(blockDim.x * blockDim.y * blockDim.z));
return (static_cast<uint32_t>(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<uint32_t>((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
projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake
+8 -1
View File
@@ -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)
projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp
+3 -1
View File
@@ -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");