saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-30 20:18:08 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Check whether nvcc supports -Wstrict-aliasing before adding the flag. (#7509)

Browse source 
* Check whether nvcc supports -Wstrict-aliasing before adding the compiler flag in CMakeList.txt.

* Removed reinterpret_cast to not cause strict aliasing violation errors or require -Wno-strict-aliasing when it is not available.
This commit is contained in:
satyajandhyala 2021-05-01 00:14:50 -07:00 committed by GitHub
parent 00882ce495
commit 9f1e61be92
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
3 changed files with 31 additions and 9 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

14
cmake/CMakeLists.txt
View file

@ -446,6 +446,17 @@ endif()
#must after OpenMP settings
find_package(Threads)
macro(check_nvcc_compiler_flag _FLAG _RESULT)
execute_process(COMMAND ${onnxruntime_CUDA_HOME}/bin/nvcc "${_FLAG}" RESULT_VARIABLE NVCC_OUT ERROR_VARIABLE NVCC_ERROR)
message("NVCC_ERROR = ${NVCC_ERROR}")
message("NVCC_OUT = ${NVCC_OUT}")
if("${NVCC_OUT}" MATCHES "0")
set(${_RESULT} 1)
else()
set(${_RESULT} 0)
endif()
endmacro()
#Set global compile flags for all the source code(including third_party code like protobuf)
#This section must be before any add_subdirectory, otherwise build may fail because /MD,/MT mismatch
if (MSVC)
@ -874,6 +885,7 @@ else()
check_cxx_compiler_flag(-Wclass-memaccess HAS_CLASS_MEMACCESS)
check_cxx_compiler_flag(-Wmaybe-uninitialized HAS_MAYBE_UNINITIALIZED)
check_cxx_compiler_flag(-Wstrict-aliasing HAS_STRICT_ALIASING)
check_nvcc_compiler_flag(-Wstrict-aliasing NVCC_HAS_STRICT_ALIASING)
if(HAS_TAUTOLOGICAL_POINTER_COMPARE)
#we may have extra null pointer checkings in debug build, it's not an issue
@ -1001,7 +1013,7 @@ function(onnxruntime_set_compile_flags target_name)
target_compile_options(${target_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:--compiler-options -Wno-deprecated-copy>" "$<$<COMPILE_LANGUAGE:CXX>:-Wno-deprecated-copy>")
endif()
if(onnxruntime_USE_CUDA)
if(HAS_STRICT_ALIASING)
if((NVCC_HAS_STRICT_ALIASING AND "${target_name}" MATCHES "cuda") OR (HAS_STRIC_ALIASING AND NOT "${target_name}" MATCHES "cuda"))
target_compile_options(${target_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:-Wno-strict-aliasing>")
endif()
endif()

2
include/onnxruntime/core/platform/threadpool.h
View file

@ -314,7 +314,7 @@ class ThreadPool {
}
if (num_batches <= 0) {
num_batches = std::min<ptrdiff_t>(total, DegreeOfParallelism(tp));
num_batches = std::min<std::ptrdiff_t>(total, DegreeOfParallelism(tp));
}
if (num_batches <= 1) {

24
onnxruntime/core/providers/cuda/math/clip_impl.cu
View file

@ -19,13 +19,23 @@ void ClipImpl(cudaStream_t stream, const T* input_data, T* output_data, const T*
typedef typename ToCudaType<T>::MappedType CudaT;
int blocksPerGrid = (int)(ceil(static_cast<float>(count) / GridDim::maxThreadsPerBlock));
_Clip<CudaT><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(reinterpret_cast<const CudaT*>(input_data),
reinterpret_cast<CudaT*>(output_data),
reinterpret_cast<const CudaT*>(min),
reinterpret_cast<const CudaT*>(max),
*reinterpret_cast<CudaT*>(&min_default),
*reinterpret_cast<CudaT*>(&max_default),
count);
union ConstAliasUnion {
const T *t;
const CudaT *cudaT;
ConstAliasUnion(const T* _t) { t = _t;}
};
union AliasUnion {
T *t;
CudaT *cudaT;
AliasUnion(T* _t) { t = _t;}
};
_Clip<CudaT><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(((union ConstAliasUnion)input_data).cudaT,
((union AliasUnion)output_data).cudaT,
((union ConstAliasUnion)min).cudaT,
((union ConstAliasUnion)max).cudaT,
*((union AliasUnion)&min_default).cudaT,
*((union AliasUnion)&max_default).cudaT,
count);
}
template void ClipImpl<float>(cudaStream_t stream, const float* input_data, float* output_data, const float* min, const float* max, float min_default, float max_default, size_t count);