Do not decay function pointer type of the kernel argument passed to hipLaunchKernelGGL
and hipLaunchKernel, otherwise some type information is lost which may cause
type inference failure for the template.
This issue caused compilation error of FeatureLPPooling in Caffe2/PyTorch and this patch
fixes that.
[ROCm/hip commit: fc228c7ea6]
previous version return a long long valus *as* double, hence we may get the wrong result.
this also affect atomicAdd(double * ...), which use long long pointer to mimic double pointer.
Signed-off-by: carlushuang <carlus.huang@amd.com>
[ROCm/hip commit: d577f27d1a]
Older gcc, e.g., 5.x, see an ambiguity in some calls. Example error
output as seen with gcc 5.5:
In file included from benchmark_wino.cpp:1:
In file included from ./miopen.hpp:13:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/regex:38:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/algorithm:62:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/bits/stl_algo.h:66:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/random:51:
/usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/bits/random.tcc:1324:27: error: call to 'abs' is ambiguous
const double __y = -std::abs(__n) * __param._M_sm - 1;
^~~~~~~~
/opt/rocm/hip/include/hip/hcc_detail/hip_complex.h:345:31: note: candidate function
__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
[ROCm/hip commit: aa7ce2ef44]
Older gcc, e.g., 5.x, see an ambiguity in some calls. Example error
output as seen with gcc 5.5:
In file included from benchmark_wino.cpp:1:
In file included from ./miopen.hpp:13:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/regex:38:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/algorithm:62:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/bits/stl_algo.h:66:
In file included from /usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/random:51:
/usr/lib/gcc/x86_64-linux-gnu/5.5.0/../../../../include/c++/5.5.0/bits/random.tcc:1324:27: error: call to 'abs' is ambiguous
const double __y = -std::abs(__n) * __param._M_sm - 1;
^~~~~~~~
/opt/rocm/hip/include/hip/hcc_detail/hip_complex.h:345:31: note: candidate function
__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
[ROCm/hip commit: f0183dfaac]
hcc_detail/math_functions.h used to include hcc_detail/hip_runtime.h.
Removing it has caused regression in TensorFlow 1.8.
Put it back for backward compatibiliity.
[ROCm/hip commit: b8a0c14171]
Implement the vector value types in terms of SCALAR_TYPE##SIZE_vector_value_type for float, int, uint and sizes 2, 3, 4, 8, and 16.
[ROCm/hip commit: e0e785a45d]
Include texture_functions.h on HIP-Clang path. Use HIP vector types instead of hc::short_vectors for texture_functions, and remove HCC headers.
[ROCm/hip commit: 8c50bcfdab]
Yaxun Sam Liu
Maneesh Gupta
Alex Voicu
Aaron Enye Shi
Johannes M Dieterich
carlushuang
lcskrishna
Mark Searles
Wen-Heng (Jack) Chung
Brian Sumner