HIP based binary are not building with recent rocm release 1.7, Any modification in source is required?

[pramenku]

Hi ekondis
I tried building HIP based binary mixbench-hip-alt and mixbench-hip-ro after installing rocm driver from https://rocm.github.io/ROCmInstall.html. I am getting below error.
Is any modification is required with latest HIP changes as mentioned in ROCm/HIP#246 ?

Changes to hipLaunchKernelGGL and HCC introduces stricter checks on use of templates in kernel launches. Samples that rely on templates and that do not conform to proper syntax require code modifications.

$ make
echo '#define VERSION_INFO "'./query_version.sh'"' >version_info.h
/opt/rocm/hip/bin/hipcc -c -O2 -I/opt/rocm/hip/hip/include main-hip.cpp
/opt/rocm/hip/bin/hipcc -O2 -I/opt/rocm/hip/hip/include -DUNIX -c mix_kernels_hip.cpp -o mix_kernels_hip.o
mix_kernels_hip.cpp:99:19: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:130:34: note: in instantiation of function template specialization 'benchmark_func<short, 256, 0>'
requested here
hipLaunchKernel(HIP_KERNEL_NAME(benchmark_func< short, BLOCK_SIZE, 0 >), dim3(dimReducedGrid), dim3(dimBlock )...
^
mix_kernels_hip.cpp:99:47: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:99:75: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:99:103: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:19: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:47: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:75: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:103: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
...is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)) ){ // extr...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:51:76: warning: remainder by zero is undefined [-Wdivision-by-zero]
const int initial_index_range = memory_ratio>0 ? UNROLLED_MEMORY_ACCESSES % ((memory_ratio+1)/2) : 1;
^ ~~~~~~~~~~~~~~~~~~~~
mix_kernels_hip.cpp:99:19: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:153:34: note: in instantiation of function template specialization 'benchmark_func<float, 256, 32>'
requested here
hipLaunchKernel(HIP_KERNEL_NAME(benchmark_func< float, BLOCK_SIZE, memory_ratio >), dim3(dimGrid), dim3(dimBloc...
^
mix_kernels_hip.cpp:211:2: note: in instantiation of function template specialization 'runbench<32>' requested here
runbench<32>(cd, size);
^
mix_kernels_hip.cpp:99:47: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:99:75: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:99:103: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:19: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:47: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:75: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:100:103: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
...is_equal(r4, GPU_INF(T)) && is_equal(r5, GPU_INF(T)) && is_equal(r6, GPU_INF(T)) && is_equal(r7, GPU_INF(T)) ){ // extr...
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:99:19: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:157:34: note: in instantiation of function template specialization 'benchmark_func<double, 256, 32>'
requested here
hipLaunchKernel(HIP_KERNEL_NAME(benchmark_func< double, BLOCK_SIZE, memory_ratio >), dim3(dimGrid), dim3(dimBlo...
^
mix_kernels_hip.cpp:211:2: note: in instantiation of function template specialization 'runbench<32>' requested here
runbench<32>(cd, size);
^
mix_kernels_hip.cpp:99:47: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
mix_kernels_hip.cpp:99:75: error: 'std::numeric_limits::infinity': no overloaded function has restriction specifiers
that are compatible with the ambient context 'benchmark_func'
if( is_equal(r0, GPU_INF(T)) && is_equal(r1, GPU_INF(T)) && is_equal(r2, GPU_INF(T)) && is_equal(r3, GPU_INF(T)) &&
^
mix_kernels_hip.cpp:16:23: note: expanded from macro 'GPU_INF'
#define GPU_INF(_T) std::numeric_limits<_T>::infinity()
^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
1 warning and 20 errors generated.
Died at /opt/rocm/hip/bin/hipcc line 498.
Makefile:112: recipe for target 'mix_kernels_hip.o' failed
make: *** [mix_kernels_hip.o] Error 1

[ekondis]

mixbench-hip-ro compiles and runs OK. The problem stems from std::numeric_limits<_T>::infinity().

This is not an critical part of the code though. As a workaround you may replace GPU_INF definition by an arbitrary value, e.g. 9999. This is not essential for this benchmark.

[pramenku]

Thanks ekondis . Are you going to put workaround in source also so that user need not to do change manually after git clone the source.

[ekondis]

Perhaps, but I would try to avoid such a workaround.
Just posted new issue for HIP: ROCm/HIP#374

[pramenku]

Thanks ekondis. It's always good to avoid workaround. :)