Add device function version for HIP platform.

[aaronenyeshi]

I'm running into an error with the latest rocPRIM develop branch on HIP-Clang. The error is due to function half_to_native in test_utils.hpp which doesn't have a host function but it being called by host function.

In file included from /work/rocprim/test/rocprim/test_hip_device_scan.cpp:35:
/work/rocprim/test/rocprim/test_utils.hpp:606:19: error: no matching function for call to 'half_to_native'
        ASSERT_EQ(half_to_native(result[i]), half_to_native(expected[i])) << "where index = " << i;
                  ^~~~~~~~~~~~~~
/work/rocprim/build/gtest/include/gtest/gtest.h:2078:48: note: expanded from macro 'ASSERT_EQ'
# define ASSERT_EQ(val1, val2) GTEST_ASSERT_EQ(val1, val2)
                                               ^~~~
/work/rocprim/build/gtest/include/gtest/gtest.h:2061:55: note: expanded from macro 'GTEST_ASSERT_EQ'
                      EqHelper<GTEST_IS_NULL_LITERAL_(val1)>::Compare, \
                                                      ^~~~
/work/rocprim/build/gtest/include/gtest/internal/gtest-internal.h:155:7: note: expanded from macro 'GTEST_IS_NULL_LITERAL_'
      x,                                             \
      ^
/work/rocprim/build/gtest/include/gtest/gtest_pred_impl.h:168:23: note: expanded from macro 'ASSERT_PRED_FORMAT2'
  GTEST_PRED_FORMAT2_(pred_format, v1, v2, GTEST_FATAL_FAILURE_)
                      ^~~~~~~~~~~
/work/rocprim/build/gtest/include/gtest/gtest_pred_impl.h:149:17: note: expanded from macro 'GTEST_PRED_FORMAT2_'
  GTEST_ASSERT_(pred_format(#v1, #v2, v1, v2), \
                ^~~~~~~~~~~
/work/rocprim/build/gtest/include/gtest/gtest_pred_impl.h:77:52: note: expanded from macro 'GTEST_ASSERT_'
  if (const ::testing::AssertionResult gtest_ar = (expression)) \
                                                   ^~~~~~~~~~
/work/rocprim/test/rocprim/test_utils.hpp:57:15: note: candidate function not viable: call to __device__ function from __host__ function
rocprim::half half_to_native(const rocprim::half& x)
              ^

As you can see in the final line, the host versions of half_to_native is removed due to the #else clause. In HIP-Clang its very strict (same as CUDA) where host functions cannot call functions which have only device attribute. Is it possible to use ROCPRIM_HOST_DEVICE instead on lines 56 and 62 of test_utils.hpp?

[aaronenyeshi]

Here is another function which also has the host and device attribute mismatch:

In file included from /work/rocprim/test/rocprim/test_hip_device_scan.cpp:33:
In file included from /work/rocprim/rocprim/include/rocprim/rocprim.hpp:50:
In file included from /work/rocprim/rocprim/include/rocprim/device/../block/block_store.hpp:31:
/work/rocprim/rocprim/include/rocprim/block/block_store_func.hpp:269:36: error: reference to __host__ function 'operator float<float, nullptr>' in __device__ function
             thread_iter[offset] = items[item];
                                   ^
/work/rocprim/rocprim/include/rocprim/device/../block/block_store.hpp:434:9: note: in instantiation of function template specialization 'rocprim::block_store_direct_striped<256, test_utils::bounds_checking_iterator<float>, __half, 16>' requested here
        block_store_direct_striped<BlockSize>(flat_id, block_output, items, valid);
        ^
/work/rocprim/rocprim/include/rocprim/device/detail/device_scan_reduce_then_scan.hpp:406:14: note: in instantiation of function template specialization 'rocprim::block_store<__half, 256, 16, rocprim::block_store_method::block_store_transpose>::store<test_utils::bounds_checking_iterator<float> >' requested here
            .store(
             ^
/work/rocprim/rocprim/include/rocprim/device/device_scan_hip.hpp:102:5: note: in instantiation of function template specialization 'rocprim::detail::final_scan_kernel_impl<false, rocprim::detail::default_scan_config<0, __half>, rocprim::constant_iterator<__half, long>, test_utils::bounds_checking_iterator<float>, rocprim::plus<float>, __half>' requested here
    final_scan_kernel_impl<Exclusive, Config>(
    ^
/work/rocprim/rocprim/include/rocprim/device/device_scan_hip.hpp:273:37: note: in instantiation of function template specialization 'rocprim::detail::final_scan_kernel<false, rocprim::detail::default_scan_config<0, __half>, rocprim::constant_iterator<__half, long>, test_utils::bounds_checking_iterator<float>, rocprim::plus<float>, __half>' requested here
            HIP_KERNEL_NAME(detail::final_scan_kernel<
                                    ^
/work/rocprim/rocprim/include/rocprim/device/device_scan_hip.hpp:524:20: note: in instantiation of function template specialization 'rocprim::detail::scan_impl<false, false, rocprim::default_config, rocprim::constant_iterator<__half, long>, test_utils::bounds_checking_iterator<float>, __half, rocprim::plus<float> >' requested here
    return detail::scan_impl<false, std::is_arithmetic<result_type>::value, Config>(
                   ^
/work/rocprim/test/rocprim/test_hip_device_scan.cpp:127:18: note: (skipping 7 contexts in backtrace; use -ftemplate-backtrace-limit=0 to see all)
        rocprim::inclusive_scan(
                 ^
/work/rocprim/build/gtest/include/gtest/internal/gtest-internal.h:747:57: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<RocprimDeviceScanTests, testing::internal::TemplateSel<RocprimDeviceScanTests_InclusiveScanEmptyInput_Test>, testing::internal::Types5<DeviceScanParams<signed char, long, rocprim::plus<long>, false>, DeviceScanParams<float, double, rocprim::minimum<double>, false>, DeviceScanParams<test_utils::custom_test_type<double>, test_utils::custom_test_type<double>, rocprim::plus<test_utils::custom_test_type<double> >, true>, DeviceScanParams<__half, __half, test_utils::half_maximum, false>, DeviceScanParams<__half, float, rocprim::plus<float>, false> > >::Register' requested here
                                 typename Types::Tail>::Register(prefix,
                                                        ^
/work/rocprim/build/gtest/include/gtest/internal/gtest-internal.h:747:57: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<RocprimDeviceScanTests, testing::internal::TemplateSel<RocprimDeviceScanTests_InclusiveScanEmptyInput_Test>, testing::internal::Types6<DeviceScanParams<short, int, rocprim::plus<int>, false>, DeviceScanParams<signed char, long, rocprim::plus<long>, false>, DeviceScanParams<float, double, rocprim::minimum<double>, false>, DeviceScanParams<test_utils::custom_test_type<double>, test_utils::custom_test_type<double>, rocprim::plus<test_utils::custom_test_type<double> >, true>, DeviceScanParams<__half, __half, test_utils::half_maximum, false>, DeviceScanParams<__half, float, rocprim::plus<float>, false> > >::Register' requested here
/work/rocprim/build/gtest/include/gtest/internal/gtest-internal.h:747:57: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<RocprimDeviceScanTests, testing::internal::TemplateSel<RocprimDeviceScanTests_InclusiveScanEmptyInput_Test>, testing::internal::Types7<DeviceScanParams<double, double, rocprim::plus<double>, true>, DeviceScanParams<short, int, rocprim::plus<int>, false>, DeviceScanParams<signed char, long, rocprim::plus<long>, false>, DeviceScanParams<float, double, rocprim::minimum<double>, false>, DeviceScanParams<test_utils::custom_test_type<double>, test_utils::custom_test_type<double>, rocprim::plus<test_utils::custom_test_type<double> >, true>, DeviceScanParams<__half, __half, test_utils::half_maximum, false>, DeviceScanParams<__half, float, rocprim::plus<float>, false> > >::Register' requested here
/work/rocprim/build/gtest/include/gtest/internal/gtest-internal.h:747:57: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<RocprimDeviceScanTests, testing::internal::TemplateSel<RocprimDeviceScanTests_InclusiveScanEmptyInput_Test>, testing::internal::Types8<DeviceScanParams<int, int, rocprim::plus<int>, false>, DeviceScanParams<double, double, rocprim::plus<double>, true>, DeviceScanParams<short, int, rocprim::plus<int>, false>, DeviceScanParams<signed char, long, rocprim::plus<long>, false>, DeviceScanParams<float, double, rocprim::minimum<double>, false>, DeviceScanParams<test_utils::custom_test_type<double>, test_utils::custom_test_type<double>, rocprim::plus<test_utils::custom_test_type<double> >, true>, DeviceScanParams<__half, __half, test_utils::half_maximum, false>, DeviceScanParams<__half, float, rocprim::plus<float>, false> > >::Register' requested here
/work/rocprim/test/rocprim/test_hip_device_scan.cpp:103:1: note: in instantiation of member function 'testing::internal::TypeParameterizedTest<RocprimDeviceScanTests, testing::internal::TemplateSel<RocprimDeviceScanTests_InclusiveScanEmptyInput_Test>, testing::internal::Types9<DeviceScanParams<unsigned short, unsigned short, rocprim::plus<unsigned short>, false>, DeviceScanParams<int, int, rocprim::plus<int>, false>, DeviceScanParams<double, double, rocprim::plus<double>, true>, DeviceScanParams<short, int, rocprim::plus<int>, false>, DeviceScanParams<signed char, long, rocprim::plus<long>, false>, DeviceScanParams<float, double, rocprim::minimum<double>, false>, DeviceScanParams<test_utils::custom_test_type<double>, test_utils::custom_test_type<double>, rocprim::plus<test_utils::custom_test_type<double> >, true>, DeviceScanParams<__half, __half, test_utils::half_maximum, false>, DeviceScanParams<__half, float, rocprim::plus<float>, false> > >::Register' requested here
TYPED_TEST(RocprimDeviceScanTests, InclusiveScanEmptyInput)
^
/work/rocprim/build/gtest/include/gtest/gtest-typed-test.h:216:27: note: expanded from macro 'TYPED_TEST'
              CaseName)>::Register("",                                        \
                          ^
/opt/rocm/hip/include/hip/hcc_detail/hip_fp16.h:221:17: note: 'operator float<float, nullptr>' declared here
                operator T() const { return data; }
                ^

[ex-rzr]

As you can see in the final line, the host versions of half_to_native is removed due to the #else clause.

But the #else clause must be executed. Does HIP-Clang define __HCC_ACCELERATOR__ or __HIP_DEVICE_COMPILE__ during host code compilation?

2. This looks like a bug in HIP

https://github.com/ROCm-Developer-Tools/HIP/blob/master/include/hip/hcc_detail/hip_fp16.h#L221

At least conversion to integral types is a __device__ function.

[aaronenyeshi]

@ex-rzr Could you help report any HIP bugs to HIP Issues? I am looking into why the host code compilation is failing in this case. Thank you

[ex-rzr]

@aaronenyeshi, sorry for the delay. I've never used HIP-clang before. Could you provide instructions where to get and how to build it?

[aaronenyeshi]

So far, it is only internal. I will discuss with my team. Thanks

[aaronenyeshi]

@ex-rzr What email can I contact you at and discuss this? Thanks

[ex-rzr]

#63 fixes this (in develop branch).

There are other changes related to HIP-clang support.

See this change: d259b36#diff-8f4678af31293d182b2d4028abadd827